diff --git a/.github/workflows/musa-ci.yml b/.github/workflows/musa-ci.yml
new file mode 100644
index 0000000..69c58c5
--- /dev/null
+++ b/.github/workflows/musa-ci.yml
@@ -0,0 +1,259 @@
+name: MUSA CI
+
+on:
+  pull_request:
+  workflow_dispatch:
+
+permissions:
+  contents: read
+  actions: read
+
+concurrency:
+  group: musa-ci-${{ github.event.pull_request.number || github.ref }}
+  cancel-in-progress: true
+
+jobs:
+  test:
+    name: Unit Tests
+    runs-on: [self-hosted, musa, s5000]
+    timeout-minutes: 90
+    env:
+      CI_ASSETS: /data/ci-assets
+      PYTHONUNBUFFERED: "1"
+
+    steps:
+      - name: Checkout PR head
+        uses: actions/checkout@v4
+        with:
+          ref: ${{ github.event.pull_request.head.sha || github.sha }}
+          fetch-depth: 1
+          submodules: false
+
+      - name: Configure MUSA toolkit environment
+        run: |
+          set -euo pipefail
+          echo "/usr/local/musa/bin" >> "${GITHUB_PATH}"
+          {
+            echo "MUSA_HOME=/usr/local/musa"
+            echo "LD_LIBRARY_PATH=/usr/local/musa/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}"
+            echo "CPLUS_INCLUDE_PATH=/usr/include/c++/11:/usr/include/x86_64-linux-gnu/c++/11:/usr/include${CPLUS_INCLUDE_PATH:+:${CPLUS_INCLUDE_PATH}}"
+          } >> "${GITHUB_ENV}"
+
+          echo "MUSA_HOME=/usr/local/musa"
+          echo "PATH includes /usr/local/musa/bin in subsequent steps"
+          echo "LD_LIBRARY_PATH=/usr/local/musa/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}"
+          echo "CPLUS_INCLUDE_PATH=/usr/include/c++/11:/usr/include/x86_64-linux-gnu/c++/11:/usr/include${CPLUS_INCLUDE_PATH:+:${CPLUS_INCLUDE_PATH}}"
+
+      - name: Prepare ONNX Runtime headers
+        run: |
+          set -euo pipefail
+          test -d "${CI_ASSETS}/onnxruntime"
+          rm -rf third_party/onnxruntime
+          cp -a "${CI_ASSETS}/onnxruntime" third_party/onnxruntime
+
+      - name: Build plugin and wheel
+        run: |
+          set -euo pipefail
+          echo "Testing commit: $(git rev-parse HEAD)"
+          bash build.sh
+
+      - name: Prepare ONNX Runtime SDK
+        run: |
+          set -euo pipefail
+          test -d "${CI_ASSETS}/ort-sdk"
+          rm -rf build/ort-sdk
+          cp -a "${CI_ASSETS}/ort-sdk" build/ort-sdk
+
+      - name: Install wheel
+        run: |
+          set -euo pipefail
+          python -m pip install --no-deps "onnxruntime==1.26.0"
+
+          shopt -s nullglob
+          wheels=(dist/onnxruntime_ep_musa-*.whl)
+          if (( ${#wheels[@]} != 1 )); then
+            echo "Expected exactly one MUSA wheel, found ${#wheels[@]}" >&2
+            printf '  %s\n' "${wheels[@]}" >&2
+            exit 1
+          fi
+          python -m pip install --no-deps --force-reinstall "${wheels[0]}"
+
+      - name: Run provider tests
+        run: bash test/run_all.sh
+
+      - name: Package build artifacts
+        run: |
+          set -euo pipefail
+          mkdir -p "${RUNNER_TEMP}/musa-build"
+          tar -C build -czf "${RUNNER_TEMP}/musa-build/musa-build.tar.gz" \
+            Release/libonnxruntime_providers_musa_plugin.so \
+            ort-sdk
+          tar -tzf "${RUNNER_TEMP}/musa-build/musa-build.tar.gz" >/dev/null
+
+      - name: Upload build artifacts
+        uses: actions/upload-artifact@v4
+        with:
+          name: musa-build-${{ github.run_id }}-${{ github.run_attempt }}
+          path: ${{ runner.temp }}/musa-build/musa-build.tar.gz
+          if-no-files-found: error
+          retention-days: 7
+
+  model-zoo:
+    name: Model Zoo Inference Performance Test
+    needs: test
+    runs-on: [self-hosted, musa, s5000]
+    timeout-minutes: 720
+    env:
+      CI_ASSETS: /data/ci-assets
+      MODEL_ZOO_WORK: /data/ci-work/onnx-model-zoo-${{ github.run_id }}-${{ github.run_attempt }}
+      PERF_RESULTS_DIR: /data/ci-work/onnx-model-zoo-${{ github.run_id }}-${{ github.run_attempt }}/performance-results
+      BASE_SOURCE_DIR: ${{ github.workspace }}/.ci/base-src
+      PYTHONUNBUFFERED: "1"
+
+    steps:
+      - name: Checkout PR head
+        uses: actions/checkout@v4
+        with:
+          ref: ${{ github.event.pull_request.head.sha || github.sha }}
+          fetch-depth: 1
+          submodules: false
+
+      - name: Configure MUSA toolkit environment
+        run: |
+          set -euo pipefail
+          echo "/usr/local/musa/bin" >> "${GITHUB_PATH}"
+          {
+            echo "MUSA_HOME=/usr/local/musa"
+            echo "LD_LIBRARY_PATH=/usr/local/musa/lib${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}"
+            echo "CPLUS_INCLUDE_PATH=/usr/include/c++/11:/usr/include/x86_64-linux-gnu/c++/11:/usr/include${CPLUS_INCLUDE_PATH:+:${CPLUS_INCLUDE_PATH}}"
+          } >> "${GITHUB_ENV}"
+
+      - name: Download build artifacts
+        uses: actions/download-artifact@v4
+        with:
+          name: musa-build-${{ github.run_id }}-${{ github.run_attempt }}
+          path: ${{ runner.temp }}/musa-build
+
+      - name: Extract build artifacts
+        run: |
+          set -euo pipefail
+          mkdir -p build
+          tar -C build -xzf "${RUNNER_TEMP}/musa-build/musa-build.tar.gz"
+
+          plugin="${GITHUB_WORKSPACE}/build/Release/libonnxruntime_providers_musa_plugin.so"
+          sdk_root="${GITHUB_WORKSPACE}/build/ort-sdk/onnxruntime-linux-x64-1.26.0"
+          test -f "${plugin}"
+          test -f "${sdk_root}/include/onnxruntime_c_api.h"
+          echo "Head commit: $(git rev-parse HEAD)"
+          sha256sum "${plugin}"
+
+      - name: Checkout PR base
+        uses: actions/checkout@v4
+        with:
+          ref: ${{ github.event.pull_request.base.sha || github.event.repository.default_branch }}
+          path: .ci/base-src
+          fetch-depth: 1
+          submodules: false
+
+      - name: Build base plugin
+        run: |
+          set -euo pipefail
+          cd "${BASE_SOURCE_DIR}"
+          rm -rf third_party/onnxruntime
+          cp -a "${CI_ASSETS}/onnxruntime" third_party/onnxruntime
+
+          echo "Base commit: $(git rev-parse HEAD)"
+          bash build.sh --clean --no-wheel
+
+          base_plugin="${BASE_SOURCE_DIR}/build/Release/libonnxruntime_providers_musa_plugin.so"
+          test -f "${base_plugin}"
+          sha256sum "${base_plugin}"
+
+      - name: Prepare model zoo
+        run: |
+          set -euo pipefail
+          test -d "${CI_ASSETS}/onnx-model-zoo"
+          rm -rf "${MODEL_ZOO_WORK}"
+          mkdir -p "${MODEL_ZOO_WORK}"
+          cp -al "${CI_ASSETS}/onnx-model-zoo/." "${MODEL_ZOO_WORK}/"
+          rm -rf "${MODEL_ZOO_WORK}/cpp_runners"
+          cp -a "${CI_ASSETS}/onnx-model-zoo/cpp_runners" "${MODEL_ZOO_WORK}/cpp_runners"
+          rm -f "${MODEL_ZOO_WORK}/musa_ep_performance.md"
+          rm -f "${MODEL_ZOO_WORK}/musa_ep_performance_with_multi_stream.md"
+          mkdir -p "${PERF_RESULTS_DIR}"
+          test -f "${MODEL_ZOO_WORK}/JD/real_models/platform/unirank-v2-8layer-combo-v6-online_graph2.onnx"
+
+      - name: Build model zoo runners
+        run: |
+          set -euo pipefail
+          ORT_ROOT="${GITHUB_WORKSPACE}/build/ort-sdk/onnxruntime-linux-x64-1.26.0" \
+            bash "${MODEL_ZOO_WORK}/cpp_runners/build_onnx_runner.sh"
+
+      - name: Benchmark base single-stream
+        run: |
+          set -euo pipefail
+          bash scripts/ci/run_model_zoo_benchmark.sh \
+            base \
+            "${BASE_SOURCE_DIR}/build/Release/libonnxruntime_providers_musa_plugin.so" \
+            single \
+            "${MODEL_ZOO_WORK}" \
+            "${PERF_RESULTS_DIR}"
+
+      - name: Benchmark base multi-stream
+        run: |
+          set -euo pipefail
+          bash scripts/ci/run_model_zoo_benchmark.sh \
+            base \
+            "${BASE_SOURCE_DIR}/build/Release/libonnxruntime_providers_musa_plugin.so" \
+            multi \
+            "${MODEL_ZOO_WORK}" \
+            "${PERF_RESULTS_DIR}"
+
+      - name: Benchmark head single-stream
+        run: |
+          set -euo pipefail
+          bash scripts/ci/run_model_zoo_benchmark.sh \
+            head \
+            "${GITHUB_WORKSPACE}/build/Release/libonnxruntime_providers_musa_plugin.so" \
+            single \
+            "${MODEL_ZOO_WORK}" \
+            "${PERF_RESULTS_DIR}"
+
+      - name: Benchmark head multi-stream
+        run: |
+          set -euo pipefail
+          bash scripts/ci/run_model_zoo_benchmark.sh \
+            head \
+            "${GITHUB_WORKSPACE}/build/Release/libonnxruntime_providers_musa_plugin.so" \
+            multi \
+            "${MODEL_ZOO_WORK}" \
+            "${PERF_RESULTS_DIR}"
+
+      - name: Compare base and head performance
+        run: |
+          set -euo pipefail
+          python3 scripts/ci/compare_model_zoo_performance.py \
+            --baseline-single "${PERF_RESULTS_DIR}/base-single.md" \
+            --candidate-single "${PERF_RESULTS_DIR}/head-single.md" \
+            --baseline-multi "${PERF_RESULTS_DIR}/base-multi.md" \
+            --candidate-multi "${PERF_RESULTS_DIR}/head-multi.md" \
+            --threshold 10 \
+            --output "${PERF_RESULTS_DIR}/model-zoo-performance-comparison.md"
+
+      - name: Upload model zoo reports
+        if: always()
+        uses: actions/upload-artifact@v4
+        with:
+          name: model-zoo-reports-${{ github.run_id }}-${{ github.run_attempt }}
+          path: ${{ env.PERF_RESULTS_DIR }}/
+          if-no-files-found: ignore
+          retention-days: 7
+
+      - name: Cleanup model zoo work directory
+        if: always()
+        run: |
+          set -euo pipefail
+          case "${MODEL_ZOO_WORK}" in
+            /data/ci-work/onnx-model-zoo-*) rm -rf "${MODEL_ZOO_WORK}" ;;
+            *) echo "Refusing to remove unexpected path: ${MODEL_ZOO_WORK}" >&2; exit 1 ;;
+          esac
diff --git a/musa/ep/src/fusion/concat_matmul_fusion.cc b/musa/ep/src/fusion/concat_matmul_fusion.cc
index 2041b39..e93165d 100644
--- a/musa/ep/src/fusion/concat_matmul_fusion.cc
+++ b/musa/ep/src/fusion/concat_matmul_fusion.cc
@@ -481,10 +481,54 @@ void RunMudnnConcat(const std::vector<Ort::ConstValue>& concat_inputs,
                     void* concat_data, ONNXTensorElementDataType elem_type,
                     musaStream_t stream) {
   ::musa::dnn::Handle* handle = MudnnHandleOrThrow(stream);
-  std::vector<::musa::dnn::Tensor> input_tensors(concat_inputs.size());
-  for (size_t i = 0; i < concat_inputs.size(); ++i) {
-    SetupMudnnTensor(input_tensors[i], concat_inputs[i].GetTensorRawData(),
-                     TensorShape(concat_inputs[i]), elem_type);
+  const size_t axis_index = static_cast<size_t>(axis);
+
+  // muDNN Concat rejects tensors with a zero extent, while ONNX permits empty
+  // inputs along the concatenation axis.  Drop those inputs before invoking
+  // muDNN; they contribute no elements and therefore do not change the
+  // concatenation result.
+  std::vector<Ort::ConstValue> non_empty_inputs;
+  non_empty_inputs.reserve(concat_inputs.size());
+  for (Ort::ConstValue input : concat_inputs) {
+    const std::vector<int64_t> shape = TensorShape(input);
+    if (axis_index >= shape.size()) {
+      throw std::runtime_error("ConcatMatMul concat axis is out of range");
+    }
+    if (shape[axis_index] < 0) {
+      throw std::runtime_error(
+          "ConcatMatMul concat input has an unresolved dimension");
+    }
+    if (shape[axis_index] != 0) {
+      non_empty_inputs.push_back(input);
+    }
+  }
+
+  if (non_empty_inputs.empty()) {
+    throw std::runtime_error(
+        "ConcatMatMul concat has no non-empty inputs for a non-empty output");
+  }
+
+  // With only one non-empty input, avoid relying on whether muDNN accepts a
+  // one-input Concat.  Empty inputs have already been removed, so the output
+  // is layout-identical to that input and a device-to-device copy is enough.
+  if (non_empty_inputs.size() == 1) {
+    const size_t bytes =
+        static_cast<size_t>(NumElements(concat_shape)) * ElementSize(elem_type);
+    const musaError_t status =
+        musaMemcpyAsync(concat_data, non_empty_inputs[0].GetTensorRawData(),
+                        bytes, musaMemcpyDeviceToDevice, stream);
+    if (status != musaSuccess) {
+      throw std::runtime_error(
+          std::string("ConcatMatMul concat copy failed: ") +
+          MusaErrorString(status));
+    }
+    return;
+  }
+
+  std::vector<::musa::dnn::Tensor> input_tensors(non_empty_inputs.size());
+  for (size_t i = 0; i < non_empty_inputs.size(); ++i) {
+    SetupMudnnTensor(input_tensors[i], non_empty_inputs[i].GetTensorRawData(),
+                     TensorShape(non_empty_inputs[i]), elem_type);
   }
 
   ::musa::dnn::Tensor output_tensor;
diff --git a/musa/ep/src/fusion/rms_norm_fusion.cc b/musa/ep/src/fusion/rms_norm_fusion.cc
index 6818767..48f8cb8 100644
--- a/musa/ep/src/fusion/rms_norm_fusion.cc
+++ b/musa/ep/src/fusion/rms_norm_fusion.cc
@@ -91,16 +91,6 @@ size_t GetMappedIndex(const std::unordered_map<std::string, size_t>& indices,
   return it->second;
 }
 
-bool IsFloatGpuTensor(Ort::ConstValue value) {
-  auto info = value.GetTensorTypeAndShapeInfo();
-  ONNXTensorElementDataType elem_type = info.GetElementType();
-  return (elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT ||
-          elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT16 ||
-          elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_BFLOAT16 ||
-          elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_DOUBLE) &&
-         IsGpuMemory(value.GetTensorMemoryInfo());
-}
-
 bool IsRmsNormStorageType(ONNXTensorElementDataType elem_type) {
   return elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT ||
          elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_FLOAT16 ||
@@ -108,6 +98,11 @@ bool IsRmsNormStorageType(ONNXTensorElementDataType elem_type) {
          elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_DOUBLE;
 }
 
+bool IsRmsNormFloatTensor(Ort::ConstValue value) {
+  auto info = value.GetTensorTypeAndShapeInfo();
+  return IsRmsNormStorageType(info.GetElementType());
+}
+
 class DeviceTempBuffer {
  public:
   DeviceTempBuffer(size_t bytes, musaStream_t stream)
@@ -142,15 +137,14 @@ std::vector<int64_t> RmsNormInvVarShape(
   return inv_var_shape;
 }
 
-bool TryMudnnRmsNorm(Ort::ConstValue input, const void* gamma_data,
+bool TryMudnnRmsNorm(const void* input_data, const void* gamma_data,
                      Ort::UnownedValue output,
                      const std::vector<int64_t>& input_shape,
                      const std::vector<int64_t>& gamma_shape,
                      ONNXTensorElementDataType elem_type, float epsilon,
                      musaStream_t stream) {
-  if (!IsGpuMemory(input.GetTensorMemoryInfo()) ||
-      !IsGpuMemory(output.GetTensorMemoryInfo()) || gamma_data == nullptr ||
-      input_shape.empty() ||
+  if (input_data == nullptr || !IsGpuMemory(output.GetTensorMemoryInfo()) ||
+      gamma_data == nullptr || input_shape.empty() ||
       elem_type == ONNX_TENSOR_ELEMENT_DATA_TYPE_DOUBLE) {
     return false;
   }
@@ -176,8 +170,7 @@ bool TryMudnnRmsNorm(Ort::ConstValue input, const void* gamma_data,
   ::musa::dnn::Tensor output_tensor;
   ::musa::dnn::Tensor gamma_tensor;
   ::musa::dnn::Tensor inv_var_tensor;
-  if (!SetMudnnTensor(input_tensor, input.GetTensorRawData(), input_shape,
-                      elem_type) ||
+  if (!SetMudnnTensor(input_tensor, input_data, input_shape, elem_type) ||
       !SetMudnnTensor(output_tensor, output.GetTensorMutableRawData(),
                       input_shape, elem_type) ||
       !SetMudnnTensor(gamma_tensor, gamma_data, gamma_shape, elem_type) ||
@@ -214,9 +207,9 @@ struct RmsNormFusionCompute : FusionNodeCompute {
       Ort::KernelContext ctx(kernel_context);
       Ort::ConstValue input = ctx.GetInput(input_index);
       Ort::ConstValue gamma = ctx.GetInput(gamma_index);
-      if (!IsFloatGpuTensor(input)) {
+      if (!IsRmsNormFloatTensor(input)) {
         return Ort::GetApi().CreateStatus(
-            ORT_NOT_IMPLEMENTED, "RmsNorm requires a MUSA floating input");
+            ORT_NOT_IMPLEMENTED, "RmsNorm requires a floating-point input");
       }
 
       auto input_info = input.GetTensorTypeAndShapeInfo();
@@ -261,15 +254,17 @@ struct RmsNormFusionCompute : FusionNodeCompute {
       }
 
       const musaStream_t stream = GetComputeStream(ctx);
+      DeviceInputBuffer input_buffer;
       DeviceInputBuffer gamma_buffer;
+      RETURN_IF_ERROR(input_buffer.Bind(input, stream));
       RETURN_IF_ERROR(gamma_buffer.Bind(gamma, stream));
-      if (TryMudnnRmsNorm(input, gamma_buffer.data(), output, input_shape,
-                          gamma_shape, input_info.GetElementType(), epsilon,
-                          stream)) {
+      if (TryMudnnRmsNorm(input_buffer.data(), gamma_buffer.data(), output,
+                          input_shape, gamma_shape, input_info.GetElementType(),
+                          epsilon, stream)) {
         return nullptr;
       }
       return LaunchStatus(
-          LaunchMusaRmsNormKernel(input.GetTensorRawData(), gamma_buffer.data(),
+          LaunchMusaRmsNormKernel(input_buffer.data(), gamma_buffer.data(),
                                   output.GetTensorMutableRawData(), rows,
                                   norm_size, epsilon, elem_type, stream));
     } catch (const Ort::Exception& ex) {
@@ -304,11 +299,10 @@ struct CastRmsNormFusionCompute : FusionNodeCompute {
       if (input_info.GetElementType() !=
               ONNX_TENSOR_ELEMENT_DATA_TYPE_BFLOAT16 ||
           gamma_info.GetElementType() !=
-              ONNX_TENSOR_ELEMENT_DATA_TYPE_BFLOAT16 ||
-          !IsGpuMemory(input.GetTensorMemoryInfo())) {
+              ONNX_TENSOR_ELEMENT_DATA_TYPE_BFLOAT16) {
         return Ort::GetApi().CreateStatus(
             ORT_NOT_IMPLEMENTED,
-            "CastRmsNorm requires MUSA bfloat16 input and gamma");
+            "CastRmsNorm requires bfloat16 input and gamma");
       }
 
       std::vector<int64_t> input_shape = input_info.GetShape();
@@ -345,15 +339,17 @@ struct CastRmsNormFusionCompute : FusionNodeCompute {
       }
 
       const musaStream_t stream = GetComputeStream(ctx);
+      DeviceInputBuffer input_buffer;
       DeviceInputBuffer gamma_buffer;
+      RETURN_IF_ERROR(input_buffer.Bind(input, stream));
       RETURN_IF_ERROR(gamma_buffer.Bind(gamma, stream));
-      if (TryMudnnRmsNorm(input, gamma_buffer.data(), output, input_shape,
-                          gamma_shape, input_info.GetElementType(), epsilon,
-                          stream)) {
+      if (TryMudnnRmsNorm(input_buffer.data(), gamma_buffer.data(), output,
+                          input_shape, gamma_shape, input_info.GetElementType(),
+                          epsilon, stream)) {
         return nullptr;
       }
       return LaunchStatus(LaunchMusaCastRmsNormBf16Kernel(
-          input.GetTensorRawData(), gamma_buffer.data(),
+          input_buffer.data(), gamma_buffer.data(),
           output.GetTensorMutableRawData(), rows, norm_size, epsilon, stream));
     } catch (const Ort::Exception& ex) {
       Ort::Status status(ex);
diff --git a/musa/ep/src/fusion/split_concat_fusion_matcher.cc b/musa/ep/src/fusion/split_concat_fusion_matcher.cc
index 645dae4..d94cdb8 100644
--- a/musa/ep/src/fusion/split_concat_fusion_matcher.cc
+++ b/musa/ep/src/fusion/split_concat_fusion_matcher.cc
@@ -62,6 +62,18 @@ bool CanFuseSplitConcat(
         !IsFloatTensorValueInfo(reshape_inputs[0])) {
       return false;
     }
+
+    // The runtime implementation consumes the tensor before Reshape and
+    // reconstructs the Reshape dimensions from the statically known output
+    // shape.  Do not select a graph whose Reshape shape is only known at
+    // runtime: CreateSplitConcatFusion cannot safely derive sequence and part
+    // width for that graph and would otherwise fail during Compile.
+    auto reshape_shape = GetTensorShape(reshape_outputs[0]);
+    if (!reshape_shape.has_value() || reshape_shape->size() != 3 ||
+        (*reshape_shape)[1] <= 0 || (*reshape_shape)[2] <= 0 ||
+        (*reshape_shape)[2] % static_cast<int64_t>(split_outputs.size()) != 0) {
+      return false;
+    }
   }
 
   auto split_axis_attr = GetIntAttribute(split_node, "axis");
@@ -79,6 +91,7 @@ bool CanFuseSplitConcat(
   }
 
   std::vector<int64_t> split_sizes;
+  const bool has_explicit_split_sizes = split_inputs.size() == 2;
   if (split_inputs.size() == 2) {
     auto split_initializer = ReadIntInitializerNoLimit(split_inputs[1]);
     if (!split_initializer.has_value() ||
@@ -103,6 +116,12 @@ bool CanFuseSplitConcat(
     }
     split_total += split_size;
   }
+  if (reshape_node && has_explicit_split_sizes) {
+    const auto reshape_shape = GetTensorShape(reshape_node.GetOutputs()[0]);
+    if (!reshape_shape.has_value() || (*reshape_shape)[2] != split_total) {
+      return false;
+    }
+  }
   for (size_t i = 0; i < split_outputs.size(); ++i) {
     Ort::ConstValueInfo split_output = split_outputs[i];
     if (Name(concat_inputs[i]) != Name(split_output) ||
diff --git a/musa/ep/src/kernels/nn/conv.cc b/musa/ep/src/kernels/nn/conv.cc
index 6d2dc15..8b3e2aa 100644
--- a/musa/ep/src/kernels/nn/conv.cc
+++ b/musa/ep/src/kernels/nn/conv.cc
@@ -157,9 +157,11 @@ OrtStatus* Conv::Compute(Ort::KernelContext& ctx) const {
     return Ort::GetApi().CreateStatus(ORT_NOT_IMPLEMENTED,
                                       "Conv only supports 2D NCHW tensors");
   }
-  if (group_ != 1 || w_shape[1] != x_shape[1]) {
-    return Ort::GetApi().CreateStatus(ORT_NOT_IMPLEMENTED,
-                                      "Conv only supports group=1");
+  if (group_ <= 0 || x_shape[1] != w_shape[1] * group_ ||
+      w_shape[0] % group_ != 0) {
+    return Ort::GetApi().CreateStatus(
+        ORT_INVALID_ARGUMENT,
+        "Conv group must divide output channels and match input channels");
   }
 
   std::vector<int64_t> pads = pads_;
@@ -213,6 +215,7 @@ OrtStatus* Conv::Compute(Ort::KernelContext& ctx) const {
   MusaConv2DParams params{};
   params.n = x_shape[0];
   params.c = x_shape[1];
+  params.group = group_;
   params.h = x_shape[2];
   params.w = x_shape[3];
   params.m = w_shape[0];
diff --git a/musa/ep/src/kernels/nn/conv_impl.mu b/musa/ep/src/kernels/nn/conv_impl.mu
index cf96abf..9072890 100644
--- a/musa/ep/src/kernels/nn/conv_impl.mu
+++ b/musa/ep/src/kernels/nn/conv_impl.mu
@@ -24,7 +24,12 @@ __global__ void Conv2DFloatKernel(const float* input,
     const int64_t n = remaining;
 
     float acc = bias == nullptr ? 0.0f : bias[oc];
-    for (int64_t ic = 0; ic < params.c; ++ic) {
+    const int64_t channels_per_group = params.c / params.group;
+    const int64_t outputs_per_group = params.m / params.group;
+    const int64_t group_index = oc / outputs_per_group;
+    for (int64_t ic_in_group = 0; ic_in_group < channels_per_group;
+         ++ic_in_group) {
+      const int64_t ic = group_index * channels_per_group + ic_in_group;
       for (int64_t kh = 0; kh < params.kernel_h; ++kh) {
         const int64_t ih =
             oh * params.stride_h + kh * params.dilation_h - params.pad_h;
@@ -40,7 +45,7 @@ __global__ void Conv2DFloatKernel(const float* input,
           const int64_t input_index =
               ((n * params.c + ic) * params.h + ih) * params.w + iw;
           const int64_t weight_index =
-              ((oc * params.c + ic) * params.kernel_h + kh) *
+              ((oc * channels_per_group + ic_in_group) * params.kernel_h + kh) *
                   params.kernel_w +
               kw;
           acc += input[input_index] * weight[weight_index];
diff --git a/musa/ep/src/kernels/shared_inc/device_kernel_types.h b/musa/ep/src/kernels/shared_inc/device_kernel_types.h
index 22dda5e..764e894 100644
--- a/musa/ep/src/kernels/shared_inc/device_kernel_types.h
+++ b/musa/ep/src/kernels/shared_inc/device_kernel_types.h
@@ -291,6 +291,7 @@ struct MusaTopKParams {
 struct MusaConv2DParams {
   int64_t n;
   int64_t c;
+  int64_t group;
   int64_t h;
   int64_t w;
   int64_t m;
diff --git a/scripts/ci/compare_model_zoo_performance.py b/scripts/ci/compare_model_zoo_performance.py
new file mode 100755
index 0000000..c6edf44
--- /dev/null
+++ b/scripts/ci/compare_model_zoo_performance.py
@@ -0,0 +1,379 @@
+#!/usr/bin/env python3
+"""Compare PR base and head model-zoo performance reports.
+
+Single-stream latency and multi-stream QPS are both compared with the same
+percentage threshold. A candidate fails when it is slower than the base by
+more than the threshold.
+"""
+from __future__ import annotations
+
+import argparse
+import math
+import os
+import re
+import sys
+from dataclasses import dataclass
+from pathlib import Path
+
+
+@dataclass(frozen=True)
+class ReportRow:
+    model: str
+    status: str
+    value: float | None
+
+
+@dataclass(frozen=True)
+class Comparison:
+    model: str
+    base_status: str
+    head_status: str
+    base_value: float | None
+    head_value: float | None
+    change_pct: float | None
+    result: str
+
+    @property
+    def failed(self) -> bool:
+        return self.result in (
+            "REGRESSION",
+            "CANDIDATE_SKIPPED",
+            "STATUS_FAILED",
+            "METRIC_MISSING",
+        )
+
+
+_NUMBER_RE = re.compile(r"[-+]?(?:\d+(?:\.\d*)?|\.\d+)(?:[eE][-+]?\d+)?")
+
+
+def _split_markdown_row(line: str) -> list[str]:
+    return [cell.strip() for cell in line.strip().strip("|").split("|")]
+
+
+def _parse_metric(value: str) -> float | None:
+    lowered = value.lower()
+    if any(word in lowered for word in ("unknown", "skipped", "failed", "n/a")):
+        return None
+    match = _NUMBER_RE.search(value)
+    if match is None:
+        return None
+    parsed = float(match.group())
+    if not math.isfinite(parsed) or parsed <= 0:
+        return None
+    return parsed
+
+
+def parse_report(path: Path, metric_column: str) -> dict[str, ReportRow]:
+    lines = path.read_text(encoding="utf-8").splitlines()
+    header_index = -1
+    model_index = status_index = metric_index = -1
+
+    for index, line in enumerate(lines):
+        if not line.strip().startswith("|"):
+            continue
+        cells = _split_markdown_row(line)
+        if "Model" not in cells or "Status" not in cells or metric_column not in cells:
+            continue
+        header_index = index
+        model_index = cells.index("Model")
+        status_index = cells.index("Status")
+        metric_index = cells.index(metric_column)
+        break
+
+    if header_index < 0:
+        raise ValueError(f"{path}: table with {metric_column!r} column not found")
+
+    rows: dict[str, ReportRow] = {}
+    for line in lines[header_index + 1 :]:
+        if not line.strip().startswith("|"):
+            break
+        cells = _split_markdown_row(line)
+        if not cells or max(model_index, status_index, metric_index) >= len(cells):
+            continue
+        if set(cells[0]) <= set("-: "):
+            continue
+        model = cells[model_index]
+        if not model:
+            continue
+        if model in rows:
+            raise ValueError(f"{path}: duplicate model row: {model}")
+        rows[model] = ReportRow(
+            model=model,
+            status=cells[status_index],
+            value=_parse_metric(cells[metric_index]),
+        )
+
+    if not rows:
+        raise ValueError(f"{path}: no model rows found")
+    return rows
+
+
+def compare_reports(
+    base: dict[str, ReportRow],
+    head: dict[str, ReportRow],
+    mode: str,
+    threshold_pct: float,
+) -> list[Comparison]:
+    missing_from_head = sorted(set(base) - set(head))
+    missing_from_base = sorted(set(head) - set(base))
+    if missing_from_head or missing_from_base:
+        details: list[str] = []
+        if missing_from_head:
+            details.append(f"missing from head: {', '.join(missing_from_head)}")
+        if missing_from_base:
+            details.append(f"missing from base: {', '.join(missing_from_base)}")
+        raise ValueError("model sets differ; " + "; ".join(details))
+
+    comparisons: list[Comparison] = []
+    for model in base:
+        base_row = base[model]
+        head_row = head[model]
+        result = "PASS"
+        change_pct: float | None = None
+
+        if head_row.status not in ("PASS", "SKIPPED"):
+            result = "STATUS_FAILED"
+        elif base_row.status == "SKIPPED" and head_row.status == "SKIPPED":
+            result = "NOT_MEASURED"
+        elif head_row.status == "SKIPPED":
+            result = "CANDIDATE_SKIPPED"
+        elif base_row.status == "SKIPPED":
+            result = "BASE_SKIPPED"
+        elif base_row.status != "PASS":
+            result = "STATUS_FAILED"
+        elif base_row.value is None or head_row.value is None:
+            result = "METRIC_MISSING"
+        else:
+            change_pct = (head_row.value / base_row.value - 1.0) * 100.0
+            slower_pct = change_pct if mode == "single" else -change_pct
+            if slower_pct > threshold_pct + 1e-9:
+                result = "REGRESSION"
+
+        comparisons.append(
+            Comparison(
+                model=model,
+                base_status=base_row.status,
+                head_status=head_row.status,
+                base_value=base_row.value,
+                head_value=head_row.value,
+                change_pct=change_pct,
+                result=result,
+            )
+        )
+    return comparisons
+
+
+def _format_number(value: float | None) -> str:
+    if value is None:
+        return "n/a"
+    return f"{value:.6g}"
+
+
+def _format_change(value: float | None) -> str:
+    if value is None:
+        return "n/a"
+    return f"{value:+.2f}%"
+
+
+def _status_label(row: Comparison, mode: str) -> str:
+    if row.result == "PASS":
+        return "PASS"
+    if row.result == "REGRESSION":
+        metric = "latency" if mode == "single" else "QPS"
+        return f"REGRESSION ({metric} {_format_change(row.change_pct)})"
+    return row.result
+
+
+def _table_for_mode(
+    title: str,
+    comparisons: list[Comparison],
+    mode: str,
+    base_metric: str,
+    head_metric: str,
+    change_metric: str,
+) -> list[str]:
+    lines = [
+        f"## {title}",
+        "",
+        f"| Model | Base {base_metric} | Head {head_metric} | {change_metric} | Result |",
+        "|---|---|---|---|---|",
+    ]
+    for row in comparisons:
+        lines.append(
+            "| "
+            + " | ".join(
+                [
+                    row.model.replace("|", "\\|"),
+                    _format_number(row.base_value),
+                    _format_number(row.head_value),
+                    _format_change(row.change_pct),
+                    _status_label(row, mode),
+                ]
+            )
+            + " |"
+        )
+    lines.append("")
+    return lines
+
+
+def build_markdown(
+    single: list[Comparison],
+    multi: list[Comparison],
+    threshold_pct: float,
+) -> tuple[str, str]:
+    failed = [row for row in single + multi if row.failed]
+    unmeasured = [row for row in single + multi if row.result == "NOT_MEASURED"]
+    model_count = len(single)
+
+    summary_lines = [
+        "## Model Zoo Performance Gate",
+        "",
+        f"- Models: {model_count}",
+        f"- Threshold: more than {threshold_pct:g}% slower than base",
+        "- Single stream metric: average latency (lower is better)",
+        "- Multi stream metric: QPS (higher is better)",
+        f"- Result: {'FAIL' if failed else 'PASS'}",
+    ]
+    if failed:
+        summary_lines.extend(["", "### Regressions", ""])
+        for row in failed:
+            summary_lines.append(
+                f"- {row.model}: {row.result}, change {_format_change(row.change_pct)}"
+            )
+    if unmeasured:
+        summary_lines.extend(
+            ["", f"- Skipped in both versions (not measured): {len(unmeasured)}"]
+        )
+    summary = "\n".join(summary_lines) + "\n"
+
+    lines = [
+        "# Model Zoo Performance Comparison",
+        "",
+        "| Key | Value |",
+        "|---|---|",
+        f"| Threshold | {threshold_pct:g}% |",
+        "| Single stream metric | Average latency, lower is better |",
+        "| Multi stream metric | QPS, higher is better |",
+        f"| Models compared | {model_count} |",
+        f"| Regressions | {len(failed)} |",
+        f"| Result | {'FAIL' if failed else 'PASS'} |",
+        "",
+    ]
+    lines.extend(
+        _table_for_mode(
+            "Single Stream",
+            single,
+            "single",
+            "latency (ms)",
+            "latency (ms)",
+            "Latency change",
+        )
+    )
+    lines.extend(
+        _table_for_mode(
+            "Multi Stream (-s 4 -q 128)",
+            multi,
+            "multi",
+            "QPS",
+            "QPS",
+            "QPS change",
+        )
+    )
+    lines.extend(["## Result", ""])
+    if failed:
+        lines.append(
+            f"FAIL: {len(failed)} model measurement(s) regressed by more than "
+            f"{threshold_pct:g}% or could not be measured."
+        )
+    else:
+        lines.append(
+            f"PASS: no model regressed by more than {threshold_pct:g}% "
+            "in single-stream latency or multi-stream QPS."
+        )
+    lines.append("")
+    return "\n".join(lines) + "\n", summary
+
+
+def main() -> int:
+    parser = argparse.ArgumentParser(description=__doc__)
+    parser.add_argument("--baseline-single", required=True, type=Path)
+    parser.add_argument("--candidate-single", required=True, type=Path)
+    parser.add_argument("--baseline-multi", required=True, type=Path)
+    parser.add_argument("--candidate-multi", required=True, type=Path)
+    parser.add_argument(
+        "--threshold",
+        type=float,
+        default=10.0,
+        help="Maximum allowed slowdown in percent (default: 10).",
+    )
+    parser.add_argument("--output", type=Path, default=None)
+    parser.add_argument(
+        "--summary",
+        type=Path,
+        default=Path(os.environ["GITHUB_STEP_SUMMARY"])
+        if os.environ.get("GITHUB_STEP_SUMMARY")
+        else None,
+        help="Optional markdown summary file, appended if it exists.",
+    )
+    args = parser.parse_args()
+
+    if not math.isfinite(args.threshold) or args.threshold < 0:
+        parser.error("--threshold must be a non-negative number")
+
+    try:
+        single = compare_reports(
+            parse_report(args.baseline_single, "Average Latency"),
+            parse_report(args.candidate_single, "Average Latency"),
+            "single",
+            args.threshold,
+        )
+        multi = compare_reports(
+            parse_report(args.baseline_multi, "QPS"),
+            parse_report(args.candidate_multi, "QPS"),
+            "multi",
+            args.threshold,
+        )
+    except (OSError, ValueError) as exc:
+        print(f"Failed to compare reports: {exc}", file=sys.stderr)
+        return 1
+
+    single_models = {row.model for row in single}
+    multi_models = {row.model for row in multi}
+    if single_models != multi_models:
+        missing_from_multi = sorted(single_models - multi_models)
+        missing_from_single = sorted(multi_models - single_models)
+        print(
+            "Single-stream and multi-stream model sets differ; "
+            f"single-only={missing_from_multi}, multi-only={missing_from_single}",
+            file=sys.stderr,
+        )
+        return 1
+
+    markdown, summary = build_markdown(single, multi, args.threshold)
+    if args.output is not None:
+        args.output.parent.mkdir(parents=True, exist_ok=True)
+        args.output.write_text(markdown, encoding="utf-8")
+        print(f"Saved performance comparison to {args.output}")
+    else:
+        print(markdown, end="")
+
+    if args.summary is not None:
+        with args.summary.open("a", encoding="utf-8") as stream:
+            stream.write(summary)
+
+    failed = [row for row in single + multi if row.failed]
+    if failed:
+        print(f"Performance gate failed: {len(failed)} regression(s)", file=sys.stderr)
+        for row in failed:
+            print(
+                f"  {row.model}: {row.result}, change {_format_change(row.change_pct)}",
+                file=sys.stderr,
+            )
+        return 1
+
+    print("Performance gate passed")
+    return 0
+
+
+if __name__ == "__main__":
+    sys.exit(main())
diff --git a/scripts/ci/run_model_zoo_benchmark.sh b/scripts/ci/run_model_zoo_benchmark.sh
new file mode 100755
index 0000000..b210b23
--- /dev/null
+++ b/scripts/ci/run_model_zoo_benchmark.sh
@@ -0,0 +1,116 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+usage() {
+	cat <<'EOF'
+Usage: run_model_zoo_benchmark.sh LABEL PLUGIN MODE ZOO_DIR RESULTS_DIR
+
+Run one model-zoo performance pass and preserve its report and error logs.
+
+Arguments:
+  LABEL        Version label used in result file names, for example base or head.
+  PLUGIN       Absolute path to libonnxruntime_providers_musa_plugin.so.
+  MODE         single or multi.
+  ZOO_DIR      Prepared onnx-model-zoo working directory.
+  RESULTS_DIR  Directory that receives the report and error logs.
+
+Modes:
+  single  run_all_onnx.sh --ep musa
+  multi   run_all_onnx.sh --ep musa -s 4 -q 128
+EOF
+}
+
+if [[ $# -ne 5 ]]; then
+	usage >&2
+	exit 2
+fi
+
+LABEL="$1"
+PLUGIN="$2"
+MODE="$3"
+ZOO_DIR="$4"
+RESULTS_DIR="$5"
+
+if [[ ! "$LABEL" =~ ^[A-Za-z0-9_.-]+$ ]]; then
+	echo "Invalid label: $LABEL" >&2
+	exit 2
+fi
+if [[ ! -f "$PLUGIN" ]]; then
+	echo "MUSA plugin does not exist: $PLUGIN" >&2
+	exit 1
+fi
+if [[ ! -d "$ZOO_DIR" ]]; then
+	echo "Model zoo directory does not exist: $ZOO_DIR" >&2
+	exit 1
+fi
+
+case "$MODE" in
+	single)
+		REPORT_NAME="musa_ep_performance.md"
+		RUN_ARGS=(--ep musa)
+		;;
+	multi)
+		REPORT_NAME="musa_ep_performance_with_multi_stream.md"
+		RUN_ARGS=(--ep musa -s 4 -q 128)
+		;;
+	*)
+		echo "Mode must be single or multi: $MODE" >&2
+		exit 2
+		;;
+esac
+
+mkdir -p "$RESULTS_DIR"
+export MUSA_EP_LIBRARY="$PLUGIN"
+export MUSA_HOME="${MUSA_HOME:-/usr/local/musa}"
+export MUSA_VISIBLE_DEVICES="${MUSA_VISIBLE_DEVICES:-0}"
+export LD_LIBRARY_PATH="$MUSA_HOME/lib:$MUSA_HOME/lib64:${LD_LIBRARY_PATH:-}"
+
+echo "benchmark_version=${LABEL}"
+echo "benchmark_mode=${MODE}"
+echo "benchmark_plugin=${MUSA_EP_LIBRARY}"
+echo "MUSA_VISIBLE_DEVICES=${MUSA_VISIBLE_DEVICES}"
+sha256sum "$MUSA_EP_LIBRARY"
+
+cd "$ZOO_DIR"
+set +e
+bash run_all_onnx.sh "${RUN_ARGS[@]}"
+run_status=$?
+set -e
+
+report_path="$ZOO_DIR/$REPORT_NAME"
+result_path="$RESULTS_DIR/$LABEL-$MODE.md"
+if [[ -f "$report_path" ]]; then
+	cp "$report_path" "$result_path"
+	echo "Saved benchmark report to $result_path"
+else
+	echo "Benchmark report was not generated: $report_path" >&2
+	exit 1
+fi
+
+if [[ -d "$ZOO_DIR/error_logs" ]]; then
+	error_results="$RESULTS_DIR/$LABEL-$MODE-error-logs"
+	rm -rf "$error_results"
+	cp -a "$ZOO_DIR/error_logs" "$error_results"
+	echo "Saved benchmark error logs to $error_results"
+fi
+
+if [[ "$run_status" -ne 0 ]]; then
+	echo "Benchmark failed with exit code $run_status" >&2
+	exit "$run_status"
+fi
+
+if ! grep -Fq '| Fail | 0 |' "$result_path"; then
+	echo "Benchmark reported one or more failed models: $result_path" >&2
+	grep -F '| Fail |' "$result_path" >&2 || true
+	exit 1
+fi
+
+if [[ "$MODE" == "single" ]]; then
+	grep -Fq '| Mode | classic |' "$result_path"
+else
+	grep -Fq '| Querys | 128 |' "$result_path"
+	grep -Fq '| Streams | 4 |' "$result_path"
+fi
+
+grep -F '| Model count |' "$result_path"
+echo "Benchmark completed: version=${LABEL} mode=${MODE}"
diff --git a/test/fusion/test_concat_matmul_fusion.py b/test/fusion/test_concat_matmul_fusion.py
index affca87..2d4f9f4 100644
--- a/test/fusion/test_concat_matmul_fusion.py
+++ b/test/fusion/test_concat_matmul_fusion.py
@@ -262,6 +262,60 @@ def test_concat_matmul_fusion_zero_k_non_empty_output(tmp_path):
     assert any(str(op).startswith("MUSAExecutionProvider_") for op in musa_ops)
 
 
+def test_concat_matmul_fusion_skips_empty_concat_input(tmp_path):
+    rng = np.random.default_rng(11)
+    x0 = rng.standard_normal((2, 3)).astype(np.float32)
+    x1 = np.empty((2, 0), dtype=np.float32)
+    x2 = rng.standard_normal((2, 2)).astype(np.float32)
+    b = rng.standard_normal((5, 4)).astype(np.float32)
+
+    nodes = [
+        helper.make_node("Concat", ["X0", "X1", "X2"], ["C"], axis=1),
+        helper.make_node("MatMul", ["C", "B"], ["Y"]),
+    ]
+    feeds = {"X0": x0, "X1": x1, "X2": x2, "B": b}
+    model = build_graph_model(
+        nodes,
+        feeds,
+        [("Y", TensorProto.FLOAT)],
+        name="concat_matmul_empty_axis_input_graph",
+    )
+
+    outputs = run_model_and_compare(model, feeds, rtol=1e-3, atol=1e-3)
+    np.testing.assert_allclose(outputs[0], np.concatenate([x0, x2], axis=1) @ b,
+                               rtol=1e-3, atol=1e-3)
+    musa_ops = _profile_musa_ops(
+        model, feeds, tmp_path, "concat_matmul_empty_axis_input"
+    )
+    assert any(str(op).startswith("MUSAExecutionProvider_") for op in musa_ops)
+
+
+def test_concat_matmul_fusion_only_one_non_empty_concat_input(tmp_path):
+    rng = np.random.default_rng(12)
+    x0 = rng.standard_normal((2, 3)).astype(np.float32)
+    x1 = np.empty((2, 0), dtype=np.float32)
+    b = rng.standard_normal((3, 4)).astype(np.float32)
+
+    nodes = [
+        helper.make_node("Concat", ["X0", "X1"], ["C"], axis=1),
+        helper.make_node("MatMul", ["C", "B"], ["Y"]),
+    ]
+    feeds = {"X0": x0, "X1": x1, "B": b}
+    model = build_graph_model(
+        nodes,
+        feeds,
+        [("Y", TensorProto.FLOAT)],
+        name="concat_matmul_single_non_empty_input_graph",
+    )
+
+    outputs = run_model_and_compare(model, feeds, rtol=1e-3, atol=1e-3)
+    np.testing.assert_allclose(outputs[0], x0 @ b, rtol=1e-3, atol=1e-3)
+    musa_ops = _profile_musa_ops(
+        model, feeds, tmp_path, "concat_matmul_single_non_empty_input"
+    )
+    assert any(str(op).startswith("MUSAExecutionProvider_") for op in musa_ops)
+
+
 def test_concat_matmul_fusion_float16_profiles_fused(tmp_path):
     rng = np.random.default_rng(4)
     x0 = rng.standard_normal((2, 3, 4, 5)).astype(np.float16)
diff --git a/test/fusion/test_rms_norm_fusion.py b/test/fusion/test_rms_norm_fusion.py
index 83a0c28..103e5f8 100644
--- a/test/fusion/test_rms_norm_fusion.py
+++ b/test/fusion/test_rms_norm_fusion.py
@@ -277,6 +277,22 @@ def test_rms_norm_fusion(tmp_path):
     _assert_rms_norm_fused(op_names)
 
 
+def test_rms_norm_fusion_accepts_cpu_bound_input(tmp_path):
+    """The fused kernel must copy a normal host feed before launching on MUSA."""
+    model = _build_rms_norm_model()
+    x = np.linspace(-1.0, 1.0, num=2 * 3 * 5, dtype=np.float32).reshape(2, 3, 5)
+
+    (actual,) = run_model_and_compare(model, {"X": x}, rtol=1e-5, atol=1e-5)
+    gamma = np.linspace(0.5, 1.5, num=5, dtype=np.float32)
+    expected = _rms_norm_reference(x, gamma, 1.0e-6)
+    np.testing.assert_allclose(actual, expected, rtol=1e-5, atol=1e-5)
+
+    op_names = _profile_op_names(
+        model, {"X": x}, str(tmp_path / "rms_norm_cpu_bound_input")
+    )
+    _assert_rms_norm_fused(op_names)
+
+
 def test_pow_rms_norm_fusion(tmp_path):
     model = _build_pow_rms_norm_model()
     x = np.linspace(-2.0, 2.0, num=2 * 3 * 5, dtype=np.float32).reshape(2, 3, 5)
diff --git a/test/fusion/test_split_concat_fusion.py b/test/fusion/test_split_concat_fusion.py
index e5cd1fa..b3ca7cc 100644
--- a/test/fusion/test_split_concat_fusion.py
+++ b/test/fusion/test_split_concat_fusion.py
@@ -163,3 +163,46 @@ def test_split_concat_fusion_absorbs_optional_transpose():
     node_names = _profile_musa_node_names(model_bytes, {"X": x})
     for name in ("SplitConcatSplit", "SplitConcatConcat", "SplitConcatTranspose"):
         assert not any(event.startswith(name) for event in node_names), node_names
+
+
+def test_split_concat_fusion_skips_dynamic_upstream_reshape():
+    """A runtime Reshape shape must not make Compile reject the whole graph."""
+    rng = np.random.default_rng(109)
+    x = rng.standard_normal((2, 12)).astype(np.float32)
+    reshape_shape = np.array([2, 3, 4], dtype=np.int64)
+    split_outputs = [f"S{i}" for i in range(4)]
+    graph = helper.make_graph(
+        [
+            # Deliberately leave the intermediate Reshape value without a
+            # ValueInfo. Its dimensions are supplied at runtime, so the
+            # SplitConcat matcher must decline this fusion candidate.
+            helper.make_node("Reshape", ["X", "shape"], ["R"], name="DynamicReshape"),
+            helper.make_node(
+                "Split", ["R"], split_outputs, axis=2, name="DynamicSplit"
+            ),
+            helper.make_node(
+                "Concat", split_outputs, ["Y"], axis=0, name="DynamicConcat"
+            ),
+        ],
+        "split_concat_dynamic_reshape_graph",
+        [
+            helper.make_tensor_value_info("X", TensorProto.FLOAT, [2, 12]),
+            helper.make_tensor_value_info("shape", TensorProto.INT64, [3]),
+        ],
+        [helper.make_tensor_value_info("Y", TensorProto.FLOAT, None)],
+    )
+    model = helper.make_model(graph, opset_imports=[helper.make_opsetid("", 17)])
+    model.ir_version = min(model.ir_version, 10)
+    model_bytes = model.SerializeToString()
+
+    musa_outputs = run_model_and_compare(
+        model_bytes,
+        {"X": x, "shape": reshape_shape},
+        rtol=1e-6,
+        atol=1e-6,
+    )
+    expected = np.concatenate(
+        np.split(x.reshape(2, 3, 4), 4, axis=2),
+        axis=0,
+    )
+    np.testing.assert_allclose(musa_outputs[0], expected, rtol=0, atol=0)
diff --git a/test/ops/test_conv.py b/test/ops/test_conv.py
index c79f1e9..8f60dde 100644
--- a/test/ops/test_conv.py
+++ b/test/ops/test_conv.py
@@ -46,6 +46,37 @@ def test_conv_opset11_float_nchw_with_bias():
     )
 
 
+def test_conv_opset11_float_nchw_grouped_with_bias():
+    x = np.random.default_rng(7).standard_normal((1, 4, 5, 5)).astype(np.float32)
+    w = np.random.default_rng(8).standard_normal((6, 2, 3, 3)).astype(np.float32)
+    b = np.random.default_rng(9).standard_normal((6,)).astype(np.float32)
+    run_and_compare(
+        "Conv",
+        inputs={"X": x, "W": w, "B": b},
+        outputs=[("Y", TensorProto.FLOAT)],
+        attrs={"group": 2, "pads": [1, 1, 1, 1], "strides": [1, 1]},
+        opset=11,
+        rtol=1e-4,
+        atol=1e-4,
+    )
+
+
+def test_conv_opset11_float_nchw_depthwise_fallback_with_bias():
+    # H=1 and kH=1 deliberately selects Conv's custom device fallback path.
+    x = np.random.default_rng(10).standard_normal((1, 4, 1, 7)).astype(np.float32)
+    w = np.random.default_rng(11).standard_normal((4, 1, 1, 3)).astype(np.float32)
+    b = np.random.default_rng(12).standard_normal((4,)).astype(np.float32)
+    run_and_compare(
+        "Conv",
+        inputs={"X": x, "W": w, "B": b},
+        outputs=[("Y", TensorProto.FLOAT)],
+        attrs={"group": 4, "pads": [0, 1, 0, 1], "strides": [1, 1]},
+        opset=11,
+        rtol=1e-4,
+        atol=1e-4,
+    )
+
+
 def test_conv_lowered_conv1d_h1_k1():
     x = np.random.default_rng(0).standard_normal((2, 64, 1, 8)).astype(np.float32)
     w = np.random.default_rng(1).standard_normal((16, 64, 1, 1)).astype(np.float32)