diff --git a/.github/workflows/build-gfx11-rocm.yml b/.github/workflows/build-gfx11-rocm.yml
new file mode 100644
index 000000000000..a309aae70357
--- /dev/null
+++ b/.github/workflows/build-gfx11-rocm.yml
@@ -0,0 +1,578 @@
+name: Build gfx11 + ROCm
+
+on:
+  push:
+    branches: [gfx11]
+  pull_request:
+    types: [opened, synchronize, reopened]
+  workflow_dispatch:
+    inputs:
+      rocm_version:
+        description: 'ROCm version to use (e.g., 7.14.0a20260608) or "latest" to auto-detect'
+        required: false
+        default: 'latest'
+      create_release:
+        description: 'Publish a dated GitHub Release (b<YYYYMMDD>) with the built binaries. Set true only for the nightly dispatch.'
+        required: false
+        default: 'false'
+      roofline:
+        description: 'Build with per-op roofline profiling support (GGML_HIP_ROOFLINE) and bundle the profiler libs (librocprofiler-sdk + libdw). Set false for lean builds.'
+        required: false
+        default: 'true'
+
+# Tracks the latest multiarch ROCm nightly (auto-detected). The earlier pin to
+# 7.14.0a20260608 worked around a libhsa-runtime64 regression (build 1b2a555677)
+# that segfaulted in GpuAgent::InitDma on the gfx115x runners; that issue
+# (https://github.com/ROCm/TheRock/issues/5763) is now resolved. Pass an explicit
+# rocm_version (e.g. 7.14.0a20260608) to pin a specific nightly again.
+#
+# Nightly is driven externally (GitHub cron only fires from the default branch).
+# On a host you control, with `gh` authenticated (token scope: repo + workflow):
+#   0 13 * * *  gh workflow run build-gfx11-rocm.yml --repo ROCm/llama.cpp --ref gfx11 -f rocm_version=latest
+
+env:
+  ROCM_VERSION: ${{ github.event.inputs.rocm_version || 'latest' }}
+  # Default ON so push/PR/cron (non-dispatch) runs keep roofline, matching prior
+  # always-on behavior; a workflow_dispatch can pass 'false' for a lean build.
+  ROOFLINE: ${{ github.event.inputs.roofline || 'true' }}
+
+jobs:
+  build-ubuntu:
+    runs-on: ubuntu-24.04
+    # Single multiarch build: one fat binary covering all current CI arches,
+    # sourced from TheRock's multiarch tarball (arch-neutral host + per-arch
+    # Tensile DBs). gfx1100-1103 (RDNA3 desktop + Hawk Point/Phoenix 760M/780M),
+    # gfx1150/1151/1153 (RDNA3.5 Strix APUs).
+    env:
+      GPU_TARGETS: gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1153
+    outputs:
+      rocm_version: ${{ steps.set-outputs.outputs.rocm_version }}
+      llamacpp_commit_hash: ${{ steps.set-outputs.outputs.llamacpp_commit_hash }}
+
+    steps:
+    - name: Free disk space
+      # Remove unused runner files to free up disk space
+      run: curl -fsSL https://raw.githubusercontent.com/kou/arrow/e49d8ae15583ceff03237571569099a6ad62be32/ci/scripts/util_free_space.sh | bash
+
+    - name: Checkout repository
+      uses: actions/checkout@v4
+
+    - name: Set up Node.js (for llama-ui build from source)
+      uses: actions/setup-node@v6
+      with:
+        node-version: "24"
+
+    - name: Clean up existing ROCm directory (safety precaution)
+      run: |
+        if [ -d "/opt/rocm" ]; then
+          echo "Removing existing /opt/rocm directory..."
+          sudo rm -rf /opt/rocm
+        fi
+        if [ -f "rocm.tar.gz" ]; then
+          rm -f rocm.tar.gz
+        fi
+        echo "Cleanup completed successfully"
+
+    - name: Install build dependencies
+      run: |
+        echo "Installing build dependencies..."
+        sudo apt update
+        sudo apt install -y cmake ninja-build unzip curl
+        echo "Verifying installations..."
+        cmake --version
+        ninja --version
+        echo "Build dependencies installation completed"
+
+    - name: Download and extract multiarch ROCm directly to /opt/rocm
+      run: |
+        rocm_version="${{ env.ROCM_VERSION }}"
+        base_url="https://rocm.nightlies.amd.com/tarball-multi-arch"
+
+        if [ "$rocm_version" = "latest" ]; then
+          echo "Auto-detecting latest multiarch ROCm version"
+          # The multiarch host serves an HTML index (not S3 XML); scrape the
+          # multiarch tarball names from it.
+          files=$(curl -s "$base_url/" \
+            | grep -oE 'therock-dist-linux-multiarch-[0-9]+\.[0-9]+\.[0-9]+(a|rc)[0-9]+\.tar\.gz' \
+            | sort -u)
+
+          latest_file=""
+          latest_major=0
+          latest_minor=0
+          latest_patch=0
+          latest_rc=0
+          latest_is_alpha=false
+
+          while IFS= read -r file; do
+            if [[ "$file" =~ therock-dist-linux-multiarch-([0-9]+\.[0-9]+\.[0-9]+(a|rc)[0-9]+)\.tar\.gz ]]; then
+              version="${BASH_REMATCH[1]}"
+              major=$(echo "$version" | cut -d. -f1)
+              minor=$(echo "$version" | cut -d. -f2)
+              patch=$(echo "$version" | cut -d. -f3 | sed 's/\(a\|rc\).*//')
+              rc=$(echo "$version" | sed 's/.*\(a\|rc\)//')
+              is_alpha=false
+              if [[ "$version" =~ a ]]; then
+                is_alpha=true
+              fi
+
+              is_newer=false
+              if [ "$major" -gt "$latest_major" ]; then
+                is_newer=true
+              elif [ "$major" -eq "$latest_major" ] && [ "$minor" -gt "$latest_minor" ]; then
+                is_newer=true
+              elif [ "$major" -eq "$latest_major" ] && [ "$minor" -eq "$latest_minor" ] && [ "$patch" -gt "$latest_patch" ]; then
+                is_newer=true
+              elif [ "$major" -eq "$latest_major" ] && [ "$minor" -eq "$latest_minor" ] && [ "$patch" -eq "$latest_patch" ]; then
+                if [ "$is_alpha" = true ] && [ "$latest_is_alpha" = false ]; then
+                  is_newer=true
+                elif [ "$is_alpha" = "$latest_is_alpha" ] && [ "$rc" -gt "$latest_rc" ]; then
+                  is_newer=true
+                fi
+              fi
+
+              if [ "$is_newer" = true ]; then
+                latest_file="$file"
+                latest_major="$major"
+                latest_minor="$minor"
+                latest_patch="$patch"
+                latest_rc="$rc"
+                latest_is_alpha="$is_alpha"
+              fi
+            fi
+          done <<< "$files"
+
+          echo "Found latest file: $latest_file"
+
+          if [[ "$latest_file" =~ therock-dist-linux-multiarch-([0-9]+\.[0-9]+\.[0-9]+(a|rc)[0-9]+)\.tar\.gz ]]; then
+            rocm_version="${BASH_REMATCH[1]}"
+            echo "Detected latest ROCm version: $rocm_version"
+          else
+            echo "Failed to extract ROCm version from latest file: $latest_file"
+            echo "Expected pattern: therock-dist-linux-multiarch-<version>.tar.gz"
+            exit 1
+          fi
+        fi
+
+        rocm_url="$base_url/therock-dist-linux-multiarch-${rocm_version}.tar.gz"
+        echo "DETECTED_ROCM_VERSION=$rocm_version" >> $GITHUB_ENV
+
+        # The multiarch tarball (~11.5 GB) ships device code for ALL 26 GPU
+        # arches: every arch's .kpack plus per-arch rocBLAS/hipBLASLt Tensile DBs.
+        # This consumer build only needs the 7 RDNA3/3.5 arches in GPU_TARGETS and
+        # uses the GEMM (Tensile) path, which works without .kpack files. So we
+        # stream-extract and prune at the tar level: drop ALL .kpack, and drop the
+        # Tensile DBs of every arch not in our target set. This keeps the runner
+        # disk footprint small (the 11.5 GB is streamed, never stored) and yields
+        # a lean multiarch package. tar matches --exclude on pre-strip member
+        # names, hence the leading "./".
+        drop_arches="gfx900 gfx906 gfx908 gfx90a gfx942 gfx950 \
+          gfx1010 gfx1011 gfx1012 gfx1030 gfx1031 gfx1032 gfx1033 gfx1034 gfx1035 gfx1036 \
+          gfx1152 gfx1200 gfx1201"
+        excludes=(--exclude='./.kpack' --exclude='./.kpack/*')
+        for a in $drop_arches; do
+          excludes+=("--exclude=./lib/*/library/${a}")
+          excludes+=("--exclude=./lib/*/library/${a}/*")
+          excludes+=("--exclude=./lib/*/library/*${a}*")
+        done
+
+        echo "Streaming multiarch ROCm from: $rocm_url (pruning .kpack + non-target arches)"
+        sudo mkdir -p /opt/rocm
+        curl -sL "$rocm_url" | sudo tar --use-compress-program=gzip -xf - \
+          -C /opt/rocm --strip-components=1 "${excludes[@]}"
+
+        echo "Retained rocBLAS Tensile arch dirs:"
+        ls /opt/rocm/lib/rocblas/library/ 2>/dev/null || echo "(none)"
+        echo "Retained hipBLASLt Tensile arch dirs:"
+        ls /opt/rocm/lib/hipblaslt/library/ 2>/dev/null || echo "(none)"
+
+    - name: Set ROCm environment variables
+      run: |
+        echo "Setting ROCm environment variables..."
+        echo "HIP_PATH=/opt/rocm" >> $GITHUB_ENV
+        echo "ROCM_PATH=/opt/rocm" >> $GITHUB_ENV
+        echo "HIP_PLATFORM=amd" >> $GITHUB_ENV
+        echo "HIP_CLANG_PATH=/opt/rocm/llvm/bin" >> $GITHUB_ENV
+        echo "HIP_INCLUDE_PATH=/opt/rocm/include" >> $GITHUB_ENV
+        echo "HIP_LIB_PATH=/opt/rocm/lib" >> $GITHUB_ENV
+        echo "HIP_DEVICE_LIB_PATH=/opt/rocm/lib/llvm/amdgcn/bitcode" >> $GITHUB_ENV
+        echo "/opt/rocm/bin:/opt/rocm/llvm/bin:$PATH" >> $GITHUB_PATH
+        echo "LD_LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib64:/opt/rocm/llvm/lib:${LD_LIBRARY_PATH:-}" >> $GITHUB_ENV
+        echo "LIBRARY_PATH=/opt/rocm/lib:/opt/rocm/lib64:${LIBRARY_PATH:-}" >> $GITHUB_ENV
+        echo "CPATH=/opt/rocm/include:${CPATH:-}" >> $GITHUB_ENV
+        echo "PKG_CONFIG_PATH=/opt/rocm/lib/pkgconfig:${PKG_CONFIG_PATH:-}" >> $GITHUB_ENV
+        echo "ROCm environment variables set successfully"
+
+    - name: Record llama.cpp commit hash
+      run: |
+        commit_hash=$(git rev-parse --short=5 HEAD)
+        echo "LLAMACPP_COMMIT_HASH=$commit_hash" >> $GITHUB_ENV
+        echo "llama.cpp commit hash (5 digits): $commit_hash"
+        echo "Current llama.cpp commit:"
+        git log --oneline -1
+
+    - name: Setup ccache
+      # Auto-consumed by the build: llama.cpp's CMake sets ccache as the compile
+      # launcher when GGML_CCACHE is ON (default) and ccache is on PATH, so no
+      # cmake flags are needed. Persists across runs via actions/cache. Nightly
+      # ROCm bumps change the compiler and will miss (expected); incremental
+      # commits on the same ROCm hit. Only write the cache from pushes to gfx11
+      # so PRs read but don't thrash the shared store.
+      uses: ggml-org/ccache-action@v1.2.21
+      with:
+        key: gfx11-rocm-multiarch
+        evict-old-files: 7d
+        save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/gfx11' }}
+
+    - name: Build Llama.cpp + ROCm
+      run: |
+        gpu_targets="${{ env.GPU_TARGETS }}"
+        echo "Building multiarch binary (GPU_TARGETS=$gpu_targets)"
+
+        mkdir build
+        cd build
+
+        cmake .. -G Ninja \
+          -DCMAKE_C_COMPILER=/opt/rocm/llvm/bin/clang \
+          -DCMAKE_CXX_COMPILER=/opt/rocm/llvm/bin/clang++ \
+          -DCMAKE_CXX_FLAGS="-I/opt/rocm/include" \
+          -DCMAKE_CROSSCOMPILING=ON \
+          -DCMAKE_BUILD_TYPE=Release \
+          -DGPU_TARGETS="$gpu_targets" \
+          -DBUILD_SHARED_LIBS=ON \
+          -DLLAMA_BUILD_TESTS=OFF \
+          -DGGML_HIP=ON \
+          -DGGML_OPENMP=OFF \
+          -DGGML_CUDA_FORCE_CUBLAS=OFF \
+          -DGGML_RPC=ON \
+          -DGGML_HIP_ROCWMMA_FATTN=OFF \
+          -DGGML_HIP_ROOFLINE=${{ env.ROOFLINE == 'true' && 'ON' || 'OFF' }} \
+          -DLLAMA_BUILD_BORINGSSL=ON \
+          -DGGML_NATIVE=OFF \
+          -DGGML_STATIC=OFF \
+          -DCMAKE_SYSTEM_NAME=Linux
+
+        cmake --build . -j $(nproc)
+
+    - name: Copy ROCm core libs to build directory
+      run: |
+        build_bin_path="build/bin"
+
+        # Copy the rocblas/library folder and all its contents
+        rocblas_lib_path="/opt/rocm/lib/rocblas/library"
+        if [ -d "$rocblas_lib_path" ]; then
+          echo "Copying rocblas/library folder and all contents..."
+          dest_rocblas_path="$build_bin_path/rocblas/library"
+          mkdir -p "$(dirname "$dest_rocblas_path")"
+          cp -r "$rocblas_lib_path" "$(dirname "$dest_rocblas_path")/"
+          echo "Copied: rocblas/library folder with all contents"
+        else
+          echo "Warning: rocblas/library folder not found at: $rocblas_lib_path"
+        fi
+
+        # Copy the hipblaslt/library folder and all its contents
+        hipblaslt_lib_path="/opt/rocm/lib/hipblaslt/library"
+        if [ -d "$hipblaslt_lib_path" ]; then
+          echo "Copying hipblaslt/library folder and all contents..."
+          dest_hipblaslt_path="$build_bin_path/hipblaslt/library"
+          mkdir -p "$(dirname "$dest_hipblaslt_path")"
+          cp -r "$hipblaslt_lib_path" "$(dirname "$dest_hipblaslt_path")/"
+          echo "Copied: hipblaslt/library folder with all contents"
+        else
+          echo "Warning: hipblaslt/library folder not found at: $hipblaslt_lib_path"
+        fi
+
+        # Copy required ROCm libraries to build directory
+        # If artifacts from ROCm or Llama.cpp change, you may need to update this list.
+        echo "Copying required ROCm libraries to build directory..."
+        cp -v /opt/rocm/lib/libhipblas.so* "$build_bin_path/" 2>/dev/null || echo "libhipblas.so* not found"
+        cp -v /opt/rocm/lib/librocblas.so* "$build_bin_path/" 2>/dev/null || echo "librocblas.so* not found"
+        cp -v /opt/rocm/lib/libamdhip64.so* "$build_bin_path/" 2>/dev/null || echo "libamdhip64.so* not found"
+        cp -v /opt/rocm/lib/librocsolver.so* "$build_bin_path/" 2>/dev/null || echo "librocsolver.so* not found"
+        cp -v /opt/rocm/lib/libroctx64.so* "$build_bin_path/" 2>/dev/null || echo "libroctx64.so* not found"
+        cp -v /opt/rocm/lib/libhipblaslt.so* "$build_bin_path/" 2>/dev/null || echo "libhipblaslt.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_liblzma.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_liblzma.so* not found"
+        cp -v /opt/rocm/lib/librocprofiler-register.so* "$build_bin_path/" 2>/dev/null || echo "librocprofiler-register.so* not found"
+        # Roofline builds dlopen librocprofiler-sdk.so.1 at runtime (ggml-cuda-roofline.cpp);
+        # the release must ship it plus its one bundle-missing dep libdw
+        # (librocm_sysdeps_dw.so.1 — a hard DT_NEEDED of the sdk lib). The $ORIGIN RPATH
+        # step below lets the dlopen resolve both locally with no LD_LIBRARY_PATH.
+        if [ "${{ env.ROOFLINE }}" = "true" ]; then
+          cp -v /opt/rocm/lib/librocprofiler-sdk.so* "$build_bin_path/" 2>/dev/null || echo "librocprofiler-sdk.so* not found"
+          cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_dw.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_dw.so* not found"
+        fi
+        cp -v /opt/rocm/lib/libamd_comgr.so* "$build_bin_path/" 2>/dev/null || echo "libamd_comgr.so* not found"
+        cp -v /opt/rocm/lib/libamd_comgr_loader.so* "$build_bin_path/" 2>/dev/null || echo "libamd_comgr_loader.so* not found"
+        cp -v /opt/rocm/lib/libhsa-runtime64.so* "$build_bin_path/" 2>/dev/null || echo "libhsa-runtime64.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_numa.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_numa.so* not found"
+        cp -v /opt/rocm/lib/librocroller.so* "$build_bin_path/" 2>/dev/null || echo "librocroller.so* not found"
+        cp -v /opt/rocm/lib/librocm_kpack.so* "$build_bin_path/" 2>/dev/null || echo "librocm_kpack.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_z.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_z.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_zstd.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_zstd.so* not found"
+        cp -v /opt/rocm/lib/llvm/lib/libLLVM.so* "$build_bin_path/" 2>/dev/null || echo "libLLVM.so* not found"
+        cp -v /opt/rocm/lib/llvm/lib/libclang-cpp.so* "$build_bin_path/" 2>/dev/null || echo "libclang-cpp.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_elf.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_elf.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_drm.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_drm.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_drm_amdgpu.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_drm_amdgpu.so* not found"
+        cp -v /opt/rocm/lib/rocm_sysdeps/lib/librocm_sysdeps_bz2.so* "$build_bin_path/" 2>/dev/null || echo "librocm_sysdeps_bz2.so* not found"
+
+        # Bundle libatomic (gcc runtime dep of llama-cli). Some self-hosted GPU
+        # runners don't have libatomic1 installed, so ship it with the artifact;
+        # the $ORIGIN RPATH step below makes the binary load this bundled copy.
+        cp -v /usr/lib/x86_64-linux-gnu/libatomic.so.1* "$build_bin_path/" 2>/dev/null \
+          || cp -v "$(gcc -print-file-name=libatomic.so.1)" "$build_bin_path/" 2>/dev/null \
+          || echo "libatomic.so.1 not found"
+
+        echo "Finished copying required ROCm libraries"
+
+    - name: Set RPATH for portable distribution
+      run: |
+        sudo apt-get install -y patchelf
+        cd build/bin
+        # Set RPATH to $ORIGIN so all libraries (including the comgr stub loader) find deps locally
+        for file in *.so* llama-*; do
+          [ -f "$file" ] && [ ! -L "$file" ] && patchelf --set-rpath '$ORIGIN' "$file" 2>/dev/null || true
+        done
+
+    - name: List build artifacts (including ROCm files)
+      run: |
+        cd build/bin
+        echo "Final build artifacts (including ROCm library files):"
+        ls -la
+
+    - name: Upload build artifacts
+      uses: actions/upload-artifact@v4
+      with:
+        name: llama-ubuntu-rocm-multiarch-x64
+        path: build/bin/
+        retention-days: 30
+
+    - name: Set job outputs
+      id: set-outputs
+      run: |
+        rocm_version="${DETECTED_ROCM_VERSION:-${{ env.ROCM_VERSION }}}"
+        echo "rocm_version=$rocm_version" >> $GITHUB_OUTPUT
+        echo "llamacpp_commit_hash=${LLAMACPP_COMMIT_HASH}" >> $GITHUB_OUTPUT
+        echo "Final rocm_version: $rocm_version"
+        echo "Final llamacpp_commit_hash: ${LLAMACPP_COMMIT_HASH}"
+
+  test-gfx:
+    needs: build-ubuntu
+    if: needs.build-ubuntu.result == 'success'
+    # Single hardware test of the multiarch artifact on gfx1151. This is the
+    # end-to-end safety net for the Tensile-only multiarch package: a real
+    # llama-cli inference exercising the rocBLAS/hipBLASLt GEMM path on-device.
+    runs-on: linux-gfx1151-gpu-rocm
+
+    steps:
+    - name: Checkout repository
+      uses: actions/checkout@v4
+
+    - name: Download build artifacts
+      uses: actions/download-artifact@v4
+      with:
+        name: llama-ubuntu-rocm-multiarch-x64
+        path: llama-binaries
+
+    - name: Download test model
+      run: |
+        # Pinned to a fixed GitHub release asset instead of huggingface.co: some
+        # self-hosted runners cannot reach huggingface.co (curl 35, connection
+        # reset), but they can reach GitHub. Source:
+        # https://github.com/jimw567/llamacpp-test-assets/releases/tag/test-assets
+        model_url="https://github.com/jimw567/llamacpp-test-assets/releases/download/test-assets/Qwen3-0.6B-Q4_0.gguf"
+        model_path="Qwen3-0.6B-Q4_0.gguf"
+        echo "Downloading test model from: $model_url"
+        curl -fL --retry 5 --retry-all-errors --retry-delay 5 -o "$model_path" "$model_url"
+        if [ -f "$model_path" ]; then
+          file_size=$(stat -c%s "$model_path")
+          echo "Model downloaded successfully. Size: $file_size bytes"
+        else
+          echo "Failed to download model"
+          exit 1
+        fi
+
+    - name: Set up library path
+      run: |
+        echo "LD_LIBRARY_PATH=$(pwd)/llama-binaries:$LD_LIBRARY_PATH" >> $GITHUB_ENV
+
+    - name: Run llama-cli test
+      run: |
+        llama_cli_path="./llama-binaries/llama-cli"
+        model_path="Qwen3-0.6B-Q4_0.gguf"
+        output_file="llama_output.txt"
+
+        chmod +x "$llama_cli_path"
+
+        if [ ! -f "$llama_cli_path" ]; then
+          echo "llama-cli not found at: $llama_cli_path"
+          echo "Available files in llama-binaries:"
+          find llama-binaries -type f
+          exit 1
+        fi
+
+        # Use a prompt with a single correct answer and greedy decoding
+        # (--temp 0) so the result is deterministic and verifiable.
+        prompt="What is 2 + 2? Reply with only the number."
+        echo "Running llama-cli test for gfx1151 (multiarch artifact)..."
+        echo "Command: $llama_cli_path -m \"$model_path\" -ngl 99 --temp 0 -p \"$prompt\" -st -v"
+
+        # Bound the run: a healthy 0.6B inference finishes in seconds. If the
+        # GPU kernel launch hangs (driver/runner issue) the step would otherwise
+        # never end and GitHub would never publish the logs. timeout forces a
+        # clean failure and lets the captured output flush for diagnosis.
+        set +e
+        timeout 180 "$llama_cli_path" -m "$model_path" -ngl 99 --temp 0 -p "$prompt" -st -v > "$output_file" 2>&1
+        exit_code=$?
+        set -e
+
+        echo "=== LLAMA-CLI OUTPUT ==="
+        if [ -f "$output_file" ] && [ -s "$output_file" ]; then
+          cat "$output_file"
+        else
+          echo "(empty)"
+        fi
+        echo "=== END OUTPUT ==="
+        echo "Process exit code: $exit_code"
+
+        if [ $exit_code -eq 124 ]; then
+          echo "❌ llama-cli timed out after 180s (GPU kernel likely hung)"
+          echo "=== GPU / driver diagnostics ==="
+          rocminfo 2>&1 | grep -iE 'Name|Marketing|gfx|Uuid' | head -40 || echo "rocminfo failed"
+          (dmesg 2>/dev/null | tail -40) || echo "dmesg unavailable"
+          rm -f "$output_file"
+          exit 1
+        fi
+
+        if [ $exit_code -ne 0 ]; then
+          echo "❌ llama-cli exited with error code: $exit_code"
+          echo "Checking for missing library dependencies..."
+          ldd "$llama_cli_path" || echo "ldd command failed"
+          rm -f "$output_file"
+          exit 1
+        fi
+
+        # Functional checks against current llama.cpp output:
+        #  (1) the ROCm GPU was selected,
+        #  (2) the model layers were offloaded to it (incl. the output layer),
+        #  (3) the model computed the correct answer to "What is 2 + 2?".
+        # With greedy decoding (--temp 0) the answer is deterministic, so this
+        # verifies the GPU math path end to end, not just that text was emitted.
+        found_device=false
+        found_offload=false
+        found_answer=false
+
+        if grep -q "using device ROCm0" "$output_file"; then
+          found_device=true
+        fi
+
+        # Qwen3-0.6B has 28 transformer layers + output = 29 GPU assignments per
+        # pass (verbose mode logs them more than once; require at least one pass).
+        layers_on_gpu=$(grep -c "assigned to device ROCm0" "$output_file" || true)
+        if [ "$layers_on_gpu" -ge 29 ] && grep -q "offloading output layer to GPU" "$output_file"; then
+          found_offload=true
+        fi
+
+        # The parsed assistant answer ("content" field, printed with -v) must
+        # contain 4 and no other digit, so "4" / "The answer is 4." pass while
+        # "5", "14", "22" fail. This is the deterministic correctness check.
+        # llama-cli now runs via an in-process server (upstream #24948); its
+        # stream ends with a verbose stop chunk carrying an empty "content":"",
+        # so drop empty matches before taking the last real answer chunk.
+        answer_content=$(grep -oE '"content":"[^"]*"' "$output_file" | grep -v '"content":""' | tail -1)
+        if echo "$answer_content" | grep -qE '4' && ! echo "$answer_content" | grep -qE '[0-35-9]'; then
+          found_answer=true
+        fi
+
+        echo "=== TEST RESULTS ==="
+        echo "ROCm GPU selected ('using device ROCm0'): $(if [ "$found_device" = true ]; then echo 'FOUND'; else echo 'NOT FOUND'; fi)"
+        echo "Layers offloaded to GPU ('assigned to device ROCm0' x$layers_on_gpu + output layer): $(if [ "$found_offload" = true ]; then echo 'FOUND'; else echo 'NOT FOUND'; fi)"
+        echo "Correct answer to '2 + 2' (parsed content: ${answer_content:-<none>}): $(if [ "$found_answer" = true ]; then echo 'FOUND'; else echo 'NOT FOUND'; fi)"
+
+        rm -f "$output_file"
+
+        if [ "$found_device" = true ] && [ "$found_offload" = true ] \
+           && [ "$found_answer" = true ]; then
+          echo "✅ Test PASSED - GPU offload + correct deterministic answer verified"
+        else
+          echo "❌ Test FAILED - Missing expected outputs"
+          exit 1
+        fi
+
+  create-release:
+    needs: [build-ubuntu, test-gfx]
+    runs-on: ubuntu-24.04
+    permissions:
+      contents: write
+    # Publish only on the nightly dispatch (external cron passes
+    # -f create_release=true). Push/PR and manual runs never release.
+    # Require the build to succeed and the gfx1151 hardware test to pass.
+    if: |
+      always() &&
+      needs.build-ubuntu.result == 'success' &&
+      needs.test-gfx.result == 'success' &&
+      github.event_name == 'workflow_dispatch' &&
+      github.event.inputs.create_release == 'true'
+    steps:
+    - name: Checkout repository
+      uses: actions/checkout@v4
+
+    - name: Download all build artifacts
+      uses: actions/download-artifact@v4
+      with:
+        path: ./all-artifacts
+
+    - name: Generate dated release tag
+      id: generate-tag
+      env:
+        GITHUB_TOKEN: ${{ github.token }}
+      run: |
+        # Date tag mirrors the ROCm nightly scheme (e.g. b20260609 for the
+        # 7.14.0a20260609-era build). One release per UTC day.
+        TAG="b$(date -u '+%Y%m%d')"
+        echo "tag=${TAG}" >> $GITHUB_OUTPUT
+
+        if gh release view "$TAG" --repo "$GITHUB_REPOSITORY" >/dev/null 2>&1; then
+          echo "Release $TAG already exists; skipping creation."
+          echo "tag_exists=true" >> $GITHUB_OUTPUT
+        else
+          echo "Release $TAG does not exist; will create."
+          echo "tag_exists=false" >> $GITHUB_OUTPUT
+        fi
+        echo "Release tag: $TAG"
+
+    - name: Create multiarch archive
+      if: steps.generate-tag.outputs.tag_exists == 'false'
+      run: |
+        TAG="${{ steps.generate-tag.outputs.tag }}"
+        root="$PWD"
+        artifact_dir="./all-artifacts/llama-ubuntu-rocm-multiarch-x64"
+        archive="llama-${TAG}-ubuntu-rocm-multiarch-x64"
+        if [ -d "$artifact_dir" ]; then
+          echo "Creating ${archive}.tar.gz"
+          tar -czf "$root/${archive}.tar.gz" -C "$artifact_dir" .
+        else
+          echo "ERROR: artifact dir not found: $artifact_dir"
+          exit 1
+        fi
+        ls -la *.tar.gz
+
+    - name: Create GitHub Release
+      if: steps.generate-tag.outputs.tag_exists == 'false'
+      env:
+        GITHUB_TOKEN: ${{ github.token }}
+      run: |
+        TAG="${{ steps.generate-tag.outputs.tag }}"
+        ROCM_VERSION="${{ needs.build-ubuntu.outputs.rocm_version }}"
+        LLAMACPP_COMMIT_HASH="${{ needs.build-ubuntu.outputs.llamacpp_commit_hash }}"
+        gh release create "$TAG" \
+          --repo "$GITHUB_REPOSITORY" \
+          --title "$TAG" \
+          --notes "**Build**: $TAG
+        **OS**: ubuntu
+        **GPU Target(s)**: gfx1100, gfx1101, gfx1102, gfx1103, gfx1150, gfx1151, gfx1153 (single multiarch binary)
+        **ROCm Version**: $ROCM_VERSION (multiarch)
+        **Llama.cpp Commit**: $LLAMACPP_COMMIT_HASH
+        **Build Date**: $(date -u '+%Y-%m-%d %H:%M:%S UTC')
+
+        Prebuilt llama.cpp ROCm binaries — one multiarch package covering the RDNA3 gfx110X family incl. Hawk Point/Phoenix (Radeon 760M/780M) and the RDNA3.5 gfx115x APUs (gfx1150/gfx1151/gfx1153). Built from TheRock's multiarch ROCm runtime, pruned to the CI target arches (per-arch Tensile databases bundled)." \
+          *.tar.gz
diff --git a/.github/workflows/code-style.yml b/.github/workflows/code-style.yml
index 50b598b84ddd..a1022a21ebce 100644
--- a/.github/workflows/code-style.yml
+++ b/.github/workflows/code-style.yml
@@ -5,9 +5,11 @@ on:
   push:
     branches:
       - master
+      - gfx11 # ROCm fork-local, DO NOT UPSTREAM
   pull_request:
     branches:
       - master
+      - gfx11 # ROCm fork-local, DO NOT UPSTREAM
 
 concurrency:
   group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }}
@@ -15,7 +17,9 @@ concurrency:
 
 jobs:
   model-naming:
-    runs-on: [self-hosted, fast]
+    # ROCm fork-local: upstream uses [self-hosted, fast] runners that don't
+    # exist in this fork. DO NOT UPSTREAM this runner change.
+    runs-on: ubuntu-24.04
     steps:
       - uses: actions/checkout@v6
       - name: Check model naming conventions
diff --git a/.github/workflows/editorconfig.yml b/.github/workflows/editorconfig.yml
index 59159cd41444..5d204abb0a29 100644
--- a/.github/workflows/editorconfig.yml
+++ b/.github/workflows/editorconfig.yml
@@ -5,9 +5,11 @@ on:
   push:
     branches:
       - master
+      - gfx11 # ROCm fork-local, DO NOT UPSTREAM
   pull_request:
     branches:
       - master
+      - gfx11 # ROCm fork-local, DO NOT UPSTREAM
 
 concurrency:
   group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }}
@@ -15,7 +17,9 @@ concurrency:
 
 jobs:
   editorconfig:
-    runs-on: [self-hosted, fast]
+    # ROCm fork-local: upstream uses [self-hosted, fast] runners that don't
+    # exist in this fork. DO NOT UPSTREAM this runner change.
+    runs-on: ubuntu-24.04
     steps:
       - uses: actions/checkout@v6
       - uses: editorconfig-checker/action-editorconfig-checker@840e866d93b8e032123c23bac69dece044d4d84c # v2.2.0
diff --git a/.github/workflows/hip-quality-check.yml b/.github/workflows/hip-quality-check.yml
index 5d23f01cf80a..d64ac2cc0834 100644
--- a/.github/workflows/hip-quality-check.yml
+++ b/.github/workflows/hip-quality-check.yml
@@ -65,6 +65,8 @@ jobs:
             -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
             -DGPU_TARGETS=gfx942 \
             -DGGML_HIP=ON \
+            -DLLAMA_BUILD_UI=OFF \
+            -DLLAMA_USE_PREBUILT_UI=OFF \
             -DGGML_HIP_EXPORT_METRICS=Off \
             -DCMAKE_HIP_FLAGS="-Werror -Wno-tautological-compare" \
             -DCMAKE_BUILD_TYPE=Release
@@ -78,6 +80,8 @@ jobs:
             -DCMAKE_HIP_COMPILER="$(hipconfig -l)/clang" \
             -DGPU_TARGETS=gfx908 \
             -DGGML_HIP=ON \
+            -DLLAMA_BUILD_UI=OFF \
+            -DLLAMA_USE_PREBUILT_UI=OFF \
             -DGGML_HIP_EXPORT_METRICS=On \
             -DCMAKE_HIP_FLAGS="" \
             -DCMAKE_BUILD_TYPE=Release
diff --git a/.github/workflows/python-lint.yml b/.github/workflows/python-lint.yml
index 0424f372a147..76eab915265d 100644
--- a/.github/workflows/python-lint.yml
+++ b/.github/workflows/python-lint.yml
@@ -4,6 +4,7 @@ on:
   push:
     branches:
       - master
+      - gfx11 # ROCm fork-local, DO NOT UPSTREAM
     paths: [
       '.github/workflows/python-lint.yml',
       '**/*.py'
@@ -21,7 +22,9 @@ concurrency:
 
 jobs:
   flake8-lint:
-    runs-on: [self-hosted, fast]
+    # ROCm fork-local: upstream uses [self-hosted, fast] runners that don't
+    # exist in this fork. DO NOT UPSTREAM this runner change.
+    runs-on: ubuntu-24.04
     name: Lint
     steps:
       - name: Check out source repository
diff --git a/README.md b/README.md
index e98f2b7f18b1..46e805cf8838 100644
--- a/README.md
+++ b/README.md
@@ -1,604 +1,4 @@
-# llama.cpp
-
-![llama](https://raw.githubusercontent.com/ggml-org/llama.brand/refs/heads/master/cover/llama-cpp/cover-llama-cpp-dark.svg)
-
-[![License: MIT](https://img.shields.io/badge/license-MIT-blue.svg)](https://opensource.org/licenses/MIT)
-[![Release](https://img.shields.io/github/v/release/ggml-org/llama.cpp)](https://github.com/ggml-org/llama.cpp/releases)
-[![Server](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/server.yml)
-[![Docker](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/docker.yml)
-[![Winget](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml/badge.svg)](https://github.com/ggml-org/llama.cpp/actions/workflows/winget.yml)
-
-[Manifesto](https://github.com/ggml-org/llama.cpp/discussions/205) / [ggml](https://github.com/ggml-org/ggml) / [ops](https://github.com/ggml-org/llama.cpp/blob/master/docs/ops.md)
-
-LLM inference in C/C++
-
-## Recent API changes
-
-- [Changelog for `libllama` API](https://github.com/ggml-org/llama.cpp/issues/9289)
-- [Changelog for `llama-server` REST API](https://github.com/ggml-org/llama.cpp/issues/9291)
-
-## Hot topics
-
-- **Hugging Face cache migration: models downloaded with `-hf` are now stored in the standard Hugging Face cache directory, enabling sharing with other HF tools.**
-- **[guide : using the new WebUI of llama.cpp](https://github.com/ggml-org/llama.cpp/discussions/16938)**
-- [guide : running gpt-oss with llama.cpp](https://github.com/ggml-org/llama.cpp/discussions/15396)
-- [[FEEDBACK] Better packaging for llama.cpp to support downstream consumers 🤗](https://github.com/ggml-org/llama.cpp/discussions/15313)
-- Support for the `gpt-oss` model with native MXFP4 format has been added | [PR](https://github.com/ggml-org/llama.cpp/pull/15091) | [Collaboration with NVIDIA](https://blogs.nvidia.com/blog/rtx-ai-garage-openai-oss) | [Comment](https://github.com/ggml-org/llama.cpp/discussions/15095)
-- Multimodal support arrived in `llama-server`: [#12898](https://github.com/ggml-org/llama.cpp/pull/12898) | [documentation](./docs/multimodal.md)
-- VS Code extension for FIM completions: https://github.com/ggml-org/llama.vscode
-- Vim/Neovim plugin for FIM completions: https://github.com/ggml-org/llama.vim
-- Hugging Face Inference Endpoints now support GGUF out of the box! https://github.com/ggml-org/llama.cpp/discussions/9669
-- Hugging Face GGUF editor: [discussion](https://github.com/ggml-org/llama.cpp/discussions/9268) | [tool](https://huggingface.co/spaces/CISCai/gguf-editor)
-- WebGPU support is now available in the browser, see a blog/demo introducing it [here](https://reeselevine.github.io/llamas-on-the-web/).
-
-----
-
-## Quick start
-
-Getting started with llama.cpp is straightforward. Here are several ways to install it on your machine:
-
-- Install `llama.cpp` using [brew, nix, winget, or conda-forge](docs/install.md)
-- Run with Docker - see our [Docker documentation](docs/docker.md)
-- Download pre-built binaries from the [releases page](https://github.com/ggml-org/llama.cpp/releases)
-- Build from source by cloning this repository - check out [our build guide](docs/build.md)
-
-Once installed, you'll need a model to work with. Head to the [Obtaining and quantizing models](#obtaining-and-quantizing-models) section to learn more.
-
-Example command:
-
-```sh
-# Use a local model file
-llama-cli -m my_model.gguf
-
-# Or download and run a model directly from Hugging Face
-llama-cli -hf ggml-org/gemma-3-1b-it-GGUF
-
-# Launch OpenAI-compatible API server
-llama-server -hf ggml-org/gemma-3-1b-it-GGUF
-```
-
-## Description
-
-The main goal of `llama.cpp` is to enable LLM inference with minimal setup and state-of-the-art performance on a wide
-range of hardware - locally and in the cloud.
-
-- Plain C/C++ implementation without any dependencies
-- Apple silicon is a first-class citizen - optimized via ARM NEON, Accelerate and Metal frameworks
-- AVX, AVX2, AVX512 and AMX support for x86 architectures
-- RVV, ZVFH, ZFH, ZICBOP and ZIHINTPAUSE support for RISC-V architectures
-- 1.5-bit, 2-bit, 3-bit, 4-bit, 5-bit, 6-bit, and 8-bit integer quantization for faster inference and reduced memory use
-- Custom CUDA kernels for running LLMs on NVIDIA GPUs (support for AMD GPUs via HIP and Moore Threads GPUs via MUSA)
-- Vulkan and SYCL backend support
-- CPU+GPU hybrid inference to partially accelerate models larger than the total VRAM capacity
-
-The `llama.cpp` project is the main playground for developing new features for the [ggml](https://github.com/ggml-org/ggml) library.
-
-<details>
-<summary>Models</summary>
-
-Typically finetunes of the base models below are supported as well.
-
-Instructions for adding support for new models: [HOWTO-add-model.md](docs/development/HOWTO-add-model.md)
-
-#### Text-only
-
-- [X] LLaMA 🦙
-- [x] LLaMA 2 🦙🦙
-- [x] LLaMA 3 🦙🦙🦙
-- [X] [Mistral 7B](https://huggingface.co/mistralai/Mistral-7B-v0.1)
-- [x] [Mixtral MoE](https://huggingface.co/models?search=mistral-ai/Mixtral)
-- [x] [DBRX](https://huggingface.co/databricks/dbrx-instruct)
-- [x] [Jamba](https://huggingface.co/ai21labs)
-- [X] [Falcon](https://huggingface.co/models?search=tiiuae/falcon)
-- [X] [Chinese LLaMA / Alpaca](https://github.com/ymcui/Chinese-LLaMA-Alpaca) and [Chinese LLaMA-2 / Alpaca-2](https://github.com/ymcui/Chinese-LLaMA-Alpaca-2)
-- [X] [Vigogne (French)](https://github.com/bofenghuang/vigogne)
-- [X] [BERT](https://github.com/ggml-org/llama.cpp/pull/5423)
-- [X] [Koala](https://bair.berkeley.edu/blog/2023/04/03/koala/)
-- [X] [Baichuan 1 & 2](https://huggingface.co/models?search=baichuan-inc/Baichuan) + [derivations](https://huggingface.co/hiyouga/baichuan-7b-sft)
-- [X] [Aquila 1 & 2](https://huggingface.co/models?search=BAAI/Aquila)
-- [X] [Starcoder models](https://github.com/ggml-org/llama.cpp/pull/3187)
-- [X] [Refact](https://huggingface.co/smallcloudai/Refact-1_6B-fim)
-- [X] [MPT](https://github.com/ggml-org/llama.cpp/pull/3417)
-- [X] [Bloom](https://github.com/ggml-org/llama.cpp/pull/3553)
-- [x] [Yi models](https://huggingface.co/models?search=01-ai/Yi)
-- [X] [StableLM models](https://huggingface.co/stabilityai)
-- [x] [Deepseek models](https://huggingface.co/models?search=deepseek-ai/deepseek)
-- [x] [Qwen models](https://huggingface.co/models?search=Qwen/Qwen)
-- [x] [PLaMo-13B](https://github.com/ggml-org/llama.cpp/pull/3557)
-- [x] [Phi models](https://huggingface.co/models?search=microsoft/phi)
-- [x] [PhiMoE](https://github.com/ggml-org/llama.cpp/pull/11003)
-- [x] [GPT-2](https://huggingface.co/gpt2)
-- [x] [Orion 14B](https://github.com/ggml-org/llama.cpp/pull/5118)
-- [x] [InternLM2](https://huggingface.co/models?search=internlm2)
-- [x] [CodeShell](https://github.com/WisdomShell/codeshell)
-- [x] [Gemma](https://ai.google.dev/gemma)
-- [x] [Mamba](https://github.com/state-spaces/mamba)
-- [x] [Grok-1](https://huggingface.co/keyfan/grok-1-hf)
-- [x] [Xverse](https://huggingface.co/models?search=xverse)
-- [x] [Command-R models](https://huggingface.co/models?search=CohereForAI/c4ai-command-r)
-- [x] [SEA-LION](https://huggingface.co/models?search=sea-lion)
-- [x] [GritLM-7B](https://huggingface.co/GritLM/GritLM-7B) + [GritLM-8x7B](https://huggingface.co/GritLM/GritLM-8x7B)
-- [x] [OLMo](https://allenai.org/olmo)
-- [x] [OLMo 2](https://allenai.org/olmo)
-- [x] [OLMoE](https://huggingface.co/allenai/OLMoE-1B-7B-0924)
-- [x] [Granite models](https://huggingface.co/collections/ibm-granite/granite-code-models-6624c5cec322e4c148c8b330)
-- [x] [GPT-NeoX](https://github.com/EleutherAI/gpt-neox) + [Pythia](https://github.com/EleutherAI/pythia)
-- [x] [Snowflake-Arctic MoE](https://huggingface.co/collections/Snowflake/arctic-66290090abe542894a5ac520)
-- [x] [Smaug](https://huggingface.co/models?search=Smaug)
-- [x] [Poro 34B](https://huggingface.co/LumiOpen/Poro-34B)
-- [x] [Bitnet b1.58 models](https://huggingface.co/1bitLLM)
-- [x] [Flan T5](https://huggingface.co/models?search=flan-t5)
-- [x] [Open Elm models](https://huggingface.co/collections/apple/openelm-instruct-models-6619ad295d7ae9f868b759ca)
-- [x] [ChatGLM3-6b](https://huggingface.co/THUDM/chatglm3-6b) + [ChatGLM4-9b](https://huggingface.co/THUDM/glm-4-9b) + [GLMEdge-1.5b](https://huggingface.co/THUDM/glm-edge-1.5b-chat) + [GLMEdge-4b](https://huggingface.co/THUDM/glm-edge-4b-chat)
-- [x] [GLM-4-0414](https://huggingface.co/collections/THUDM/glm-4-0414-67f3cbcb34dd9d252707cb2e)
-- [x] [SmolLM](https://huggingface.co/collections/HuggingFaceTB/smollm-6695016cad7167254ce15966)
-- [x] [EXAONE-3.0-7.8B-Instruct](https://huggingface.co/LGAI-EXAONE/EXAONE-3.0-7.8B-Instruct)
-- [x] [FalconMamba Models](https://huggingface.co/collections/tiiuae/falconmamba-7b-66b9a580324dd1598b0f6d4a)
-- [x] [Jais](https://huggingface.co/inceptionai/jais-13b-chat)
-- [x] [Bielik-11B-v2.3](https://huggingface.co/collections/speakleash/bielik-11b-v23-66ee813238d9b526a072408a)
-- [x] [RWKV-7](https://huggingface.co/collections/shoumenchougou/rwkv7-gxx-gguf)
-- [x] [RWKV-6](https://github.com/BlinkDL/RWKV-LM)
-- [x] [QRWKV-6](https://huggingface.co/recursal/QRWKV6-32B-Instruct-Preview-v0.1)
-- [x] [GigaChat-20B-A3B](https://huggingface.co/ai-sage/GigaChat-20B-A3B-instruct)
-- [X] [Trillion-7B-preview](https://huggingface.co/trillionlabs/Trillion-7B-preview)
-- [x] [Ling models](https://huggingface.co/collections/inclusionAI/ling-67c51c85b34a7ea0aba94c32)
-- [x] [Liquid LFM2 models](https://huggingface.co/collections/LiquidAI/lfm2)
-- [x] [Liquid LFM2.5 models](https://huggingface.co/collections/LiquidAI/lfm25)
-- [x] [Liquid Nanos](https://huggingface.co/collections/LiquidAI/liquid-nanos)
-- [x] [Hunyuan models](https://huggingface.co/collections/tencent/hunyuan-dense-model-6890632cda26b19119c9c5e7)
-- [x] [BailingMoeV2 (Ring/Ling 2.0) models](https://huggingface.co/collections/inclusionAI/ling-v2-68bf1dd2fc34c306c1fa6f86)
-- [x] [Mellum models](https://huggingface.co/JetBrains/models?search=mellum)
-
-#### Multimodal
-
-- [x] [LLaVA 1.5 models](https://huggingface.co/collections/liuhaotian/llava-15-653aac15d994e992e2677a7e), [LLaVA 1.6 models](https://huggingface.co/collections/liuhaotian/llava-16-65b9e40155f60fd046a5ccf2)
-- [x] [BakLLaVA](https://huggingface.co/models?search=SkunkworksAI/Bakllava)
-- [x] [Obsidian](https://huggingface.co/NousResearch/Obsidian-3B-V0.5)
-- [x] [ShareGPT4V](https://huggingface.co/models?search=Lin-Chen/ShareGPT4V)
-- [x] [MobileVLM 1.7B/3B models](https://huggingface.co/models?search=mobileVLM)
-- [x] [Yi-VL](https://huggingface.co/models?search=Yi-VL)
-- [x] [Mini CPM](https://huggingface.co/models?search=MiniCPM)
-- [x] [Moondream](https://huggingface.co/vikhyatk/moondream2)
-- [x] [Bunny](https://github.com/BAAI-DCAI/Bunny)
-- [x] [GLM-EDGE](https://huggingface.co/models?search=glm-edge)
-- [x] [Qwen2-VL](https://huggingface.co/collections/Qwen/qwen2-vl-66cee7455501d7126940800d)
-- [x] [LFM2-VL](https://huggingface.co/collections/LiquidAI/lfm2-vl-68963bbc84a610f7638d5ffa)
-
-</details>
-
-<details>
-<summary>Bindings</summary>
-
-- Python: [ddh0/easy-llama](https://github.com/ddh0/easy-llama)
-- Python: [abetlen/llama-cpp-python](https://github.com/abetlen/llama-cpp-python)
-- Go: [go-skynet/go-llama.cpp](https://github.com/go-skynet/go-llama.cpp)
-- Node.js: [withcatai/node-llama-cpp](https://github.com/withcatai/node-llama-cpp)
-- JS/TS (llama.cpp server client): [lgrammel/modelfusion](https://modelfusion.dev/integration/model-provider/llamacpp)
-- JS/TS (Programmable Prompt Engine CLI): [offline-ai/cli](https://github.com/offline-ai/cli)
-- JavaScript/Wasm (works in browser): [tangledgroup/llama-cpp-wasm](https://github.com/tangledgroup/llama-cpp-wasm)
-- Typescript/Wasm (nicer API, available on npm): [ngxson/wllama](https://github.com/ngxson/wllama)
-- Ruby: [yoshoku/llama_cpp.rb](https://github.com/yoshoku/llama_cpp.rb)
-- Ruby: [docusealco/rllama](https://github.com/docusealco/rllama)
-- Rust (more features): [edgenai/llama_cpp-rs](https://github.com/edgenai/llama_cpp-rs)
-- Rust (nicer API): [mdrokz/rust-llama.cpp](https://github.com/mdrokz/rust-llama.cpp)
-- Rust (more direct bindings): [utilityai/llama-cpp-rs](https://github.com/utilityai/llama-cpp-rs)
-- Rust (automated build from crates.io): [ShelbyJenkins/llm_client](https://github.com/ShelbyJenkins/llm_client)
-- C#/.NET: [SciSharp/LLamaSharp](https://github.com/SciSharp/LLamaSharp)
-- C#/VB.NET (more features - community license): [LM-Kit.NET](https://docs.lm-kit.com/lm-kit-net/index.html)
-- Scala 3: [donderom/llm4s](https://github.com/donderom/llm4s)
-- Clojure: [phronmophobic/llama.clj](https://github.com/phronmophobic/llama.clj)
-- React Native: [mybigday/llama.rn](https://github.com/mybigday/llama.rn)
-- Java: [kherud/java-llama.cpp](https://github.com/kherud/java-llama.cpp)
-- Java: [QuasarByte/llama-cpp-jna](https://github.com/QuasarByte/llama-cpp-jna)
-- Zig: [deins/llama.cpp.zig](https://github.com/Deins/llama.cpp.zig)
-- Flutter/Dart: [netdur/llama_cpp_dart](https://github.com/netdur/llama_cpp_dart)
-- Flutter: [xuegao-tzx/Fllama](https://github.com/xuegao-tzx/Fllama)
-- PHP (API bindings and features built on top of llama.cpp): [distantmagic/resonance](https://github.com/distantmagic/resonance) [(more info)](https://github.com/ggml-org/llama.cpp/pull/6326)
-- Guile Scheme: [guile_llama_cpp](https://savannah.nongnu.org/projects/guile-llama-cpp)
-- Swift [srgtuszy/llama-cpp-swift](https://github.com/srgtuszy/llama-cpp-swift)
-- Swift [ShenghaiWang/SwiftLlama](https://github.com/ShenghaiWang/SwiftLlama)
-- Delphi [Embarcadero/llama-cpp-delphi](https://github.com/Embarcadero/llama-cpp-delphi)
-- Go (no CGo needed): [hybridgroup/yzma](https://github.com/hybridgroup/yzma)
-- Android: [llama.android](/examples/llama.android)
-
-</details>
-
-<details>
-<summary>UIs</summary>
-
-*(to have a project listed here, it should clearly state that it depends on `llama.cpp`)*
-
-- [AI Sublime Text plugin](https://github.com/yaroslavyaroslav/OpenAI-sublime-text) (MIT)
-- [BonzAI App](https://apps.apple.com/us/app/bonzai-your-local-ai-agent/id6752847988) (proprietary)
-- [cztomsik/ava](https://github.com/cztomsik/ava) (MIT)
-- [Dot](https://github.com/alexpinel/Dot) (GPL)
-- [eva](https://github.com/ylsdamxssjxxdd/eva) (MIT)
-- [iohub/collama](https://github.com/iohub/coLLaMA) (Apache-2.0)
-- [janhq/jan](https://github.com/janhq/jan) (AGPL)
-- [johnbean393/Sidekick](https://github.com/johnbean393/Sidekick) (MIT)
-- [KanTV](https://github.com/zhouwg/kantv?tab=readme-ov-file) (Apache-2.0)
-- [KodiBot](https://github.com/firatkiral/kodibot) (GPL)
-- [llama.vim](https://github.com/ggml-org/llama.vim) (MIT)
-- [LARS](https://github.com/abgulati/LARS) (AGPL)
-- [Llama Assistant](https://github.com/vietanhdev/llama-assistant) (GPL)
-- [LlamaLib](https://github.com/undreamai/LlamaLib) (Apache-2.0)
-- [LLMFarm](https://github.com/guinmoon/LLMFarm?tab=readme-ov-file) (MIT)
-- [LLMUnity](https://github.com/undreamai/LLMUnity) (MIT)
-- [LMStudio](https://lmstudio.ai/) (proprietary)
-- [LocalAI](https://github.com/mudler/LocalAI) (MIT)
-- [LostRuins/koboldcpp](https://github.com/LostRuins/koboldcpp) (AGPL)
-- [MindMac](https://mindmac.app) (proprietary)
-- [MindWorkAI/AI-Studio](https://github.com/MindWorkAI/AI-Studio) (FSL-1.1-MIT)
-- [Mobile-Artificial-Intelligence/maid](https://github.com/Mobile-Artificial-Intelligence/maid) (MIT)
-- [Mozilla-Ocho/llamafile](https://github.com/Mozilla-Ocho/llamafile) (Apache-2.0)
-- [nat/openplayground](https://github.com/nat/openplayground) (MIT)
-- [nomic-ai/gpt4all](https://github.com/nomic-ai/gpt4all) (MIT)
-- [ollama/ollama](https://github.com/ollama/ollama) (MIT)
-- [oobabooga/text-generation-webui](https://github.com/oobabooga/text-generation-webui) (AGPL)
-- [PocketPal AI](https://github.com/a-ghorbani/pocketpal-ai) (MIT)
-- [psugihara/FreeChat](https://github.com/psugihara/FreeChat) (MIT)
-- [ptsochantaris/emeltal](https://github.com/ptsochantaris/emeltal) (MIT)
-- [pythops/tenere](https://github.com/pythops/tenere) (AGPL)
-- [ramalama](https://github.com/containers/ramalama) (MIT)
-- [semperai/amica](https://github.com/semperai/amica) (MIT)
-- [withcatai/catai](https://github.com/withcatai/catai) (MIT)
-- [Autopen](https://github.com/blackhole89/autopen) (GPL)
-
-</details>
-
-<details>
-<summary>Tools</summary>
-
-- [akx/ggify](https://github.com/akx/ggify) – download PyTorch models from Hugging Face Hub and convert them to GGML
-- [akx/ollama-dl](https://github.com/akx/ollama-dl) – download models from the Ollama library to be used directly with llama.cpp
-- [crashr/gppm](https://github.com/crashr/gppm) – launch llama.cpp instances utilizing NVIDIA Tesla P40 or P100 GPUs with reduced idle power consumption
-- [gpustack/gguf-parser](https://github.com/gpustack/gguf-parser-go/tree/main/cmd/gguf-parser) - review/check the GGUF file and estimate the memory usage
-- [Styled Lines](https://marketplace.unity.com/packages/tools/generative-ai/styled-lines-llama-cpp-model-292902) (proprietary licensed, async wrapper of inference part for game development in Unity3d with pre-built Mobile and Web platform wrappers and a model example)
-- [unslothai/unsloth](https://github.com/unslothai/unsloth) – 🦥 exports/saves fine-tuned and trained models to GGUF (Apache-2.0)
-
-</details>
-
-<details>
-<summary>Infrastructure</summary>
-
-- [Paddler](https://github.com/intentee/paddler) - Open-source LLMOps platform for hosting and scaling AI in your own infrastructure
-- [GPUStack](https://github.com/gpustack/gpustack) - Manage GPU clusters for running LLMs
-- [llama_cpp_canister](https://github.com/onicai/llama_cpp_canister) - llama.cpp as a smart contract on the Internet Computer, using WebAssembly
-- [llama-swap](https://github.com/mostlygeek/llama-swap) - transparent proxy that adds automatic model switching with llama-server
-- [Kalavai](https://github.com/kalavai-net/kalavai-client) - Crowdsource end to end LLM deployment at any scale
-- [llmaz](https://github.com/InftyAI/llmaz) - ☸️ Easy, advanced inference platform for large language models on Kubernetes.
-- [LLMKube](https://github.com/defilantech/llmkube) - Kubernetes operator for llama.cpp with multi-GPU and Apple Silicon Metal
-  support"
-</details>
-
-<details>
-<summary>Games</summary>
-
-- [Lucy's Labyrinth](https://github.com/MorganRO8/Lucys_Labyrinth) - A simple maze game where agents controlled by an AI model will try to trick you.
-
-</details>
-
-
-## Supported backends
-
-| Backend | Target devices |
-| --- | --- |
-| [Metal](docs/build.md#metal-build) | Apple Silicon |
-| [BLAS](docs/build.md#blas-build) | All |
-| [BLIS](docs/backend/BLIS.md) | All |
-| [SYCL](docs/backend/SYCL.md) | Intel GPU |
-| [OpenVINO [In Progress]](docs/backend/OPENVINO.md) | Intel CPUs, GPUs, and NPUs |
-| [MUSA](docs/build.md#musa) | Moore Threads GPU |
-| [CUDA](docs/build.md#cuda) | Nvidia GPU |
-| [HIP](docs/build.md#hip) | AMD GPU |
-| [ZenDNN](docs/build.md#zendnn) | AMD CPU |
-| [Vulkan](docs/build.md#vulkan) | GPU |
-| [CANN](docs/build.md#cann) | Ascend NPU |
-| [OpenCL](docs/backend/OPENCL.md) | Adreno GPU |
-| [IBM zDNN](docs/backend/zDNN.md) | IBM Z & LinuxONE |
-| [WebGPU](docs/build.md#webgpu) | All |
-| [RPC](https://github.com/ggml-org/llama.cpp/tree/master/tools/rpc) | All |
-| [Hexagon [In Progress]](docs/backend/snapdragon/README.md) | Snapdragon |
-| [VirtGPU](docs/backend/VirtGPU.md) | VirtGPU APIR |
-
-## Obtaining and quantizing models
-
-The [Hugging Face](https://huggingface.co) platform hosts a [number of LLMs](https://huggingface.co/models?library=gguf&sort=trending) compatible with `llama.cpp`:
-
-- [Trending](https://huggingface.co/models?library=gguf&sort=trending)
-- [LLaMA](https://huggingface.co/models?sort=trending&search=llama+gguf)
-
-You can either manually download the GGUF file or directly use any `llama.cpp`-compatible models from [Hugging Face](https://huggingface.co/) or other model hosting sites, by using this CLI argument: `-hf <user>/<model>[:quant]`. For example:
-
-```sh
-llama-cli -hf ggml-org/gemma-3-1b-it-GGUF
-```
-
-By default, the CLI would download from Hugging Face, you can switch to other options with the environment variable `MODEL_ENDPOINT`. The `MODEL_ENDPOINT` must point to a Hugging Face compatible API endpoint.
-
-After downloading a model, use the CLI tools to run it locally - see below.
-
-`llama.cpp` requires the model to be stored in the [GGUF](https://github.com/ggml-org/ggml/blob/master/docs/gguf.md) file format. Models in other data formats can be converted to GGUF using the `convert_*.py` Python scripts in this repo.
-
-The Hugging Face platform provides a variety of online tools for converting, quantizing and hosting models with `llama.cpp`:
-
-- Use the [GGUF-my-repo space](https://huggingface.co/spaces/ggml-org/gguf-my-repo) to convert to GGUF format and quantize model weights to smaller sizes
-- Use the [GGUF-my-LoRA space](https://huggingface.co/spaces/ggml-org/gguf-my-lora) to convert LoRA adapters to GGUF format (more info: https://github.com/ggml-org/llama.cpp/discussions/10123)
-- Use the [GGUF-editor space](https://huggingface.co/spaces/CISCai/gguf-editor) to edit GGUF meta data in the browser (more info: https://github.com/ggml-org/llama.cpp/discussions/9268)
-- Use the [Inference Endpoints](https://ui.endpoints.huggingface.co/) to directly host `llama.cpp` in the cloud (more info: https://github.com/ggml-org/llama.cpp/discussions/9669)
-
-To learn more about model quantization, [read this documentation](tools/quantize/README.md)
-
-## [`llama-cli`](tools/cli)
-
-#### A CLI tool for accessing and experimenting with most of `llama.cpp`'s functionality.
-
-- <details open>
-    <summary>Run in conversation mode</summary>
-
-    Models with a built-in chat template will automatically activate conversation mode. If this doesn't occur, you can manually enable it by adding `-cnv` and specifying a suitable chat template with `--chat-template NAME`
-
-    ```bash
-    llama-cli -m model.gguf
-
-    # > hi, who are you?
-    # Hi there! I'm your helpful assistant! I'm an AI-powered chatbot designed to assist and provide information to users like you. I'm here to help answer your questions, provide guidance, and offer support on a wide range of topics. I'm a friendly and knowledgeable AI, and I'm always happy to help with anything you need. What's on your mind, and how can I assist you today?
-    #
-    # > what is 1+1?
-    # Easy peasy! The answer to 1+1 is... 2!
-    ```
-
-    </details>
-
-- <details>
-    <summary>Run in conversation mode with custom chat template</summary>
-
-    ```bash
-    # use the "chatml" template (use -h to see the list of supported templates)
-    llama-cli -m model.gguf -cnv --chat-template chatml
-
-    # use a custom template
-    llama-cli -m model.gguf -cnv --in-prefix 'User: ' --reverse-prompt 'User:'
-    ```
-
-    </details>
-
-- <details>
-    <summary>Constrain the output with a custom grammar</summary>
-
-    ```bash
-    llama-cli -m model.gguf -n 256 --grammar-file grammars/json.gbnf -p 'Request: schedule a call at 8pm; Command:'
-
-    # {"appointmentTime": "8pm", "appointmentDetails": "schedule a a call"}
-    ```
-
-    The [grammars/](grammars/) folder contains a handful of sample grammars. To write your own, check out the [GBNF Guide](grammars/README.md).
-
-    For authoring more complex JSON grammars, check out https://grammar.intrinsiclabs.ai/
-
-    </details>
-
-
-## [`llama-server`](tools/server)
-
-#### A lightweight, [OpenAI API](https://github.com/openai/openai-openapi) compatible, HTTP server for serving LLMs.
-
-- <details open>
-    <summary>Start a local HTTP server with default configuration on port 8080</summary>
-
-    ```bash
-    llama-server -m model.gguf --port 8080
-
-    # Basic web UI can be accessed via browser: http://localhost:8080
-    # Chat completion endpoint: http://localhost:8080/v1/chat/completions
-    ```
-
-    </details>
-
-- <details>
-    <summary>Support multiple-users and parallel decoding</summary>
-
-    ```bash
-    # up to 4 concurrent requests, each with 4096 max context
-    llama-server -m model.gguf -c 16384 -np 4
-    ```
-
-    </details>
-
-- <details>
-    <summary>Enable speculative decoding</summary>
-
-    ```bash
-    # the draft.gguf model should be a small variant of the target model.gguf
-    llama-server -m model.gguf -md draft.gguf
-    ```
-
-    </details>
-
-- <details>
-    <summary>Serve an embedding model</summary>
-
-    ```bash
-    # use the /embedding endpoint
-    llama-server -m model.gguf --embedding --pooling cls -ub 8192
-    ```
-
-    </details>
-
-- <details>
-    <summary>Serve a reranking model</summary>
-
-    ```bash
-    # use the /reranking endpoint
-    llama-server -m model.gguf --reranking
-    ```
-
-    </details>
-
-- <details>
-    <summary>Constrain all outputs with a grammar</summary>
-
-    ```bash
-    # custom grammar
-    llama-server -m model.gguf --grammar-file grammar.gbnf
-
-    # JSON
-    llama-server -m model.gguf --grammar-file grammars/json.gbnf
-    ```
-
-    </details>
-
-
-## [`llama-perplexity`](tools/perplexity)
-
-#### A tool for measuring the [perplexity](tools/perplexity/README.md) [^1] (and other quality metrics) of a model over a given text.
-
-- <details open>
-    <summary>Measure the perplexity over a text file</summary>
-
-    ```bash
-    llama-perplexity -m model.gguf -f file.txt
-
-    # [1]15.2701,[2]5.4007,[3]5.3073,[4]6.2965,[5]5.8940,[6]5.6096,[7]5.7942,[8]4.9297, ...
-    # Final estimate: PPL = 5.4007 +/- 0.67339
-    ```
-
-    </details>
-
-- <details>
-    <summary>Measure KL divergence</summary>
-
-    ```bash
-    # TODO
-    ```
-
-    </details>
-
-[^1]: [https://huggingface.co/docs/transformers/perplexity](https://huggingface.co/docs/transformers/perplexity)
-
-## [`llama-bench`](tools/llama-bench)
-
-#### Benchmark the performance of the inference for various parameters.
-
-- <details open>
-    <summary>Run default benchmark</summary>
-
-    ```bash
-    llama-bench -m model.gguf
-
-    # Output:
-    # | model               |       size |     params | backend    | threads |          test |                  t/s |
-    # | ------------------- | ---------: | ---------: | ---------- | ------: | ------------: | -------------------: |
-    # | qwen2 1.5B Q4_0     | 885.97 MiB |     1.54 B | Metal,BLAS |      16 |         pp512 |      5765.41 ± 20.55 |
-    # | qwen2 1.5B Q4_0     | 885.97 MiB |     1.54 B | Metal,BLAS |      16 |         tg128 |        197.71 ± 0.81 |
-    #
-    # build: 3e0ba0e60 (4229)
-    ```
-
-    </details>
-
-## [`llama-simple`](examples/simple)
-
-#### A minimal example for implementing apps with `llama.cpp`. Useful for developers.
-
-- <details>
-    <summary>Basic text completion</summary>
-
-    ```bash
-    llama-simple -m model.gguf
-
-    # Hello my name is Kaitlyn and I am a 16 year old girl. I am a junior in high school and I am currently taking a class called "The Art of
-    ```
-
-    </details>
-
-
-## Contributing
-
-- Contributors can open PRs
-- Collaborators will be invited based on contributions
-- Maintainers can push to branches in the `llama.cpp` repo and merge PRs into the `master` branch
-- Any help with managing issues, PRs and projects is very appreciated!
-- See [good first issues](https://github.com/ggml-org/llama.cpp/issues?q=is%3Aissue+is%3Aopen+label%3A%22good+first+issue%22) for tasks suitable for first contributions
-- Read the [CONTRIBUTING.md](CONTRIBUTING.md) for more information
-- Make sure to read this: [Inference at the edge](https://github.com/ggml-org/llama.cpp/discussions/205)
-- A bit of backstory for those who are interested: [Changelog podcast](https://changelog.com/podcast/532)
-
-## Other documentation
-
-- [cli](tools/cli/README.md)
-- [completion](tools/completion/README.md)
-- [server](tools/server/README.md)
-- [GBNF grammars](grammars/README.md)
-
-#### Development documentation
-
-- [How to build](docs/build.md)
-- [Running on Docker](docs/docker.md)
-- [Build on Android](docs/android.md)
-- [Multi-GPU usage](docs/multi-gpu.md)
-- [Performance troubleshooting](docs/development/token_generation_performance_tips.md)
-- [GGML tips & tricks](https://github.com/ggml-org/llama.cpp/wiki/GGML-Tips-&-Tricks)
-
-#### Seminal papers and background on the models
-
-If your issue is with model generation quality, then please at least scan the following links and papers to understand the limitations of LLaMA models. This is especially important when choosing an appropriate model size and appreciating both the significant and subtle differences between LLaMA models and ChatGPT:
-- LLaMA:
-    - [Introducing LLaMA: A foundational, 65-billion-parameter large language model](https://ai.facebook.com/blog/large-language-model-llama-meta-ai/)
-    - [LLaMA: Open and Efficient Foundation Language Models](https://arxiv.org/abs/2302.13971)
-- GPT-3
-    - [Language Models are Few-Shot Learners](https://arxiv.org/abs/2005.14165)
-- GPT-3.5 / InstructGPT / ChatGPT:
-    - [Aligning language models to follow instructions](https://openai.com/research/instruction-following)
-    - [Training language models to follow instructions with human feedback](https://arxiv.org/abs/2203.02155)
-
-## XCFramework
-The XCFramework is a precompiled version of the library for iOS, visionOS, tvOS,
-and macOS. It can be used in Swift projects without the need to compile the
-library from source. For example:
-```swift
-// swift-tools-version: 5.10
-// The swift-tools-version declares the minimum version of Swift required to build this package.
-
-import PackageDescription
-
-let package = Package(
-    name: "MyLlamaPackage",
-    targets: [
-        .executableTarget(
-            name: "MyLlamaPackage",
-            dependencies: [
-                "LlamaFramework"
-            ]),
-        .binaryTarget(
-            name: "LlamaFramework",
-            url: "https://github.com/ggml-org/llama.cpp/releases/download/b5046/llama-b5046-xcframework.zip",
-            checksum: "c19be78b5f00d8d29a25da41042cb7afa094cbf6280a225abe614b03b20029ab"
-        )
-    ]
-)
-```
-The above example is using an intermediate build `b5046` of the library. This can be modified
-to use a different version by changing the URL and checksum.
-
-## Completions
-Command-line completion is available for some environments.
-
-#### Bash Completion
-```bash
-$ build/bin/llama-cli --completion-bash > ~/.llama-completion.bash
-$ source ~/.llama-completion.bash
-```
-Optionally this can be added to your `.bashrc` or `.bash_profile` to load it
-automatically. For example:
-```console
-$ echo "source ~/.llama-completion.bash" >> ~/.bashrc
-```
-
-## Dependencies
-
-- [yhirose/cpp-httplib](https://github.com/yhirose/cpp-httplib) - Single-header HTTP server, used by `llama-server` - MIT license
-- [stb-image](https://github.com/nothings/stb) - Single-header image format decoder, used by multimodal subsystem - Public domain
-- [nlohmann/json](https://github.com/nlohmann/json) - Single-header JSON library, used by various tools/examples - MIT License
-- [miniaudio.h](https://github.com/mackron/miniaudio) - Single-header audio format decoder, used by multimodal subsystem - Public domain
-- [subprocess.h](https://github.com/sheredom/subprocess.h) - Single-header process launching solution for C and C++ - Public domain
+# Deprecation warning
+> [!CAUTION]
+> The ROCm/llama.cpp repository is retired, please use the [upstream](https://github.com/ggml-org/llama.cpp) repository
+>
diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt
index 0ec62a3773d6..c4e99c9de654 100644
--- a/ggml/CMakeLists.txt
+++ b/ggml/CMakeLists.txt
@@ -219,6 +219,7 @@ option(GGML_HIP_NO_VMM                      "ggml: do not try to use HIP VMM"
 option(GGML_HIP_ROCWMMA_FATTN               "ggml: enable rocWMMA for FlashAttention"         OFF)
 option(GGML_HIP_MMQ_MFMA                    "ggml: enable MFMA MMA for CDNA in MMQ"           ON)
 option(GGML_HIP_EXPORT_METRICS              "ggml: enable kernel perf metrics output"         OFF)
+option(GGML_HIP_ROOFLINE                    "ggml: enable per-op roofline profiling"          OFF)
 option(GGML_MUSA_GRAPHS                     "ggml: use MUSA graph, experimental, unstable"    OFF)
 option(GGML_MUSA_MUDNN_COPY                 "ggml: enable muDNN for accelerated copy"         OFF)
 option(GGML_VULKAN                          "ggml: use Vulkan"                                OFF)
diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h
index ac133665d978..be8ad751dca6 100644
--- a/ggml/include/ggml.h
+++ b/ggml/include/ggml.h
@@ -2438,11 +2438,27 @@ extern "C" {
            struct ggml_tensor  * d,
            bool                  masked);
 
+    // memory layout of the ssm_conv input (sx), stored in op_params[0]
+    enum ggml_ssm_conv_layout {
+        GGML_SSM_CONV_LAYOUT_TIME_MAJOR     = 0, // sx = [d_conv-1+n_t, d_inner, n_s] (time contiguous)
+        GGML_SSM_CONV_LAYOUT_CHANNELS_MAJOR = 1, // sx = [d_inner, d_conv-1+n_t, n_s] (channels contiguous)
+    };
+
     GGML_API struct ggml_tensor * ggml_ssm_conv(
             struct ggml_context * ctx,
             struct ggml_tensor  * sx,
             struct ggml_tensor  * c);
 
+    // same as ggml_ssm_conv but sx is channels-major: [d_inner, d_conv-1+n_t, n_s]
+    // (channels contiguous). output layout is unchanged: [d_inner, n_t, n_s].
+    GGML_API struct ggml_tensor * ggml_ssm_conv_channels_major(
+            struct ggml_context * ctx,
+            struct ggml_tensor  * sx,
+            struct ggml_tensor  * c);
+
+    // input (sx) memory layout of an SSM_CONV op (see enum ggml_ssm_conv_layout)
+    GGML_API enum ggml_ssm_conv_layout ggml_ssm_conv_get_layout(const struct ggml_tensor * op);
+
     GGML_API struct ggml_tensor * ggml_ssm_scan(
             struct ggml_context * ctx,
             struct ggml_tensor  * s,
diff --git a/ggml/src/ggml-backend-meta.cpp b/ggml/src/ggml-backend-meta.cpp
index 1f29ec86712d..e084932cd94d 100644
--- a/ggml/src/ggml-backend-meta.cpp
+++ b/ggml/src/ggml-backend-meta.cpp
@@ -754,6 +754,15 @@ static struct ggml_backend_meta_split_state ggml_backend_meta_get_split_state(
     };
 
     auto handle_ssm_conv = [&](const std::vector<ggml_backend_meta_split_state> & src_ss) -> ggml_backend_meta_split_state {
+        if (ggml_ssm_conv_get_layout(tensor) == GGML_SSM_CONV_LAYOUT_CHANNELS_MAJOR) {
+            // channels-major: d_inner is sx (src0) axis 0 and c (src1) axis 1, so a d_inner
+            // split lands on different source axes; the output [d_inner, n_t, n_s] splits on axis 0.
+            if (src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_0 && src_ss[1].axis == GGML_BACKEND_SPLIT_AXIS_1) {
+                return {GGML_BACKEND_SPLIT_AXIS_0, {0}, {1}, 1};
+            }
+            return handle_generic(src_ss, /*scalar_only =*/ false);
+        }
+        // time-major: d_inner is axis 1 of both sx and c
         if (src_ss[0].axis == src_ss[1].axis) {
             if (src_ss[0].axis == GGML_BACKEND_SPLIT_AXIS_0) {
                 return {GGML_BACKEND_SPLIT_AXIS_1, {0}, {1}, 1};
diff --git a/ggml/src/ggml-cann/ggml-cann.cpp b/ggml/src/ggml-cann/ggml-cann.cpp
index 5f51ea3bb3c8..ad734b4a4a33 100644
--- a/ggml/src/ggml-cann/ggml-cann.cpp
+++ b/ggml/src/ggml-cann/ggml-cann.cpp
@@ -2684,7 +2684,12 @@ static bool ggml_backend_cann_supports_op(ggml_backend_dev_t dev, const ggml_ten
                 return true;
             }
         case GGML_OP_SSM_CONV:
-            return true;
+            // ggml_cann_ssm_conv() requires F32 src/dst and the time-major layout
+            // (the channels-major layout is only implemented on CPU/CUDA)
+            return op->type == GGML_TYPE_F32 &&
+                   op->src[0]->type == GGML_TYPE_F32 &&
+                   op->src[1]->type == GGML_TYPE_F32 &&
+                   ggml_ssm_conv_get_layout(op) == GGML_SSM_CONV_LAYOUT_TIME_MAJOR;
         case GGML_OP_CUMSUM:
             return op->src[0]->type == GGML_TYPE_F32;
         case GGML_OP_TRI:
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index df0028cf15e3..84b2b1b40850 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -9498,17 +9498,29 @@ static void ggml_compute_forward_ssm_conv_f32(
     const int ith = params->ith;
     const int nth = params->nth;
 
-    const int nc  = src1->ne[0]; // d_conv
-    const int ncs = src0->ne[0]; // d_conv - 1 + n_t
-    const int nr  = src0->ne[1]; // d_inner
-    const int n_t =  dst->ne[1]; // tokens per sequence
-    const int n_s =  dst->ne[2]; // number of sequences in the batch
+    //   time-major     : src0 = [d_conv-1+n_t, d_inner, n_s] (time contiguous)
+    //   channels-major : src0 = [d_inner, d_conv-1+n_t, n_s] (channels contiguous)
+    // output is [d_inner, n_t, n_s] (channels contiguous) in both cases
+    const bool channels_major = ggml_ssm_conv_get_layout(dst) == GGML_SSM_CONV_LAYOUT_CHANNELS_MAJOR;
+
+    const int nc  = src1->ne[0];                                 // d_conv
+    const int ncs = channels_major ? src0->ne[1] : src0->ne[0];  // d_conv - 1 + n_t
+    const int nr  = channels_major ? src0->ne[0] : src0->ne[1];  // d_inner
+    const int n_t =  dst->ne[1];                                 // tokens per sequence
+    const int n_s =  dst->ne[2];                                 // number of sequences in the batch
 
     GGML_ASSERT( dst->ne[0] == nr);
-    GGML_ASSERT(src0->nb[0] == sizeof(float));
+    GGML_ASSERT(src0->nb[0] == sizeof(float));       // input is contiguous (either layout)
     GGML_ASSERT(src1->nb[0] == sizeof(float));
     GGML_ASSERT(src0->nb[1] == src0->ne[0]*sizeof(float));
 
+    // byte strides of the input to move one channel / one time step, per layout
+    const size_t chan_nb = channels_major ? src0->nb[0] : src0->nb[1];
+    const size_t time_nb = channels_major ? src0->nb[1] : src0->nb[0];
+    const int    chan_stride = chan_nb / sizeof(float);
+    const int    time_stride = time_nb / sizeof(float);
+    GGML_UNUSED(ncs);
+
     // rows per thread
     const int dr = (nr + nth - 1)/nth;
 
@@ -9519,13 +9531,11 @@ static void ggml_compute_forward_ssm_conv_f32(
 
     for (int i3 = 0; i3 < n_s; ++i3) {
         for (int i2 = 0; i2 < n_t; ++i2) {
-            // {d_conv - 1 + n_t, d_inner, n_seqs}
-            // sliding window
-            const float * s = (const float *) ((const char *) src0->data + ir0*(src0->nb[1]) + i2*(src0->nb[0]) + i3*(src0->nb[2])); // {d_conv, d_inner, n_s}
+            // sliding window over the time axis (starting at token i2)
+            const float * s = (const float *) ((const char *) src0->data + ir0*chan_nb + i2*time_nb + i3*(src0->nb[2]));
             const float * c = (const float *) ((const char *) src1->data + ir0*(src1->nb[1])); // {d_conv, d_inner}
             float * x = (float *) ((char *) dst->data + ir0*(dst->nb[0]) + i2*(dst->nb[1]) + i3*(dst->nb[2])); // {d_inner, n_t, n_s}
 
-            // TODO: transpose the output for smaller strides for big batches?
             // d_inner
             for (int i1 = 0; i1 < ir; ++i1) {
                 // rowwise dot product
@@ -9534,7 +9544,7 @@ static void ggml_compute_forward_ssm_conv_f32(
 
                 // d_conv
                 for (int i0 = 0; i0 < nc; ++i0) {
-                    sumf += s[i0 + i1*ncs] * c[i0 + i1*nc];
+                    sumf += s[i0*time_stride + i1*chan_stride] * c[i0 + i1*nc];
                 }
                 x[i1] = sumf;
             }
diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
index 290dc4aff259..ac3da98627bc 100644
--- a/ggml/src/ggml-cuda/common.cuh
+++ b/ggml/src/ggml-cuda/common.cuh
@@ -1394,7 +1394,9 @@ struct ggml_backend_cuda_context {
     cudaEvent_t copy_event = nullptr;
 
     cudaStream_t streams[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } };
-    cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES] = {nullptr};
+    // one cuBLAS handle per (device, stream): the handle carries a workspace that must not be shared
+    // by concurrent streams, otherwise overlapped GEMMs corrupt each other's results
+    cublasHandle_t cublas_handles[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS] = { { nullptr } };
 
     int curr_stream_no = 0;
 
@@ -1457,6 +1459,11 @@ struct ggml_backend_cuda_context {
 
     ggml_cuda_stream_context concurrent_stream_context;
 
+    // dedicated buffer for the branches of overlapped concurrent regions (attention QKV, MoE shared
+    // expert), reused across layers so their scratch never aliases tensors read across the region
+    ggml_backend_buffer_t concurrent_scratch      = nullptr;
+    size_t                concurrent_scratch_size = 0;
+
     ~ggml_backend_cuda_context();
 
     cudaStream_t stream(int device, int stream) {
@@ -1472,12 +1479,13 @@ struct ggml_backend_cuda_context {
     ggml_cuda_stream_context & stream_context() { return concurrent_stream_context; }
 
     cublasHandle_t cublas_handle(int device) {
-        if (cublas_handles[device] == nullptr) {
+        cublasHandle_t & handle = cublas_handles[device][curr_stream_no];
+        if (handle == nullptr) {
             ggml_cuda_set_device(device);
-            CUBLAS_CHECK(cublasCreate(&cublas_handles[device]));
-            CUBLAS_CHECK(cublasSetMathMode(cublas_handles[device], CUBLAS_TF32_TENSOR_OP_MATH));
+            CUBLAS_CHECK(cublasCreate(&handle));
+            CUBLAS_CHECK(cublasSetMathMode(handle, CUBLAS_TF32_TENSOR_OP_MATH));
         }
-        return cublas_handles[device];
+        return handle;
     }
 
     cublasHandle_t cublas_handle() {
diff --git a/ggml/src/ggml-cuda/fattn-tile.cuh b/ggml/src/ggml-cuda/fattn-tile.cuh
index 3e07a9f7e04f..3988f4fd77f5 100644
--- a/ggml/src/ggml-cuda/fattn-tile.cuh
+++ b/ggml/src/ggml-cuda/fattn-tile.cuh
@@ -313,8 +313,19 @@ static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_am
     return 0;
 }
 
+static constexpr __host__ __device__ uint32_t ggml_cuda_fattn_tile_get_config_amd_rdna3_5(const int DKQ, const int DV, const int ncols) {
+    // With rocWMMA FlashAttention off, D=256 prefill runs on the tile kernel; on RDNA3.5 a smaller
+    // K tile with higher occupancy is faster than the shared RDNA values. Other cases fall back.
+    GGML_CUDA_FATTN_TILE_CONFIG_CASE(256, 256, 32, 256, 4, 64, 64)
+
+    return ggml_cuda_fattn_tile_get_config_amd_rdna(DKQ, DV, ncols);
+}
+
 static __host__ uint32_t ggml_cuda_fattn_tile_get_config(const int DKQ, const int DV, const int ncols, const int cc) {
     if (GGML_CUDA_CC_IS_AMD(cc)) {
+        if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
+            return ggml_cuda_fattn_tile_get_config_amd_rdna3_5(DKQ, DV, ncols);
+        }
         if (GGML_CUDA_CC_IS_RDNA(cc)) {
             return ggml_cuda_fattn_tile_get_config_amd_rdna(DKQ, DV, ncols);
         }
@@ -328,11 +339,13 @@ static __host__ uint32_t ggml_cuda_fattn_tile_get_config(const int DKQ, const in
 
 static constexpr __device__ uint32_t ggml_cuda_fattn_tile_get_config(const int DKQ, const int DV, const int ncols) {
 #ifdef GGML_USE_HIP
-#ifdef RDNA
+#ifdef RDNA3_5
+    return ggml_cuda_fattn_tile_get_config_amd_rdna3_5(DKQ, DV, ncols);
+#elif defined(RDNA)
     return ggml_cuda_fattn_tile_get_config_amd_rdna(DKQ, DV, ncols);
 #else
     return ggml_cuda_fattn_tile_get_config_amd(DKQ, DV, ncols);
-#endif // RDNA
+#endif // RDNA3_5
 #else
 #ifdef FAST_FP16_AVAILABLE
     return ggml_cuda_fattn_tile_get_config_nvidia_fp16(DKQ, DV, ncols);
diff --git a/ggml/src/ggml-cuda/ggml-cuda-roofline.cpp b/ggml/src/ggml-cuda/ggml-cuda-roofline.cpp
new file mode 100644
index 000000000000..a0f529dc8c25
--- /dev/null
+++ b/ggml/src/ggml-cuda/ggml-cuda-roofline.cpp
@@ -0,0 +1,641 @@
+// Per-op GPU roofline profiling for the HIP/ROCm backend via rocprofiler-sdk.
+//
+// Attributes each GPU kernel's device-measured duration to the ggml op that launched
+// it, using a unique external correlation id per op invocation, and writes a JSON
+// report on exit with one row per invocation listing the kernels it dispatched.
+// Aggregation (grouping identical ops, averaging, counting) is left to the consumer.
+// Activated at runtime by the environment variable GGML_ROOFLINE_OUT=<path.json>;
+// every entry point is a no-op unless it is set.
+//
+// rocprofiler-sdk is loaded with dlopen and configured with rocprofiler_force_configure
+// at runtime, and its entry points are resolved with dlsym. It is intentionally not
+// linked: its global constructors abort when the library is loaded early as part of
+// another shared object.
+
+#include "ggml-cuda-roofline.h"
+#include "ggml-impl.h"
+
+#include <rocprofiler-sdk/rocprofiler.h>
+#include <rocprofiler-sdk/registration.h>
+
+#include <hip/hip_runtime_api.h>   // hipStreamSynchronize / hipMemcpy: read back MoE routing ids
+
+#include <dlfcn.h>
+#include <cxxabi.h>
+
+#include <algorithm>
+#include <atomic>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <fstream>
+#include <mutex>
+#include <sstream>
+#include <string>
+#include <unordered_map>
+#include <unordered_set>
+#include <utility>
+#include <vector>
+
+namespace {
+
+constexpr int n_op_params = GGML_MAX_OP_PARAMS / (int) sizeof(int32_t);  // 16
+
+// Geometry and memory traffic captured for one distinct ggml op shape.
+struct op_record {
+    const char * op    = "";                        // op / unary / glu name (ggml_op_desc)
+    const char * dtype = "";                        // destination type
+    const char * quant = "";                        // src0 type (weight quantization for matmul)
+    int64_t      ne[4]                   = {0, 0, 0, 0};  // destination shape
+    int64_t      src_ne[GGML_MAX_SRC][4] = {};           // source shapes
+    const char * src_types[GGML_MAX_SRC] = {};           // source types (ggml_type_name)
+    int          n_src                   = 0;
+    int32_t      op_params[n_op_params]  = {};            // op parameters (conv / pool / rope / ... geometry)
+    int64_t      bytes                   = 0;             // total HBM traffic: destination + all sources
+    int64_t      dst_bytes               = 0;             // ggml_nbytes(destination)
+    int64_t      src_bytes[GGML_MAX_SRC] = {};            // ggml_nbytes(source)
+    // Two identities per tensor, so the consumer can tell which tensors actually move memory
+    // without seeing the graph. A *tensor id* is one ggml_tensor's address (identifies a single
+    // tensor; used to dedup a tensor read by several ops). A *storage id* is the address of the
+    // tensor that owns the underlying buffer -- ggml's view_src root -- so tensors that alias the
+    // same memory (a view and its source; an in-place op's output and the dst it was handed in as
+    // a source) share one storage id; that is how the consumer spots internal and in-place tensors.
+    // Both are needed -- the storage id alone is not enough. Distinct sources can alias one buffer
+    // (e.g. two ops reading the first and second half of the same tensor: one storage id, two
+    // tensor ids). Both halves are genuinely read, so read dedup must key on the tensor id;
+    // deduping on the storage id would count the buffer once and undercount its traffic.
+    uint64_t     dst_storage_id          = 0;             // destination's storage id (its buffer)
+    uint64_t     src_tensor_ids[GGML_MAX_SRC]  = {};      // each source's tensor id (dedup key)
+    uint64_t     src_storage_ids[GGML_MAX_SRC] = {};      // each source's storage id (its buffer)
+    int64_t      M = 0, N = 0, K = 0, n_experts = 0, top_k = 0;  // matmul dimensions
+    std::vector<op_record> fused_nodes;                  // per-node geometry when this row is a fused group (else empty)
+};
+
+// One kernel dispatch: which kernel ran, when it started, and for how long. The start
+// time is kept only to order kernels within an op (buffer records arrive unordered).
+struct dispatch {
+    uint64_t start_ns    = 0;
+    uint64_t kernel_id   = 0;
+    uint64_t duration_ns = 0;
+};
+
+std::mutex                                           g_mutex;
+std::unordered_map<uint64_t, op_record>              g_records;      // geometry id -> geometry
+std::unordered_map<uint64_t, uint64_t>               g_invocations;  // invocation id -> geometry id
+std::unordered_map<uint64_t, std::vector<dispatch>>  g_dispatches;   // invocation id -> dispatches
+std::unordered_map<uint64_t, std::string>            g_kernel_names; // kernel id -> demangled symbol
+// MoE routing is data-dependent, so the count of distinct experts actually read is captured
+// per invocation (not per shape): one entry per MUL_MAT_ID launch. See count_active_experts.
+std::unordered_map<uint64_t, int64_t>                g_invocation_experts;  // invocation id -> distinct experts routed
+// Per-expert routed-token counts for the same launch (length n_experts, index e = number of
+// (token, slot) routings that selected expert e; Σ = n_tokens*top_k). Emitted as tokens_per_expert
+// so the consumer can show the routing-load distribution (skew that drives MoE-GEMM padding).
+std::unordered_map<uint64_t, std::vector<int32_t>>   g_invocation_expert_hist;  // invocation id -> per-expert token counts
+std::atomic<uint64_t>                                g_next_invocation{1};
+thread_local uint64_t                                g_current_invocation = 0;  // id pushed by the last begin_op on this thread
+
+bool                     g_active = false;
+std::string              g_out_path;
+std::string              g_device;                          // GPU architecture, e.g. "gfx1151"
+rocprofiler_context_id_t g_context{0};
+rocprofiler_buffer_id_t  g_buffer{0};
+
+// rocprofiler-sdk entry points, resolved at runtime via dlsym.
+using fn_create_context = rocprofiler_status_t (*)(rocprofiler_context_id_t *);
+using fn_create_buffer  = rocprofiler_status_t (*)(rocprofiler_context_id_t, size_t, size_t,
+                                                   rocprofiler_buffer_policy_t,
+                                                   rocprofiler_buffer_tracing_cb_t, void *,
+                                                   rocprofiler_buffer_id_t *);
+using fn_config_buffer  = rocprofiler_status_t (*)(rocprofiler_context_id_t,
+                                                   rocprofiler_buffer_tracing_kind_t,
+                                                   const rocprofiler_tracing_operation_t *,
+                                                   size_t, rocprofiler_buffer_id_t);
+using fn_config_callback = rocprofiler_status_t (*)(rocprofiler_context_id_t,
+                                                    rocprofiler_callback_tracing_kind_t,
+                                                    const rocprofiler_tracing_operation_t *,
+                                                    size_t, rocprofiler_callback_tracing_cb_t, void *);
+using fn_start          = rocprofiler_status_t (*)(rocprofiler_context_id_t);
+using fn_flush          = rocprofiler_status_t (*)(rocprofiler_buffer_id_t);
+using fn_get_thread_id  = rocprofiler_status_t (*)(rocprofiler_thread_id_t *);
+using fn_push_id        = rocprofiler_status_t (*)(rocprofiler_context_id_t, rocprofiler_thread_id_t,
+                                                   rocprofiler_user_data_t);
+using fn_pop_id         = rocprofiler_status_t (*)(rocprofiler_context_id_t, rocprofiler_thread_id_t,
+                                                   rocprofiler_user_data_t *);
+
+fn_create_context  p_create_context  = nullptr;
+fn_create_buffer   p_create_buffer   = nullptr;
+fn_config_buffer   p_config_buffer   = nullptr;
+fn_config_callback p_config_callback = nullptr;
+fn_start           p_start           = nullptr;
+fn_flush           p_flush           = nullptr;
+fn_get_thread_id   p_get_thread_id   = nullptr;
+fn_push_id         p_push_id         = nullptr;
+fn_pop_id          p_pop_id          = nullptr;
+
+// Resolve the required entry points; returns false if any is missing. The code-object
+// callback service is optional (kernel-symbol names) and resolved separately.
+bool resolve_symbols() {
+    p_create_context  = (fn_create_context)  dlsym(RTLD_DEFAULT, "rocprofiler_create_context");
+    p_create_buffer   = (fn_create_buffer)   dlsym(RTLD_DEFAULT, "rocprofiler_create_buffer");
+    p_config_buffer   = (fn_config_buffer)   dlsym(RTLD_DEFAULT, "rocprofiler_configure_buffer_tracing_service");
+    p_start           = (fn_start)           dlsym(RTLD_DEFAULT, "rocprofiler_start_context");
+    p_flush           = (fn_flush)           dlsym(RTLD_DEFAULT, "rocprofiler_flush_buffer");
+    p_get_thread_id   = (fn_get_thread_id)   dlsym(RTLD_DEFAULT, "rocprofiler_get_thread_id");
+    p_push_id         = (fn_push_id)         dlsym(RTLD_DEFAULT, "rocprofiler_push_external_correlation_id");
+    p_pop_id          = (fn_pop_id)          dlsym(RTLD_DEFAULT, "rocprofiler_pop_external_correlation_id");
+    p_config_callback = (fn_config_callback) dlsym(RTLD_DEFAULT, "rocprofiler_configure_callback_tracing_service");
+    return p_create_context && p_create_buffer && p_config_buffer && p_start && p_flush &&
+           p_get_thread_id && p_push_id && p_pop_id;
+}
+
+std::string demangle(const char * name) {
+    if (!name) return "";
+    int    status     = 0;
+    char * demangled  = abi::__cxa_demangle(name, nullptr, nullptr, &status);
+    std::string result = (status == 0 && demangled) ? demangled : name;
+    free(demangled);
+    return result;
+}
+
+void json_escape(std::ostringstream & out, const std::string & str) {
+    for (char c : str) {
+        if (c == '"' || c == '\\') out << '\\' << c;
+        else                       out << c;
+    }
+}
+
+uint64_t hash_mix(uint64_t hash, uint64_t value) {
+    hash ^= value + 0x9e3779b97f4a7c15ULL + (hash << 6) + (hash >> 2);
+    return hash;
+}
+
+// Storage identity: ggml points a view/in-place tensor's view_src at the buffer it aliases, so
+// following the chain to the end yields the tensor that actually owns the HBM. Two tensors touch
+// the same memory iff they share this root. VIEW/RESHAPE outputs carry view_src, and so do
+// in-place ops (e.g. CPY does ggml_view_tensor(dst)), which lets the consumer tell a read that
+// actually targets the op's own output (an in-place destination) from a genuine input read.
+const ggml_tensor * roofline_storage(const ggml_tensor * t) {
+    while (t && t->view_src) t = t->view_src;
+    return t;
+}
+
+// Fill a record's geometry and single-node HBM byte fields from one ggml node.
+void fill_head_record(op_record & rec, const ggml_tensor * node) {
+    const ggml_tensor * src0 = node->src[0];
+    const ggml_tensor * src1 = node->src[1];
+
+    rec.op    = ggml_op_desc(node);
+    rec.dtype = ggml_type_name(node->type);
+    rec.quant = src0 ? ggml_type_name(src0->type) : "";
+    for (int d = 0; d < 4; d++) rec.ne[d] = node->ne[d];
+    for (int j = 0; j < GGML_MAX_SRC; j++) {
+        if (node->src[j]) {
+            for (int d = 0; d < 4; d++) rec.src_ne[j][d] = node->src[j]->ne[d];
+            rec.src_types[j] = ggml_type_name(node->src[j]->type);
+            rec.n_src = j + 1;
+        }
+    }
+    memcpy(rec.op_params, node->op_params, sizeof(rec.op_params));
+
+    // total HBM traffic for this op: write the destination and read every source.
+    rec.dst_bytes = (int64_t) ggml_nbytes(node);
+    rec.bytes     = rec.dst_bytes;
+    rec.dst_storage_id   = (uint64_t) (uintptr_t) roofline_storage(node);
+    for (int j = 0; j < GGML_MAX_SRC; j++) {
+        if (node->src[j]) {
+            rec.src_bytes[j] = (int64_t) ggml_nbytes(node->src[j]);
+            rec.bytes       += rec.src_bytes[j];
+            rec.src_tensor_ids[j]   = (uint64_t) (uintptr_t) node->src[j];
+            rec.src_storage_ids[j]  = (uint64_t) (uintptr_t) roofline_storage(node->src[j]);
+        }
+    }
+
+    if ((node->op == GGML_OP_MUL_MAT || node->op == GGML_OP_MUL_MAT_ID) && src0 && src1) {
+        rec.K = src0->ne[0];
+        rec.N = src0->ne[1];
+        if (node->op == GGML_OP_MUL_MAT) {
+            rec.M = src1->ne[1];
+        } else {
+            // MUL_MAT_ID (MoE): src0 = experts [K,N,E], src1 = input (often broadcast
+            // [K,1,n_tokens]), src2 = ids [n_expert_used, n_tokens]. top_k is the number of
+            // experts routed per token = ids->ne[0]; src1->ne[1] is 1 when the input is
+            // broadcast, so it must not be used. Fall back to src1->ne[1] only if ids absent.
+            rec.M         = src1->ne[2];
+            rec.n_experts = src0->ne[2];
+            rec.top_k     = node->src[2] ? node->src[2]->ne[0] : src1->ne[1];
+        }
+    }
+}
+
+// Distinct experts actually routed by a MoE (MUL_MAT_ID) launch. The number of expert weight
+// slabs streamed from HBM is data-dependent -- it is the count of unique ids, not the shape
+// bound min(M*top_k, E) -- and varies per launch, so it is read back for every invocation.
+//
+// The ids tensor is produced by an upstream op on the same stream and may still be in flight,
+// so the stream is synchronized before the copy. This slows the profiling run but does not
+// perturb the report: every roofline figure is derived from per-kernel device timestamps, which
+// a host-side sync between ops leaves untouched. Returns 0 (ids not counted) on any failure or
+// a non-i32 / missing ids tensor; callers then fall back to the shape bound in the consumer.
+// Also fills *counts* with the per-expert routed-token count (length n_experts, index e = number
+// of routings that selected expert e): the same single host-side pass that counts distinct experts
+// yields the full load distribution at no extra cost. *counts* is cleared and left empty on any
+// failure (missing / non-i32 ids); the return value stays the distinct-expert count as before.
+int64_t count_active_experts(const ggml_tensor * ids, int64_t n_experts, hipStream_t stream,
+                             std::vector<int32_t> & counts) {
+    counts.clear();
+    if (!ids || ids->type != GGML_TYPE_I32 || ids->data == nullptr || n_experts <= 0) return 0;
+    if (hipStreamSynchronize(stream) != hipSuccess) return 0;
+
+    std::vector<char> host(ggml_nbytes(ids));
+    if (hipMemcpy(host.data(), ids->data, host.size(), hipMemcpyDeviceToHost) != hipSuccess) return 0;
+
+    counts.assign(n_experts, 0);
+    int64_t used = 0;
+    for (int64_t i2 = 0; i2 < ids->ne[2]; ++i2) {
+        for (int64_t i1 = 0; i1 < ids->ne[1]; ++i1) {
+            const int32_t * row = (const int32_t *) (host.data() + i2 * ids->nb[2] + i1 * ids->nb[1]);
+            for (int64_t i0 = 0; i0 < ids->ne[0]; ++i0) {
+                const int32_t e = row[i0];
+                if (e >= 0 && e < n_experts) {
+                    if (counts[e] == 0) ++used;
+                    ++counts[e];
+                }
+            }
+        }
+    }
+    return used;
+}
+
+// Dedup hash of one node's geometry (destination, all sources, op params, types); distinct
+// shapes get distinct ids so the report can be deduplicated.
+uint64_t head_geometry_id(const op_record & rec, const ggml_tensor * node) {
+    const ggml_tensor * src0 = node->src[0];
+    uint64_t geometry_id = hash_mix(0, (uint64_t) node->op);
+    for (int d = 0; d < 4; d++) geometry_id = hash_mix(geometry_id, (uint64_t) rec.ne[d]);
+    for (int j = 0; j < GGML_MAX_SRC; j++) {
+        for (int d = 0; d < 4; d++) geometry_id = hash_mix(geometry_id, (uint64_t) rec.src_ne[j][d]);
+    }
+    for (int p = 0; p < n_op_params; p++) geometry_id = hash_mix(geometry_id, (uint64_t) (uint32_t) rec.op_params[p]);
+    geometry_id = hash_mix(geometry_id, (uint64_t) (src0 ? src0->type : 0));
+    geometry_id = hash_mix(geometry_id, (uint64_t) node->type);
+    return geometry_id;
+}
+
+// Buffer callback: records each kernel dispatch (kernel id + device time) under the
+// invocation that launched it. Invoked asynchronously as the tracing buffer fills.
+void buffer_callback(rocprofiler_context_id_t, rocprofiler_buffer_id_t,
+                     rocprofiler_record_header_t ** headers, size_t n_headers, void *, uint64_t) {
+    std::lock_guard<std::mutex> lock(g_mutex);
+    for (size_t i = 0; i < n_headers; i++) {
+        rocprofiler_record_header_t * header = headers[i];
+        if (header->category == ROCPROFILER_BUFFER_CATEGORY_TRACING &&
+            header->kind     == ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) {
+            auto * record = static_cast<rocprofiler_buffer_tracing_kernel_dispatch_record_t *>(header->payload);
+            const uint64_t invocation = record->correlation_id.external.value;
+            g_dispatches[invocation].push_back({record->start_timestamp,
+                                                record->dispatch_info.kernel_id,
+                                                record->end_timestamp - record->start_timestamp});
+        }
+    }
+}
+
+// Code-object callback: records kernel id -> demangled symbol as kernels load. The SDK
+// frees the name string on unload, so the demangled copy is kept.
+void code_object_callback(rocprofiler_callback_tracing_record_t record,
+                          rocprofiler_user_data_t *, void *) {
+    if (record.kind      != ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT ||
+        record.operation != ROCPROFILER_CODE_OBJECT_DEVICE_KERNEL_SYMBOL_REGISTER ||
+        record.phase     != ROCPROFILER_CALLBACK_PHASE_LOAD) {
+        return;
+    }
+    auto * symbol = static_cast<
+        rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t *>(record.payload);
+    std::lock_guard<std::mutex> lock(g_mutex);
+    g_kernel_names[symbol->kernel_id] = demangle(symbol->kernel_name);
+}
+
+void write_shape(std::ostringstream & out, const char * key, const int64_t ne[4]) {
+    out << "\"" << key << "\": [" << ne[0] << ", " << ne[1] << ", " << ne[2] << ", " << ne[3] << "]";
+}
+
+// One fused node: its geometry plus the raw per-tensor data the consumer needs to compute the
+// fused group's HBM traffic itself -- destination/source byte counts, a storage id per tensor
+// (which buffer it uses, so aliasing tensors are recognised) and a source tensor id (dedup key).
+// The consumer applies the exclusion policy (skip views, drop in-place destinations and internal
+// reads, dedup); the producer only records each tensor's size and identity.
+void write_fused_node(std::ostringstream & out, const op_record & rec) {
+    out << "{\"ggml_op\": \"" << rec.op << "\", "
+        << "\"dtype\": \"" << rec.dtype << "\", \"quant\": \"" << rec.quant << "\", ";
+    write_shape(out, "ne", rec.ne);  out << ", ";
+    out << "\"src_ne\": [";
+    for (int j = 0; j < rec.n_src; j++) {
+        if (j) out << ", ";
+        out << "[" << rec.src_ne[j][0] << ", " << rec.src_ne[j][1] << ", "
+            << rec.src_ne[j][2] << ", " << rec.src_ne[j][3] << "]";
+    }
+    out << "], \"src_types\": [";
+    for (int j = 0; j < rec.n_src; j++) {
+        if (j) out << ", ";
+        out << "\"" << (rec.src_types[j] ? rec.src_types[j] : "") << "\"";
+    }
+    out << "], \"op_params\": [";
+    for (int p = 0; p < n_op_params; p++) { if (p) out << ", "; out << rec.op_params[p]; }
+    out << "], \"dst_bytes\": " << rec.dst_bytes << ", \"dst_storage_id\": " << rec.dst_storage_id << ", ";
+    out << "\"src_bytes\": [";
+    for (int j = 0; j < rec.n_src; j++) { if (j) out << ", "; out << rec.src_bytes[j]; }
+    out << "], \"src_tensor_ids\": [";
+    for (int j = 0; j < rec.n_src; j++) { if (j) out << ", "; out << rec.src_tensor_ids[j]; }
+    out << "], \"src_storage_ids\": [";
+    for (int j = 0; j < rec.n_src; j++) { if (j) out << ", "; out << rec.src_storage_ids[j]; }
+    out << "], \"M\": " << rec.M << ", \"N\": " << rec.N << ", \"K\": " << rec.K
+        << ", \"n_experts\": " << rec.n_experts << ", \"top_k\": " << rec.top_k << "}";
+}
+
+// Write the per-invocation JSON report. Registered with atexit while profiling is active.
+void write_report() {
+    if (!g_active || !p_flush) return;
+    p_flush(g_buffer);
+
+    std::lock_guard<std::mutex> lock(g_mutex);
+    double total_us = 0.0;
+    for (auto & [invocation, dispatches] : g_dispatches) {
+        // invocation 0 is the sentinel used to shield MoE ids read-backs (see begin_op): any
+        // copy kernel it emitted is not part of any op, so keep it out of the total too.
+        if (invocation != 0) {
+            for (const auto & d : dispatches) total_us += d.duration_ns / 1e3;
+        }
+        // order kernels within an op causally (buffer records arrive unordered)
+        std::sort(dispatches.begin(), dispatches.end(),
+                  [](const dispatch & a, const dispatch & b) { return a.start_ns < b.start_ns; });
+    }
+
+    std::ostringstream out;
+    out << "{\n  \"device\": \"" << g_device << "\",\n";
+    out << "  \"total_gpu_time_us\": " << total_us << ",\n  \"rows\": [\n";
+    bool first = true;
+    for (const auto & [invocation, geometry_id] : g_invocations) {
+        auto dispatch_it = g_dispatches.find(invocation);
+        if (dispatch_it == g_dispatches.end()) continue;  // no kernels recorded for this op
+        auto record_it = g_records.find(geometry_id);
+        if (record_it == g_records.end()) continue;
+        const op_record & rec = record_it->second;
+
+        if (!first) out << ",\n";
+        first = false;
+
+        out << "    {\"ggml_op\": \"" << rec.op << "\", ";
+        if (!rec.fused_nodes.empty()) {
+            out << "\"fused_ops\": [";
+            for (size_t k = 0; k < rec.fused_nodes.size(); k++) {
+                if (k) out << ", ";
+                write_fused_node(out, rec.fused_nodes[k]);
+            }
+            out << "], ";
+        }
+        out << "\"dtype\": \"" << rec.dtype << "\", \"quant\": \"" << rec.quant << "\", ";
+        write_shape(out, "ne", rec.ne);  out << ", ";
+        out << "\"src_ne\": [";
+        for (int j = 0; j < rec.n_src; j++) {
+            if (j) out << ", ";
+            out << "[" << rec.src_ne[j][0] << ", " << rec.src_ne[j][1] << ", "
+                << rec.src_ne[j][2] << ", " << rec.src_ne[j][3] << "]";
+        }
+        out << "], ";
+        out << "\"src_types\": [";
+        for (int j = 0; j < rec.n_src; j++) {
+            if (j) out << ", ";
+            out << "\"" << (rec.src_types[j] ? rec.src_types[j] : "") << "\"";
+        }
+        out << "], ";
+        out << "\"op_params\": [";
+        for (int p = 0; p < n_op_params; p++) { if (p) out << ", "; out << rec.op_params[p]; }
+        out << "], ";
+        out << "\"bytes\": " << rec.bytes << ", ";
+        out << "\"dst_bytes\": " << rec.dst_bytes << ", \"dst_storage_id\": " << rec.dst_storage_id << ", ";
+        out << "\"src_bytes\": [";
+        for (int j = 0; j < rec.n_src; j++) { if (j) out << ", "; out << rec.src_bytes[j]; }
+        out << "], \"src_tensor_ids\": [";
+        for (int j = 0; j < rec.n_src; j++) { if (j) out << ", "; out << rec.src_tensor_ids[j]; }
+        out << "], \"src_storage_ids\": [";
+        for (int j = 0; j < rec.n_src; j++) { if (j) out << ", "; out << rec.src_storage_ids[j]; }
+        out << "], ";
+        // experts_used: distinct experts this specific launch routed, read back per invocation
+        // (0 for non-MoE ops). The consumer scales the expert-weight HBM traffic by this instead
+        // of the shape bound min(M*top_k, E), which over-counts when routing leaves experts idle.
+        auto experts_it = g_invocation_experts.find(invocation);
+        const int64_t experts_used = experts_it != g_invocation_experts.end() ? experts_it->second : 0;
+        out << "\"M\": " << rec.M << ", \"N\": " << rec.N << ", \"K\": " << rec.K
+            << ", \"n_experts\": " << rec.n_experts << ", \"top_k\": " << rec.top_k
+            << ", \"experts_used\": " << experts_used << ", ";
+        // tokens_per_expert: per-expert routed-token counts for this launch (length n_experts,
+        // index e = routings to expert e). Present only for MoE launches whose ids were read back;
+        // lets the consumer render the routing-load distribution under --show-histograms.
+        auto hist_it = g_invocation_expert_hist.find(invocation);
+        if (hist_it != g_invocation_expert_hist.end()) {
+            out << "\"tokens_per_expert\": [";
+            for (size_t e = 0; e < hist_it->second.size(); ++e) { if (e) out << ", "; out << hist_it->second[e]; }
+            out << "], ";
+        }
+        out << "\"kernels\": [";
+        bool kernel_first = true;
+        for (const auto & d : dispatch_it->second) {
+            if (!kernel_first) out << ", ";
+            kernel_first = false;
+            auto name_it = g_kernel_names.find(d.kernel_id);
+            out << "{\"name\": \"";
+            json_escape(out, name_it != g_kernel_names.end() ? name_it->second : "");
+            out << "\", \"gpu_time_us\": " << d.duration_ns / 1e3 << "}";
+        }
+        out << "]}";
+    }
+    out << "\n  ]\n}\n";
+
+    std::ofstream file(g_out_path);
+    file << out.str();
+    // runs from atexit, where the ggml logger is already torn down, so write directly
+    fprintf(stderr, "ggml-roofline: wrote %s (%zu invocations, %zu unique ops, total GPU %.1f us)\n",
+            g_out_path.c_str(), g_invocations.size(), g_records.size(), total_us);
+}
+
+int tool_init(rocprofiler_client_finalize_t, void *) {
+    if (!resolve_symbols()) {
+        GGML_LOG_ERROR("ggml-roofline: could not resolve rocprofiler-sdk symbols; disabled\n");
+        g_active = false;
+        return -1;
+    }
+    if (p_create_context(&g_context) != ROCPROFILER_STATUS_SUCCESS) return -1;
+    // 256 MiB buffer, 4 MiB watermark: batch dispatch records so the callback fires in
+    // large chunks instead of per record.
+    if (p_create_buffer(g_context, 256 * 1024 * 1024, 4 * 1024 * 1024,
+                        ROCPROFILER_BUFFER_POLICY_LOSSLESS, buffer_callback, nullptr, &g_buffer)
+            != ROCPROFILER_STATUS_SUCCESS) return -1;
+    if (p_config_buffer(g_context, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, g_buffer)
+            != ROCPROFILER_STATUS_SUCCESS) return -1;
+    // Kernel-symbol names are optional: timing works without them, kernels just omit "name".
+    if (p_config_callback) {
+        p_config_callback(g_context, ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT, nullptr, 0,
+                          code_object_callback, nullptr);
+    }
+    if (p_start(g_context) != ROCPROFILER_STATUS_SUCCESS) return -1;
+    return 0;
+}
+
+void tool_fini(void *) {}
+
+rocprofiler_tool_configure_result_t g_configure_result{
+    sizeof(rocprofiler_tool_configure_result_t), &tool_init, &tool_fini, nullptr};
+
+rocprofiler_tool_configure_result_t *
+configure(uint32_t, const char *, uint32_t, rocprofiler_client_id_t * client_id) {
+    client_id->name = "ggml-roofline";
+    g_active = true;
+    return &g_configure_result;
+}
+
+using fn_force_configure = rocprofiler_status_t (*)(
+    rocprofiler_tool_configure_result_t * (*)(uint32_t, const char *, uint32_t, rocprofiler_client_id_t *));
+
+} // namespace
+
+void ggml_cuda_roofline_init(void) {
+    static std::once_flag once;
+    std::call_once(once, [] {
+        const char * out_path = getenv("GGML_ROOFLINE_OUT");
+        if (!out_path || !out_path[0]) return;
+        g_out_path = out_path;
+
+        // Per-op attribution requires the eager node loop, so graphs must be disabled.
+        // The backend reads this env var lazily and only checks for its presence; set it
+        // here before that happens, without overriding an explicit user setting.
+        if (setenv("GGML_CUDA_DISABLE_GRAPHS", "1", 0) == 0) {
+            GGML_LOG_INFO("ggml-roofline: graphs disabled (required for per-op attribution)\n");
+        }
+
+        void * handle = dlopen("librocprofiler-sdk.so.1", RTLD_NOW | RTLD_GLOBAL);
+        if (!handle) handle = dlopen("librocprofiler-sdk.so", RTLD_NOW | RTLD_GLOBAL);
+        if (!handle) {
+            GGML_LOG_ERROR("ggml-roofline: cannot dlopen librocprofiler-sdk: %s\n", dlerror());
+            return;
+        }
+        auto force_configure = (fn_force_configure) dlsym(handle, "rocprofiler_force_configure");
+        if (!force_configure || force_configure(&configure) != ROCPROFILER_STATUS_SUCCESS) {
+            GGML_LOG_ERROR("ggml-roofline: rocprofiler_force_configure failed\n");
+            return;
+        }
+        std::atexit(write_report);
+        GGML_LOG_INFO("ggml-roofline: enabled -> %s\n", g_out_path.c_str());
+    });
+}
+
+void ggml_cuda_roofline_set_device(const char * arch) {
+    if (g_active && arch) g_device = arch;
+}
+
+void ggml_cuda_roofline_reset(void) {
+    if (!g_active) return;
+    // Drain any pending records first so warmup dispatches are accounted, then dropped.
+    if (p_flush) p_flush(g_buffer);
+    std::lock_guard<std::mutex> lock(g_mutex);
+    g_records.clear();
+    g_invocations.clear();
+    g_dispatches.clear();
+    g_invocation_experts.clear();
+    // g_next_invocation stays monotonic so a late warmup record cannot collide with a
+    // post-reset invocation id; g_kernel_names is kept (code objects do not reload).
+}
+
+void ggml_cuda_roofline_begin_op(const struct ggml_tensor * node, void * stream) {
+    if (!g_active || node == nullptr || !p_push_id) return;
+
+    op_record rec;
+    fill_head_record(rec, node);
+    const uint64_t geometry_id = head_geometry_id(rec, node);
+
+    // Each op invocation gets a unique correlation id so its kernels stay separate; the
+    // shared geometry is stored once per shape.
+    const uint64_t invocation = g_next_invocation.fetch_add(1, std::memory_order_relaxed);
+
+    // rocprofiler keeps a per-thread external-correlation-id stack; the id is captured at each
+    // kernel dispatch. This thread id is needed for that bookkeeping and for the MoE read below,
+    // so resolve it up front.
+    static thread_local rocprofiler_thread_id_t thread_id = [] {
+        rocprofiler_thread_id_t t = 0; if (p_get_thread_id) p_get_thread_id(&t); return t;
+    }();
+
+    // MoE: capture how many distinct experts this launch actually routes (see count_active_experts).
+    // Reading the ids tensor may lower to a copy kernel; push a sentinel correlation id (0, never a
+    // real invocation) around it so such a dispatch is attributed to no op and skipped by
+    // write_report, instead of polluting the previous op still on the stack.
+    if (stream && node->op == GGML_OP_MUL_MAT_ID && node->src[2]) {
+        rocprofiler_user_data_t sentinel; sentinel.value = 0;
+        p_push_id(g_context, thread_id, sentinel);
+        std::vector<int32_t> hist;
+        const int64_t used = count_active_experts(node->src[2], rec.n_experts, (hipStream_t) stream, hist);
+        rocprofiler_user_data_t popped;
+        p_pop_id(g_context, thread_id, &popped);
+        std::lock_guard<std::mutex> lock(g_mutex);
+        g_invocation_experts[invocation] = used;
+        if (!hist.empty()) g_invocation_expert_hist[invocation] = std::move(hist);
+    }
+
+    {
+        std::lock_guard<std::mutex> lock(g_mutex);
+        g_invocations.emplace(invocation, geometry_id);
+        if (g_records.find(geometry_id) == g_records.end()) g_records.emplace(geometry_id, rec);
+    }
+
+    // Tag the kernels launched until the next op with this invocation id. rocprofiler
+    // keeps a per-thread stack, so pop the previous id before pushing the new one.
+    static thread_local bool pushed = false;
+    if (pushed) {
+        rocprofiler_user_data_t previous;
+        p_pop_id(g_context, thread_id, &previous);
+    }
+    rocprofiler_user_data_t current;
+    current.value = invocation;
+    p_push_id(g_context, thread_id, current);
+    pushed = true;
+
+    // Remember which invocation this thread just tagged so a following fusion can override
+    // its record to cover the whole fused span (see ggml_cuda_roofline_fuse_ops).
+    g_current_invocation = invocation;
+}
+
+void ggml_cuda_roofline_fuse_ops(const struct ggml_cgraph * cgraph, int node_idx, int node_count) {
+    if (!g_active || cgraph == nullptr || node_count < 2 || g_current_invocation == 0) return;
+
+    const ggml_tensor * head = cgraph->nodes[node_idx];
+
+    // Keep the head op's geometry (op name, shapes, M/N/K) so the row is still labelled by the
+    // head op. The row-level byte fields stay head-only and are ignored for fused rows; the
+    // consumer derives the group's traffic from the per-node data recorded below.
+    op_record rec;
+    fill_head_record(rec, head);
+
+    // Record every fused node's geometry plus its raw per-tensor byte counts and storage/tensor
+    // ids (captured by fill_head_record while the graph is live). The consumer sums exact FLOPs
+    // across the group and computes its external HBM traffic from them -- see write_fused_node.
+    rec.fused_nodes.reserve(node_count);
+    for (int j = node_idx; j < node_idx + node_count; ++j) {
+        op_record sub;
+        fill_head_record(sub, cgraph->nodes[j]);
+        rec.fused_nodes.push_back(std::move(sub));
+    }
+
+    // Fused geometry id: head geometry plus each fused node's op and destination shape, so
+    // identical fusions deduplicate and distinct ones stay separate.
+    uint64_t geometry_id = head_geometry_id(rec, head);
+    for (int j = node_idx; j < node_idx + node_count; ++j) {
+        const ggml_tensor * n = cgraph->nodes[j];
+        geometry_id = hash_mix(geometry_id, (uint64_t) n->op);
+        for (int d = 0; d < 4; d++) geometry_id = hash_mix(geometry_id, (uint64_t) n->ne[d]);
+    }
+
+    // Re-point the current invocation at the fused record. The provisional head-only record
+    // from begin_op stays in g_records; if no non-fused invocation references it, it is
+    // simply never emitted (the report iterates g_invocations).
+    std::lock_guard<std::mutex> lock(g_mutex);
+    g_invocations[g_current_invocation] = geometry_id;
+    if (g_records.find(geometry_id) == g_records.end()) g_records.emplace(geometry_id, std::move(rec));
+}
diff --git a/ggml/src/ggml-cuda/ggml-cuda-roofline.h b/ggml/src/ggml-cuda/ggml-cuda-roofline.h
new file mode 100644
index 000000000000..10db062cbd25
--- /dev/null
+++ b/ggml/src/ggml-cuda/ggml-cuda-roofline.h
@@ -0,0 +1,42 @@
+#pragma once
+
+// Per-op GPU roofline profiling for the HIP/ROCm backend (see ggml-cuda-roofline.cpp).
+// Compiled only when GGML_HIP_ROOFLINE is defined and activated at runtime by the
+// environment variable GGML_ROOFLINE_OUT=<path.json>. Every entry point is a no-op
+// unless that variable is set.
+
+#include "ggml.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+// Load and configure rocprofiler-sdk when GGML_ROOFLINE_OUT is set. Call once during
+// backend registration, before any GPU stream is created. Idempotent; no-op otherwise.
+void ggml_cuda_roofline_init(void);
+
+// Record the GPU architecture (e.g. "gfx1151") stored in the report. No-op unless active.
+void ggml_cuda_roofline_set_device(const char * arch);
+
+// Discard everything captured so far (drains pending records first). Call after a
+// warmup run so the report covers only the measured run. No-op unless active.
+void ggml_cuda_roofline_reset(void);
+
+// Tag the GPU kernels launched for this op so their device time is attributed to it.
+// Call once per op, before its kernel(s) are dispatched. No-op unless active.
+//
+// stream is the op's CUDA/HIP stream (passed as void * to keep this header free of the HIP
+// runtime): for MoE (MUL_MAT_ID) it is synchronized so the routing ids tensor can be read
+// back and the distinct active-expert count captured for this specific launch. Pass nullptr
+// if unavailable (the MoE expert count is then skipped).
+void ggml_cuda_roofline_begin_op(const struct ggml_tensor * node, void * stream);
+
+// Override the record of the op tagged by the last begin_op so it covers a fused span of
+// node_count nodes (cgraph->nodes[node_idx .. node_idx+node_count-1]): lists every fused op
+// and reports the fused group's HBM traffic with intermediates discarded. Call right after a
+// successful fusion, before advancing the loop. No-op unless active.
+void ggml_cuda_roofline_fuse_ops(const struct ggml_cgraph * cgraph, int node_idx, int node_count);
+
+#ifdef __cplusplus
+}
+#endif
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 98816f885cf6..0ac6239f7480 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -2,6 +2,10 @@
 #include "ggml-impl.h"
 #include "ggml-backend-impl.h"
 
+#ifdef GGML_HIP_ROOFLINE
+#include "ggml-cuda/ggml-cuda-roofline.h"
+#endif
+
 #include "ggml-cuda/allreduce.cuh"
 #include "ggml-cuda/common.cuh"
 #include "ggml-cuda/acc.cuh"
@@ -85,6 +89,7 @@
 #include <cstdio>
 #include <cstdlib>
 #include <string>
+#include <unordered_set>
 #include <vector>
 
 static_assert(sizeof(half) == sizeof(ggml_fp16_t), "wrong fp16 size");
@@ -644,10 +649,15 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() {
                 CUDA_CHECK(cudaStreamDestroy(streams[i][j]));
             }
         }
-        if (cublas_handles[i] != nullptr) {
-            CUBLAS_CHECK(cublasDestroy(cublas_handles[i]));
+        for (int j = 0; j < GGML_CUDA_MAX_STREAMS; ++j) {
+            if (cublas_handles[i][j] != nullptr) {
+                CUBLAS_CHECK(cublasDestroy(cublas_handles[i][j]));
+            }
         }
     }
+    if (concurrent_scratch != nullptr) {
+        ggml_backend_buffer_free(concurrent_scratch);
+    }
 }
 
 
@@ -1755,6 +1765,27 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
         ggml_cuda_mul_mat_vec_f(ctx, src0, src1, nullptr, dst);
         return;
     }
+    // A single-column f32 weight (ne01 == 1) times a wide activation (ne11 > MMVF_MAX_BATCH_SIZE)
+    // misses mmvf above because ne11 is the batch dim, so it falls back to a wasteful f32 cuBLAS
+    // GEMM. The output is a vector, so swapping the operands is transpose-free: the activation
+    // becomes the matrix and the single-column weight the lone F32 vector, with dst reinterpreted
+    // as [ne11, 1].
+    if (ne01 == 1 && ne11 > MMVF_MAX_BATCH_SIZE && ne2 == 1 && ne3 == 1
+            && src0->type == GGML_TYPE_F32
+            && ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst)
+            && ggml_cuda_should_use_mmvf(src1->type, cc, src1->ne, src1->nb, /*ne11 =*/ 1)) {
+        // dst is [ne00=1, ne11] (one output feature, ne11 tokens); the swapped GEMV instead
+        // produces [ne11, 1] (ne11 rows, a single vector column). Both are the same ne11
+        // contiguous floats, so we only relabel the shape of a shallow copy - no data moves:
+        ggml_tensor dst_vec = *dst;
+        dst_vec.ne[0] = ne11;                    // rows: one result per token
+        dst_vec.ne[1] = 1;                       // a single output vector column
+        dst_vec.nb[1] = dst_vec.nb[0]*ne11;      // nb[0] (element stride) stays; row dim spans ne11
+        dst_vec.nb[2] = dst_vec.nb[1];           // higher dims are 1 -> same stride
+        dst_vec.nb[3] = dst_vec.nb[1];
+        ggml_cuda_mul_mat_vec_f(ctx, src1, src0, nullptr, &dst_vec);
+        return;
+    }
     if (ggml_cuda_should_use_mmf(src0->type, cc, warp_size, src0->ne, src0->nb, ne11, /*mul_mat_id =*/ false)) {
         ggml_cuda_mul_mat_f(ctx, src0, src1, nullptr, dst);
         return;
@@ -3870,8 +3901,11 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud
                         is_concurrent_event_active = false;
                         concurrent_event           = nullptr;
                     } else {
-                        GGML_ASSERT (concurrent_event->stream_mapping.find(node) != concurrent_event->stream_mapping.end());
-                        cuda_ctx->curr_stream_no = concurrent_event->stream_mapping[node];
+                        // region nodes not mapped to a concurrent stream run on the main stream:
+                        // this keeps the routed branch on the main stream while only the shared
+                        // expert forks off (the vLLM shared-expert model)
+                        auto it = concurrent_event->stream_mapping.find(node);
+                        cuda_ctx->curr_stream_no = it != concurrent_event->stream_mapping.end() ? it->second : 0;
                         GGML_LOG_DEBUG("Setting stream no to %d for node %s\n", cuda_ctx->curr_stream_no, node->name);
                     }
                 } else if (i - prev_i > 1) {
@@ -3880,7 +3914,8 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud
                     try_launch_concurrent_event(prev_node);
 
                     if (is_concurrent_event_active) {
-                        cuda_ctx->curr_stream_no = concurrent_event->stream_mapping[node];
+                        auto it = concurrent_event->stream_mapping.find(node);
+                        cuda_ctx->curr_stream_no = it != concurrent_event->stream_mapping.end() ? it->second : 0;
                         GGML_LOG_DEBUG("Setting stream no to %d for node %s\n", cuda_ctx->curr_stream_no, node->name);
                     }
                 }
@@ -3895,9 +3930,16 @@ static void ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context * cud
                     continue;
                 }
 
+#ifdef GGML_HIP_ROOFLINE
+                ggml_cuda_roofline_begin_op(node, (void *) cuda_ctx->stream());
+#endif
+
                 int nodes_to_skip = ggml_cuda_try_fuse(cuda_ctx, cgraph, i);
 
                 if (nodes_to_skip != 0) {
+#ifdef GGML_HIP_ROOFLINE
+                    ggml_cuda_roofline_fuse_ops(cgraph, i, nodes_to_skip + 1);
+#endif
 #ifdef GGML_CUDA_DEBUG
                     const int last_fused = i + nodes_to_skip;
                     GGML_LOG_INFO("nodes_fused: %d, first: %s (%s), last: %s (%s)\n",
@@ -4154,6 +4196,11 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph
     // store {fork_idx, join_idx}
     std::vector<std::pair<int, int>> concurrent_node_ranges;
 
+    // per-event lists of concurrent branch nodes to place in the dedicated scratch buffer (below),
+    // so the branches are mutually disjoint and disjoint from tensors read across the region -
+    // this replaces the fragile node interleaving that ggml-alloc/execution order can desync
+    std::vector<std::vector<const ggml_tensor *>> concurrent_groups;
+
     for (const auto & [root_node, count] : fan_out) {
         if (count >= min_fan_out && count <= max_fan_out) {
             const int root_node_idx = node_indices[root_node];
@@ -4245,10 +4292,6 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph
                 int fork_node_idx = node_indices[root_node];
                 int join_node_idx = node_indices[join_node];
 
-                int       current_branch_idx = 0;
-                int       current_node_idx   = fork_node_idx + 1;
-                const int n_branches         = nodes_per_branch.size();
-
                 int total_branch_nodes = 0;
                 for (std::vector<const ggml_tensor *> branch_nodes : nodes_per_branch) {
                     total_branch_nodes += branch_nodes.size();
@@ -4277,37 +4320,182 @@ static void ggml_backend_cuda_graph_optimize(ggml_backend_t backend, ggml_cgraph
                 GGML_LOG_DEBUG("Adding stream at node %s %p\n", root_node->name, root_node);
                 concurrent_node_ranges.emplace_back(fork_node_idx, join_node_idx);
 
-                // interleave tensors to extend lifetimes so that ggml graph doesn't recycle them
-                // example transformation:
-                // [attn-norm, QMul, QNorm, QRope, KMul, KNorm, KRope, VMul, attn] ->
-                // [attn-norm, QMul, KMul, VMul, QNorm, VNorm, QRope, KRope, attn]
-                while (current_node_idx < join_node_idx) {
-                    std::vector<const ggml_tensor *> & branch_nodes = nodes_per_branch[current_branch_idx];
-
-                    bool has_node = false;
-                    for (std::vector<const ggml_tensor *> branch_node : nodes_per_branch) {
-                        has_node |= branch_node.size() > 0;
+                // place all branch nodes in the dedicated scratch buffer (below) instead of
+                // interleaving them: ggml-alloc then keeps the branches mutually disjoint and
+                // disjoint from the fork output that every branch reads concurrently
+                std::vector<const ggml_tensor *> group;
+                for (const auto & branch_nodes : nodes_per_branch) {
+                    for (const ggml_tensor * n : branch_nodes) {
+                        group.push_back(n);
                     }
+                }
+                concurrent_groups.push_back(std::move(group));
+            }
+        }
+    }
+
+    // MoE shared-expert overlap: run the shared expert on a separate stream, overlapped with the
+    // routed experts. fork = the FFN-input norm feeding both branches, join = ggml_add(ffn_moe_out,
+    // ffn_shexp*). Operands are matched by the names set via cb() in the model graph. Decode only
+    // (gated below): prefill is compute-bound and gains nothing from the overlap.
+    const auto reach_backward = [](const ggml_tensor * start) {
+        std::unordered_set<const ggml_tensor *> seen;
+        std::vector<const ggml_tensor *> stack = { start };
+        while (!stack.empty()) {
+            const ggml_tensor * t = stack.back();
+            stack.pop_back();
+            if (!t || seen.count(t)) {
+                continue;
+            }
+            seen.insert(t);
+            for (int s = 0; s < GGML_MAX_SRC; ++s) {
+                if (t->src[s]) {
+                    stack.push_back(t->src[s]);
+                }
+            }
+        }
+        return seen;
+    };
 
-                    GGML_ASSERT(has_node);
+    for (int join_idx = 0; join_idx < cgraph->n_nodes; ++join_idx) {
+        ggml_tensor * join_node = cgraph->nodes[join_idx];
+        if (join_node->op != GGML_OP_ADD) {
+            continue;
+        }
 
-                    if (branch_nodes.empty()) {
-                        current_branch_idx = (current_branch_idx + 1) % n_branches;
-                        continue;
-                    }
+        // Only overlap during decode (single token). Overlapping the shared expert only helps when
+        // the routed branch leaves the GPU underutilized for it to run alongside; that is the case in
+        // decode (batch 1, latency/occupancy-bound) but not in prefill, where the routed matmuls are
+        // large and already saturate the GPU, so the overlap adds contention without a speedup.
+        if (ggml_nrows(join_node) > 1) {
+            continue;
+        }
 
-                    cgraph->nodes[current_node_idx] = const_cast<ggml_tensor *>(branch_nodes.front());
-                    current_node_idx++;
-                    branch_nodes.erase(branch_nodes.begin());
+        ggml_tensor * routed_out = nullptr;
+        ggml_tensor * shexp_out  = nullptr;
+        for (int s = 0; s < 2; ++s) {
+            ggml_tensor * x = join_node->src[s];
+            ggml_tensor * y = join_node->src[1 - s];
+            if (x && y && strstr(x->name, "ffn_moe_out") && strstr(y->name, "ffn_shexp")) {
+                routed_out = x;
+                shexp_out  = y;
+            }
+        }
+        if (!routed_out || !shexp_out) {
+            continue;
+        }
 
-                    // append all empty nodes
-                    while (!branch_nodes.empty() && is_noop(branch_nodes.front())) {
-                        cgraph->nodes[current_node_idx] = const_cast<ggml_tensor *>(branch_nodes.front());
-                        current_node_idx++;
-                        branch_nodes.erase(branch_nodes.begin());
-                    }
+        const std::unordered_set<const ggml_tensor *> reach_routed = reach_backward(routed_out);
+        const std::unordered_set<const ggml_tensor *> reach_shexp  = reach_backward(shexp_out);
+
+        // fork = highest-index node reachable from both branches (the ffn_norm output)
+        int fork_idx = -1;
+        for (const ggml_tensor * t : reach_routed) {
+            if (!reach_shexp.count(t)) {
+                continue;
+            }
+            auto it = node_indices.find(t);
+            if (it != node_indices.end() && it->second < join_idx && it->second > fork_idx) {
+                fork_idx = it->second;
+            }
+        }
+        if (fork_idx < 0) {
+            continue;
+        }
+
+        bool overlaps = false;
+        for (const auto & [start, end] : concurrent_node_ranges) {
+            if (!(join_idx < start || fork_idx > end)) {
+                overlaps = true;
+            }
+        }
+        if (overlaps) {
+            continue;
+        }
+
+        // partition the region (fork_idx, join_idx): shared-expert nodes -> stream 2, routed -> 1
+        std::vector<std::vector<const ggml_tensor *>> nodes_per_branch(2);
+        for (int i = fork_idx + 1; i < join_idx; ++i) {
+            const ggml_tensor * n = cgraph->nodes[i];
+            const int branch = reach_shexp.count(n) ? 1 : 0;
+            nodes_per_branch[branch].push_back(n);
+        }
+        if (nodes_per_branch[0].empty() || nodes_per_branch[1].empty()) {
+            continue;
+        }
+
+        // vLLM shared-expert model: the routed experts stay on the main stream and only the shared
+        // expert forks onto a single aux stream, joined at the add. Keeping the large routed branch
+        // on the main stream avoids migrating it and needs only one fork/join.
+        ggml_cuda_concurrent_event concurrent_event(1);
+        concurrent_event.join_node = join_node;
+        for (const ggml_tensor * n : nodes_per_branch[1]) {
+            concurrent_event.stream_mapping[n] = 1;
+        }
+
+        const ggml_tensor * fork_node = cgraph->nodes[fork_idx];
+        concurrent_event.original_order.reserve(join_idx - fork_idx - 1);
+        for (int i = fork_idx + 1; i < join_idx; ++i) {
+            concurrent_event.original_order.push_back(cgraph->nodes[i]);
+        }
+
+        std::unordered_map<const ggml_tensor *, ggml_cuda_concurrent_event> & concurrent_events = cuda_ctx->stream_context().concurrent_events;
+        if (concurrent_events.find(fork_node) != concurrent_events.end()) {
+            continue;
+        }
+        concurrent_events.emplace(fork_node, std::move(concurrent_event));
+        GGML_LOG_DEBUG("Adding shared-expert stream at node %s %p\n", fork_node->name, fork_node);
+        concurrent_node_ranges.emplace_back(fork_idx, join_idx);
+
+        // the shared-expert nodes get a dedicated buffer (below), so the graph order is left intact
+        // and no interleaving is needed to keep the branch non-overlapping
+        concurrent_groups.push_back(nodes_per_branch[1]);
+    }
+
+    // Place every concurrent branch (attention QKV and MoE shared-expert) in a dedicated buffer so
+    // its nodes never share an address with each other or with tensors read across the region (which
+    // ggml-alloc could otherwise recycle, corrupting concurrent reads). Layers run sequentially, so
+    // one buffer sized to the largest region is reused across all of them; within a region each node
+    // gets a distinct offset so the concurrent scratch stays disjoint.
+    if (!concurrent_groups.empty()) {
+        const size_t alignment = 128;
 
-                    current_branch_idx = (current_branch_idx + 1) % n_branches;
+        const auto group_footprint = [&](const std::vector<const ggml_tensor *> & group) {
+            size_t off = 0;
+            for (const ggml_tensor * n : group) {
+                if (is_noop(n) || n->view_src != nullptr) {
+                    continue;
+                }
+                off += GGML_PAD(ggml_nbytes(n), alignment);
+            }
+            return off;
+        };
+
+        size_t needed = 0;
+        for (const auto & group : concurrent_groups) {
+            needed = std::max(needed, group_footprint(group));
+        }
+
+        if (needed > 0) {
+            if (cuda_ctx->concurrent_scratch == nullptr || cuda_ctx->concurrent_scratch_size < needed) {
+                if (cuda_ctx->concurrent_scratch != nullptr) {
+                    ggml_backend_buffer_free(cuda_ctx->concurrent_scratch);
+                }
+                cuda_ctx->concurrent_scratch      = ggml_backend_buft_alloc_buffer(ggml_backend_cuda_buffer_type(cuda_ctx->device), needed);
+                cuda_ctx->concurrent_scratch_size = needed;
+            }
+
+            char * const base = (char *) ggml_backend_buffer_get_base(cuda_ctx->concurrent_scratch);
+            for (const auto & group : concurrent_groups) {
+                size_t off = 0;
+                for (const ggml_tensor * cn : group) {
+                    if (is_noop(cn) || cn->view_src != nullptr) {
+                        continue;
+                    }
+                    ggml_tensor * n = const_cast<ggml_tensor *>(cn);
+                    n->data   = base + off;
+                    n->buffer = cuda_ctx->concurrent_scratch;
+                    off += GGML_PAD(ggml_nbytes(n), alignment);
                 }
             }
         }
@@ -4891,8 +5079,10 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
             }
         }
         case GGML_OP_SSM_CONV: {
-            // assumes d_inner % threads == 0
-            return op->src[0]->ne[1] % 128 == 0;
+            // assumes d_inner % threads == 0; d_inner is on ne[1] (time-major) or ne[0] (channels-major)
+            const bool channels_major = ggml_ssm_conv_get_layout(op) == GGML_SSM_CONV_LAYOUT_CHANNELS_MAJOR;
+            const int64_t d_inner = channels_major ? op->src[0]->ne[0] : op->src[0]->ne[1];
+            return d_inner % 128 == 0;
         }
         case GGML_OP_CONT:
             return true;
@@ -5162,6 +5352,14 @@ static const ggml_backend_reg_i ggml_backend_cuda_reg_interface = {
 
 // backend registry
 ggml_backend_reg_t ggml_backend_cuda_reg() {
+#ifdef GGML_HIP_ROOFLINE
+    ggml_cuda_roofline_init();  // configure rocprofiler tracing before any stream is created
+    if (ggml_cuda_info().device_count > 0) {
+        char arch[32];
+        snprintf(arch, sizeof(arch), "gfx%x", ggml_cuda_info().devices[0].cc & 0xffff);
+        ggml_cuda_roofline_set_device(arch);
+    }
+#endif
     static ggml_backend_reg reg;
     static bool initialized = false;
 
diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh
index edf546d8f1e2..3f8976d2c18d 100644
--- a/ggml/src/ggml-cuda/mmq.cuh
+++ b/ggml/src/ggml-cuda/mmq.cuh
@@ -140,6 +140,9 @@ static constexpr __device__ int get_mmq_x_max_device() {
 }
 
 static int get_mmq_y_host(const int cc) {
+    if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
+        return 64;
+    }
     return GGML_CUDA_CC_IS_AMD(cc) ? (GGML_CUDA_CC_IS_RDNA1(cc) ? 64 : 128) :
         ((GGML_CUDA_CC_IS_NVIDIA(cc) && ggml_cuda_highest_compiled_arch(cc) >= GGML_CUDA_CC_VOLTA) ? 128 : 64);
 }
@@ -155,7 +158,9 @@ if (type == GGML_TYPE_NVFP4 || type == GGML_TYPE_MXFP4) {
 
 static constexpr __device__ int get_mmq_y_device() {
 #if defined(GGML_USE_HIP)
-#if defined(RDNA1)
+#if defined(RDNA3_5)
+    return 64;
+#elif defined(RDNA1)
     return 64;
 #else
     return 128;
@@ -296,6 +301,9 @@ static constexpr __device__ int mmq_get_granularity_device(const int /*mmq_x*/)
 
 #if defined(GGML_USE_HIP)
 static int mmq_get_nwarps_host(const int cc, const int warp_size) {
+    if (GGML_CUDA_CC_IS_RDNA3_5(cc)) {
+        return 4;
+    }
     return amd_mfma_available(cc) ? 8 : 256/warp_size;
 }
 #else
@@ -305,7 +313,9 @@ static int mmq_get_nwarps_host(const int /*cc*/, const int warp_size) {
 #endif // (GGML_USE_HIP)
 
 static constexpr __device__ int mmq_get_nwarps_device() {
-#if defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+#if defined(RDNA3_5)
+    return 4;
+#elif defined(AMD_MFMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
     return 8;
 #else
     return 256/ggml_cuda_get_physical_warp_size();
@@ -4066,6 +4076,18 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda
     int mmq_x_best  = 0;
     int ntiles_x_best = INT_MAX;
 
+    // ncols_max assumes every token can land in a single tile column range; for MoE that is the
+    // worst case of all tokens routed to one expert, but experts typically receive only
+    // ncols_dst/nchannels_y tokens. Tiling over ~2x that average keeps the common per-expert
+    // tiles filled instead of mostly empty; the 2x covers routing imbalance and larger batches
+    // converge back to ncols_max.
+    // Restricted to RDNA3.5, the only arch this MoE tile sizing was tuned/validated on.
+    int64_t ncols_to_tile = args.ncols_max;
+    if (args.expert_bounds != nullptr && GGML_CUDA_CC_IS_RDNA3_5(cc)) {
+        const int64_t ncols_per_expert = (args.ncols_dst + args.nchannels_y - 1) / args.nchannels_y;
+        ncols_to_tile = 2*ncols_per_expert < args.ncols_max ? 2*ncols_per_expert : args.ncols_max;
+    }
+
     for (int mmq_x = 8; mmq_x <= mmq_x_max && ntiles_x_best > 1; mmq_x += 8) {
         const int granularity = mmq_get_granularity_host(mmq_x, cc);
 
@@ -4073,7 +4095,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda
             continue;
         }
 
-        const int ntiles_x = (args.ncols_max + mmq_x - 1) / mmq_x;
+        const int ntiles_x = (ncols_to_tile + mmq_x - 1) / mmq_x;
 
         if (ntiles_x < ntiles_x_best) {
             mmq_x_best = mmq_x;
diff --git a/ggml/src/ggml-cuda/ssm-conv.cu b/ggml/src/ggml-cuda/ssm-conv.cu
index 1463169cf78b..d1052772de39 100644
--- a/ggml/src/ggml-cuda/ssm-conv.cu
+++ b/ggml/src/ggml-cuda/ssm-conv.cu
@@ -2,7 +2,11 @@
 #include "ssm-conv.cuh"
 #include "unary.cuh"
 
-template <bool apply_silu, size_t split_d_inner, size_t d_conv>
+// channels_major selects the input (src0) memory layout:
+//   false = time-major     : src0 = [d_conv-1+n_t, d_inner, n_s], time contiguous (nb0)
+//   true  = channels-major : src0 = [d_inner, d_conv-1+n_t, n_s], channels contiguous (nb0)
+// The output (dst) is [d_inner, n_t, n_s] (channels contiguous) in both cases.
+template <bool apply_silu, bool channels_major, size_t split_d_inner, size_t d_conv>
 static __global__ void ssm_conv_f32(const float * src0_ptr, const float * src1_ptr,
                                     const float * bias_ptr,
                                     const int src0_nb0, const int src0_nb1, const int src0_nb2, const int src1_nb1,
@@ -13,18 +17,22 @@ static __global__ void ssm_conv_f32(const float * src0_ptr, const float * src1_p
     const float * GGML_CUDA_RESTRICT src1 = src1_ptr;
     const float * GGML_CUDA_RESTRICT bias = bias_ptr;
     float       * GGML_CUDA_RESTRICT dst  = dst_ptr;
-    GGML_UNUSED(src0_nb0);
     const int tid  = threadIdx.x;
     const int bidx = blockIdx.x;
     const int bidy = blockIdx.y;
 
-    const float * x_block = (const float *) ((const char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1);
+    // byte strides of the input to move one channel / one time step, per layout
+    const int chan_nb = channels_major ? src0_nb0 : src0_nb1;
+    const int time_nb = channels_major ? src0_nb1 : src0_nb0;
+
+    const float * x_block = (const float *) ((const char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * chan_nb);
     const float * w_block = (const float *) ((const char *) src1 + bidy * split_d_inner * src1_nb1);
     float *       y_block = (float *) ((char *) dst + bidx * dst_nb2 + bidy * split_d_inner * dst_nb0);
 
-    const int stride_x = src0_nb1 / sizeof(float);
-    const int stride_w = src1_nb1 / sizeof(float);
-    const int stride_y = dst_nb1 / sizeof(float);
+    const int stride_xc = chan_nb / sizeof(float);   // stride to the next channel
+    const int stride_xt = time_nb / sizeof(float);   // stride to the next time step
+    const int stride_w  = src1_nb1 / sizeof(float);
+    const int stride_y  = dst_nb1 / sizeof(float);
 
     float x[d_conv] = { 0.0f };
     float w[d_conv] = { 0.0f };
@@ -42,10 +50,10 @@ static __global__ void ssm_conv_f32(const float * src0_ptr, const float * src1_p
 
         if (i == 0) {
             for (size_t j = 0; j < d_conv; j++) {
-                x[j] = x_block[tid * stride_x + j];
+                x[j] = x_block[tid * stride_xc + j * stride_xt];
             }
         } else {
-            x[(i - 1) % d_conv] = x_block[tid * stride_x + i + d_conv - 1];
+            x[(i - 1) % d_conv] = x_block[tid * stride_xc + (i + d_conv - 1) * stride_xt];
         }
 
 #pragma unroll
@@ -57,7 +65,7 @@ static __global__ void ssm_conv_f32(const float * src0_ptr, const float * src1_p
     }
 }
 
-template <bool apply_silu, size_t split_d_inner, size_t d_conv, int64_t split_n_t>
+template <bool apply_silu, bool channels_major, size_t split_d_inner, size_t d_conv, int64_t split_n_t>
 static __global__ void ssm_conv_long_token_f32(const float * __restrict__ src0, const float * __restrict__ src1,
                                                const float * __restrict__ bias,
                                                const int src0_nb0, const int src0_nb1, const int src0_nb2,
@@ -68,15 +76,19 @@ static __global__ void ssm_conv_long_token_f32(const float * __restrict__ src0,
     const int bidy = blockIdx.y;
     const int bidz = blockIdx.z;
 
-    const float * x_block = (const float *) ((const char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * src0_nb1 +
-                                             bidz * split_n_t * src0_nb0);
+    const int chan_nb = channels_major ? src0_nb0 : src0_nb1;
+    const int time_nb = channels_major ? src0_nb1 : src0_nb0;
+
+    const float * x_block = (const float *) ((const char *) src0 + bidx * src0_nb2 + bidy * split_d_inner * chan_nb +
+                                             bidz * split_n_t * time_nb);
     const float * w_block = (const float *) ((const char *) src1 + bidy * split_d_inner * src1_nb1);
     float *       y_block =
         (float *) ((char *) dst + bidx * dst_nb2 + bidz * split_n_t * dst_nb1 + bidy * split_d_inner * dst_nb0);
 
-    const int stride_x = src0_nb1 / sizeof(float);
-    const int stride_w = src1_nb1 / sizeof(float);
-    const int stride_y = dst_nb1 / sizeof(float);
+    const int stride_xc = chan_nb / sizeof(float);   // stride to the next channel
+    const int stride_xt = time_nb / sizeof(float);   // stride to the next time step
+    const int stride_w  = src1_nb1 / sizeof(float);
+    const int stride_y  = dst_nb1 / sizeof(float);
 
     const int64_t local_n_t = min(split_n_t, n_t - bidz * split_n_t);
     const int     n_cols    = d_conv - 1 + split_n_t;
@@ -84,20 +96,31 @@ static __global__ void ssm_conv_long_token_f32(const float * __restrict__ src0,
     extern __shared__ float smem[];
 
     constexpr int load_cols   = d_conv - 1 + split_n_t;
-    constexpr int total_elems = split_d_inner * load_cols;
-    int row = tid / load_cols;
-    int col = tid % load_cols;
+    if (channels_major) {
+        // channels are contiguous: each thread streams its own channel over the
+        // time tile, so consecutive threads read consecutive addresses (coalesced).
 #pragma unroll
-    for (int idx = 0; idx < total_elems; idx += split_d_inner) {
-        if (row < (int)split_d_inner) {
-            smem[row * n_cols + col] = x_block[row * stride_x + col];
+        for (int col = 0; col < load_cols; col++) {
+            smem[tid * n_cols + col] = x_block[tid * stride_xc + col * stride_xt];
         }
+    } else {
+        // time is contiguous: distribute the [channel, time] tile across threads
+        // so that consecutive threads read consecutive (time) addresses (coalesced).
+        constexpr int total_elems = split_d_inner * load_cols;
+        int row = tid / load_cols;
+        int col = tid % load_cols;
+#pragma unroll
+        for (int idx = 0; idx < total_elems; idx += split_d_inner) {
+            if (row < (int)split_d_inner) {
+                smem[row * n_cols + col] = x_block[row * stride_xc + col * stride_xt];
+            }
 
-        col += split_d_inner;
-        row += col / load_cols;
-        col  = col % load_cols;
-        if (idx >= total_elems - tid - split_d_inner) {
-            break;
+            col += split_d_inner;
+            row += col / load_cols;
+            col  = col % load_cols;
+            if (idx >= total_elems - tid - split_d_inner) {
+                break;
+            }
         }
     }
     __syncthreads();
@@ -123,7 +146,7 @@ static __global__ void ssm_conv_long_token_f32(const float * __restrict__ src0,
     }
 }
 
-template <bool apply_silu>
+template <bool apply_silu, bool channels_major>
 static void ssm_conv_f32_cuda(const float * src0, const float * src1, const float * bias, const int src0_nb0, const int src0_nb1,
                               const int src0_nb2, const int src1_nb1, float * dst, const int dst_nb0, const int dst_nb1,
                               const int dst_nb2, const int64_t nc, const int64_t nr, const int64_t n_t,
@@ -136,13 +159,13 @@ static void ssm_conv_f32_cuda(const float * src0, const float * src1, const floa
         if (n_t <= 32) {
             const dim3 blocks(n_s, (nr + threads - 1) / threads, 1);
             const ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(blocks, threads, 0, stream);
-            ggml_cuda_kernel_launch(ssm_conv_f32<apply_silu, threads, kNC>, launch_params, src0, src1, bias, src0_nb0, src0_nb1,
+            ggml_cuda_kernel_launch(ssm_conv_f32<apply_silu, channels_major, threads, kNC>, launch_params, src0, src1, bias, src0_nb0, src0_nb1,
                                                                         src0_nb2, src1_nb1, dst, dst_nb0, dst_nb1, dst_nb2, n_t);
         } else {
             const int64_t split_n_t = 32;
             dim3          blocks(n_s, (nr + threads - 1) / threads, (n_t + split_n_t - 1) / split_n_t);
             const size_t  smem_size = threads * (kNC - 1 + split_n_t) * sizeof(float);
-            ssm_conv_long_token_f32<apply_silu, threads, kNC, split_n_t><<<blocks, threads, smem_size, stream>>>(
+            ssm_conv_long_token_f32<apply_silu, channels_major, threads, kNC, split_n_t><<<blocks, threads, smem_size, stream>>>(
                 src0, src1, bias, src0_nb0, src0_nb1, src0_nb2, src1_nb1, dst, dst_nb0, dst_nb1, dst_nb2, n_t);
         }
     };
@@ -172,13 +195,15 @@ void ggml_cuda_op_ssm_conv(ggml_backend_cuda_context & ctx, ggml_tensor * dst, g
     // When fusing, write to silu_dst (the node downstream references).
     const struct ggml_tensor * out = fuse_silu ? silu_dst : dst;
 
-    const int64_t nc  = src1->ne[0];                // d_conv
-    const int64_t nr  = src0->ne[1];                // d_inner
-    const int64_t n_t = out->ne[1];                 // tokens per sequence
-    const int64_t n_s = out->ne[2];                 // number of sequences in the batch
+    const bool channels_major = ggml_ssm_conv_get_layout(dst) == GGML_SSM_CONV_LAYOUT_CHANNELS_MAJOR;
+
+    const int64_t nc  = src1->ne[0];                                // d_conv
+    const int64_t nr  = channels_major ? src0->ne[0] : src0->ne[1]; // d_inner
+    const int64_t n_t = out->ne[1];                                 // tokens per sequence
+    const int64_t n_s = out->ne[2];                                 // number of sequences in the batch
 
     GGML_ASSERT(out->ne[0] == nr);
-    GGML_ASSERT(src0->nb[0] == sizeof(float));
+    GGML_ASSERT(src0->nb[0] == sizeof(float));      // input is contiguous (either layout)
     GGML_ASSERT(src1->nb[0] == sizeof(float));
     GGML_ASSERT(src0->nb[1] == src0->ne[0] * sizeof(float));
 
@@ -197,10 +222,20 @@ void ggml_cuda_op_ssm_conv(ggml_backend_cuda_context & ctx, ggml_tensor * dst, g
     }
 
     if (fuse_silu) {
-        ssm_conv_f32_cuda<true>(src0_d, src1_d, bias_d, src0->nb[0], src0->nb[1], src0->nb[2], src1->nb[1], dst_d, out->nb[0], out->nb[1],
-                          out->nb[2], nc, nr, n_t, n_s, stream);
+        if (channels_major) {
+            ssm_conv_f32_cuda<true, true>(src0_d, src1_d, bias_d, src0->nb[0], src0->nb[1], src0->nb[2], src1->nb[1], dst_d, out->nb[0], out->nb[1],
+                              out->nb[2], nc, nr, n_t, n_s, stream);
+        } else {
+            ssm_conv_f32_cuda<true, false>(src0_d, src1_d, bias_d, src0->nb[0], src0->nb[1], src0->nb[2], src1->nb[1], dst_d, out->nb[0], out->nb[1],
+                              out->nb[2], nc, nr, n_t, n_s, stream);
+        }
     } else {
-        ssm_conv_f32_cuda<false>(src0_d, src1_d, bias_d, src0->nb[0], src0->nb[1], src0->nb[2], src1->nb[1], dst_d, out->nb[0], out->nb[1],
-                          out->nb[2], nc, nr, n_t, n_s, stream);
+        if (channels_major) {
+            ssm_conv_f32_cuda<false, true>(src0_d, src1_d, bias_d, src0->nb[0], src0->nb[1], src0->nb[2], src1->nb[1], dst_d, out->nb[0], out->nb[1],
+                              out->nb[2], nc, nr, n_t, n_s, stream);
+        } else {
+            ssm_conv_f32_cuda<false, false>(src0_d, src1_d, bias_d, src0->nb[0], src0->nb[1], src0->nb[2], src1->nb[1], dst_d, out->nb[0], out->nb[1],
+                              out->nb[2], nc, nr, n_t, n_s, stream);
+        }
     }
 }
diff --git a/ggml/src/ggml-hexagon/ggml-hexagon.cpp b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
index 76c71d7ee70e..3285e0933347 100644
--- a/ggml/src/ggml-hexagon/ggml-hexagon.cpp
+++ b/ggml/src/ggml-hexagon/ggml-hexagon.cpp
@@ -3250,6 +3250,11 @@ static bool ggml_hexagon_supported_ssm_conv(const struct ggml_hexagon_session *
     const struct ggml_tensor * src1 = op->src[1];
     const struct ggml_tensor * dst  = op;
 
+    // the channels-major input layout is only implemented on CPU/CUDA
+    if (ggml_ssm_conv_get_layout(op) != GGML_SSM_CONV_LAYOUT_TIME_MAJOR) {
+        return false;
+    }
+
     // Only support FP32 for now
     if (src0->type != GGML_TYPE_F32 || src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
         return false;
diff --git a/ggml/src/ggml-hip/CMakeLists.txt b/ggml/src/ggml-hip/CMakeLists.txt
index 7121193f1c84..5eb581b66b25 100644
--- a/ggml/src/ggml-hip/CMakeLists.txt
+++ b/ggml/src/ggml-hip/CMakeLists.txt
@@ -159,4 +159,20 @@ endif()
 
 target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas)
 
+# Per-op roofline profiling via rocprofiler-sdk. The library is loaded at runtime with
+# dlopen and its symbols resolved with dlsym, so only its headers and libdl are needed
+# to build; it is not linked (its constructors abort when linked into an early-loaded
+# shared object).
+if (GGML_HIP_ROOFLINE)
+    find_path(ROCPROFILER_SDK_INCLUDE_DIR rocprofiler-sdk/rocprofiler.h)
+    if (NOT ROCPROFILER_SDK_INCLUDE_DIR)
+        message(FATAL_ERROR "GGML_HIP_ROOFLINE requires rocprofiler-sdk headers (set CMAKE_PREFIX_PATH to the ROCm root)")
+    endif()
+    target_sources(ggml-hip PRIVATE ../ggml-cuda/ggml-cuda-roofline.cpp)
+    set_source_files_properties(../ggml-cuda/ggml-cuda-roofline.cpp PROPERTIES LANGUAGE CXX)
+    target_include_directories(ggml-hip PRIVATE ${ROCPROFILER_SDK_INCLUDE_DIR})
+    target_compile_definitions(ggml-hip PRIVATE GGML_HIP_ROOFLINE)
+    target_link_libraries(ggml-hip PRIVATE ${CMAKE_DL_LIBS})
+endif()
+
 target_compile_options(ggml-hip PRIVATE "$<$<COMPILE_LANGUAGE:HIP>:-ffast-math;-fno-finite-math-only>")
diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m
index 5d29250f654b..aa881fd90b25 100644
--- a/ggml/src/ggml-metal/ggml-metal-device.m
+++ b/ggml/src/ggml-metal/ggml-metal-device.m
@@ -1265,6 +1265,8 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
             }
             return has_simdgroup_mm; // TODO: over-restricted for vec-kernels
         case GGML_OP_SSM_CONV:
+            // the channels-major input layout is only implemented on CPU/CUDA
+            return has_simdgroup_reduction && ggml_ssm_conv_get_layout(op) == GGML_SSM_CONV_LAYOUT_TIME_MAJOR;
         case GGML_OP_SSM_SCAN:
             return has_simdgroup_reduction;
         case GGML_OP_RWKV_WKV6:
diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 5c96b9a9f6de..8d023a201456 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -6645,7 +6645,8 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
                    (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32) ||
                    (op->src[0]->type == GGML_TYPE_F16 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
         case GGML_OP_SSM_CONV:
-            return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32);
+            // the channels-major input layout is only implemented on CPU/CUDA
+            return (op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32 && op->type == GGML_TYPE_F32 && ggml_ssm_conv_get_layout(op) == GGML_SSM_CONV_LAYOUT_TIME_MAJOR);
         case GGML_OP_GATED_DELTA_NET:
             {
                 // Match the Vulkan backend: only F32 -> F32, S_v in {16, 32, 64, 128}.
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 5226fb184a98..7bf03446bc6a 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -5924,9 +5924,11 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons
         case GGML_OP_GATED_DELTA_NET:
             return true;
         case GGML_OP_SSM_CONV:
+            // the channels-major input layout is only implemented on CPU/CUDA
             return op->type == GGML_TYPE_F32 &&
                    op->src[0]->type == GGML_TYPE_F32 &&
-                   op->src[1]->type == GGML_TYPE_F32;
+                   op->src[1]->type == GGML_TYPE_F32 &&
+                   ggml_ssm_conv_get_layout(op) == GGML_SSM_CONV_LAYOUT_TIME_MAJOR;
         case GGML_OP_ROLL:
             return op->type == GGML_TYPE_F32;
         case GGML_OP_ARANGE:
diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
index a483d22c1a26..a4fe4dbb7f20 100644
--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
@@ -17683,7 +17683,8 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
                 return true;
             }
         case GGML_OP_SSM_CONV:
-            return op->src[0]->type == GGML_TYPE_F32;
+            // the channels-major input layout is only implemented on CPU/CUDA
+            return op->src[0]->type == GGML_TYPE_F32 && ggml_ssm_conv_get_layout(op) == GGML_SSM_CONV_LAYOUT_TIME_MAJOR;
         case GGML_OP_CONV_TRANSPOSE_1D:
             return op->src[0]->type == GGML_TYPE_F32 && op->src[1]->type == GGML_TYPE_F32;
         case GGML_OP_COL2IM_1D:
diff --git a/ggml/src/ggml-webgpu/ggml-webgpu.cpp b/ggml/src/ggml-webgpu/ggml-webgpu.cpp
index 29025e9ba4e3..92b54e361311 100644
--- a/ggml/src/ggml-webgpu/ggml-webgpu.cpp
+++ b/ggml/src/ggml-webgpu/ggml-webgpu.cpp
@@ -4354,7 +4354,8 @@ static bool ggml_backend_webgpu_device_supports_op(ggml_backend_dev_t dev, const
                           (src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16);
             break;
         case GGML_OP_SSM_CONV:
-            supports_op = op->type == GGML_TYPE_F32;
+            // the channels-major input layout is only implemented on CPU/CUDA
+            supports_op = op->type == GGML_TYPE_F32 && ggml_ssm_conv_get_layout(op) == GGML_SSM_CONV_LAYOUT_TIME_MAJOR;
             break;
         case GGML_OP_SSM_SCAN:
             supports_op = op->type == GGML_TYPE_F32 &&
diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c
index de0321d9ffd9..241695a28b96 100644
--- a/ggml/src/ggml.c
+++ b/ggml/src/ggml.c
@@ -5531,25 +5531,38 @@ struct ggml_tensor * ggml_flash_attn_back(
 
 // ggml_ssm_conv
 
-struct ggml_tensor * ggml_ssm_conv(
+// The input (sx) memory layout is selected by `layout` and stored in op_params[0].
+// The output is [d_inner, n_t, n_s] in both cases.
+static struct ggml_tensor * ggml_ssm_conv_impl(
         struct ggml_context * ctx,
         struct ggml_tensor  * sx,
-        struct ggml_tensor  * c) {
+        struct ggml_tensor  * c,
+        enum ggml_ssm_conv_layout layout) {
     GGML_ASSERT(ggml_is_3d(sx));
     GGML_ASSERT(ggml_is_matrix(c));
 
     const int64_t d_conv  = c->ne[0];
     const int64_t d_inner = c->ne[1];
-    const int64_t n_t     = sx->ne[0] - d_conv + 1; // tokens per sequence
     const int64_t n_s     = sx->ne[2];
 
+    // the time dimension index within sx depends on the layout
+    const int64_t time_ne = layout == GGML_SSM_CONV_LAYOUT_TIME_MAJOR ? sx->ne[0] : sx->ne[1];
+    const int64_t n_t     = time_ne - d_conv + 1; // tokens per sequence
+
     // TODO: maybe support other strides than 1?
-    GGML_ASSERT(sx->ne[0] == d_conv - 1 + n_t);
-    GGML_ASSERT(sx->ne[1] == d_inner);
+    if (layout == GGML_SSM_CONV_LAYOUT_TIME_MAJOR) {
+        GGML_ASSERT(sx->ne[0] == d_conv - 1 + n_t);
+        GGML_ASSERT(sx->ne[1] == d_inner);
+    } else {
+        GGML_ASSERT(sx->ne[0] == d_inner);
+        GGML_ASSERT(sx->ne[1] == d_conv - 1 + n_t);
+    }
     GGML_ASSERT(n_t >= 0);
 
     struct ggml_tensor * result = ggml_new_tensor_3d(ctx, GGML_TYPE_F32, d_inner, n_t, n_s);
 
+    ggml_set_op_params_i32(result, 0, (int32_t) layout);
+
     result->op     = GGML_OP_SSM_CONV;
     result->src[0] = sx;
     result->src[1] = c;
@@ -5557,6 +5570,25 @@ struct ggml_tensor * ggml_ssm_conv(
     return result;
 }
 
+struct ggml_tensor * ggml_ssm_conv(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * sx,
+        struct ggml_tensor  * c) {
+    return ggml_ssm_conv_impl(ctx, sx, c, GGML_SSM_CONV_LAYOUT_TIME_MAJOR);
+}
+
+struct ggml_tensor * ggml_ssm_conv_channels_major(
+        struct ggml_context * ctx,
+        struct ggml_tensor  * sx,
+        struct ggml_tensor  * c) {
+    return ggml_ssm_conv_impl(ctx, sx, c, GGML_SSM_CONV_LAYOUT_CHANNELS_MAJOR);
+}
+
+enum ggml_ssm_conv_layout ggml_ssm_conv_get_layout(const struct ggml_tensor * op) {
+    GGML_ASSERT(op->op == GGML_OP_SSM_CONV);
+    return (enum ggml_ssm_conv_layout) ggml_get_op_params_i32(op, 0);
+}
+
 // ggml_ssm_scan
 
 struct ggml_tensor * ggml_ssm_scan(
diff --git a/scripts/hip/gcn-cdna-vgpr-check.py b/scripts/hip/gcn-cdna-vgpr-check.py
index bbbce52ef39f..297f3bede8e2 100644
--- a/scripts/hip/gcn-cdna-vgpr-check.py
+++ b/scripts/hip/gcn-cdna-vgpr-check.py
@@ -143,7 +143,18 @@ def main():
         '_ZL9mul_mat_qIL9ggml_type40ELi112ELb0EEvPKcPKiS4_S4_PfS5_iiiiiiiiiiiiiiiii',
         '_ZL9mul_mat_qIL9ggml_type40ELi112ELb1EEvPKcPKiS4_S4_PfS5_iiiiiiiiiiiiiiiii',
         '_ZL9mul_mat_qIL9ggml_type40ELi128ELb0EEvPKcPKiS4_S4_PfS5_iiiiiiiiiiiiiiiii',
-        '_ZL9mul_mat_qIL9ggml_type40ELi128ELb1EEvPKcPKiS4_S4_PfS5_iiiiiiiiiiiiiiiii'
+        '_ZL9mul_mat_qIL9ggml_type40ELi128ELb1EEvPKcPKiS4_S4_PfS5_iiiiiiiiiiiiiiiii',
+        # TEMPORARY (revisit on next upstream sync): the mmq rewrite repacked
+        # kernel dims into HIP_vector_type<uint,3>, changing the mangled name so
+        # the pre-existing allowlist entries no longer match. These Q2_K variants
+        # and the rwkv_wkv_f32 recurrent kernel spill on gfx908 exactly like their
+        # already-ignored siblings. Upstream master is itself red on this check
+        # (ggml-org/llama.cpp#21020 "HIP quality check failing again"); drop these
+        # once upstream refreshes its own ignore list.
+        '_ZL12rwkv_wkv_f32ILi128EEviiiiPKfS1_S1_S1_S1_S1_Pf',
+        '_ZL9mul_mat_qIL9ggml_type10ELi48ELb0EEvPKcPKiS4_S4_PfS5_15HIP_vector_typeIjLj3EEiiiiiS7_S7_iiiS7_S7_iiiS7_',
+        '_ZL9mul_mat_qIL9ggml_type10ELi64ELb0EEvPKcPKiS4_S4_PfS5_15HIP_vector_typeIjLj3EEiiiiiS7_S7_iiiS7_S7_iiiS7_',
+        '_ZL9mul_mat_qIL9ggml_type10ELi64ELb1EEvPKcPKiS4_S4_PfS5_15HIP_vector_typeIjLj3EEiiiiiS7_S7_iiiS7_S7_iiiS7_'
     }
 
     functions = parse_log_file(log_file)
diff --git a/src/models/delta-net-base.cpp b/src/models/delta-net-base.cpp
index ad6612647736..9fab0572972c 100644
--- a/src/models/delta-net-base.cpp
+++ b/src/models/delta-net-base.cpp
@@ -463,28 +463,33 @@ ggml_tensor * llm_build_delta_net_base::build_conv_state(
     ggml_tensor * conv_states = build_rs(inp, conv_states_all, hparams.n_embd_r(), n_seqs);
     cb(conv_states, "conv_states", il);
 
-    conv_states = ggml_reshape_3d(ctx0, conv_states, conv_kernel_size - 1, conv_channels, n_seqs);
+    // channels-major layout: [conv_channels, time, n_seqs] (channels contiguous).
+    // qkv_mixed already arrives channels-major as [conv_channels, n_tokens, n_seqs],
+    // so no transpose is needed: we prepend the recurrent conv state along the time
+    // axis (dim 1) and feed ggml_ssm_conv_channels_major directly. This avoids the
+    // physical transpose (a CONT kernel) that the time-major path required.
+    conv_states = ggml_reshape_3d(ctx0, conv_states, conv_channels, conv_kernel_size - 1, n_seqs);
     cb(conv_states, "conv_states_reshaped", il);
 
-    qkv_mixed = ggml_transpose(ctx0, qkv_mixed);
-    cb(qkv_mixed, "qkv_mixed_transposed", il);
-
-    ggml_tensor * conv_input = ggml_concat(ctx0, conv_states, qkv_mixed, 0);
+    ggml_tensor * conv_input = ggml_concat(ctx0, conv_states, qkv_mixed, 1);
     cb(conv_input, "conv_input", il);
 
     const int64_t row_count = (conv_kernel_size - 1) * conv_channels;
 
     const size_t row_size  = ggml_row_size(conv_states_all->type, row_count);
 
+    // number of time steps in conv_input (state + new tokens)
+    const int64_t n_time = conv_input->ne[1];
+
     if (cparams.n_rs_seq == 0) {
-        const int64_t s_idx  = conv_input->ne[0] - conv_states->ne[0];
+        const int64_t s_idx  = n_time - (conv_kernel_size - 1);
         const int64_t s_slot = 0;
 
         ggml_tensor * conv_state_last =
             ggml_view_3d(ctx0, conv_input,
-                    conv_kernel_size - 1, conv_channels, n_seqs,
+                    conv_channels, conv_kernel_size - 1, n_seqs,
                     conv_input->nb[1], conv_input->nb[2],
-                    ggml_row_size(conv_input->type, s_idx));
+                    s_idx * conv_input->nb[1]);
         cb(conv_state_last, "conv_state_last", il);
 
         ggml_tensor * conv_state_update =
@@ -502,14 +507,14 @@ ggml_tensor * llm_build_delta_net_base::build_conv_state(
         const int64_t K = (int64_t) cparams.n_rs_seq + 1;
 
         for (int64_t t = 1; t <= K; ++t) {
-            const int64_t s_idx  = std::max<int64_t>(0, conv_input->ne[0] - conv_states->ne[0] - K + t);
+            const int64_t s_idx  = std::max<int64_t>(0, n_time - (conv_kernel_size - 1) - K + t);
             const int64_t s_slot = K - t;
 
             ggml_tensor * conv_state_last =
                 ggml_view_3d(ctx0, conv_input,
-                        conv_kernel_size - 1, conv_channels, n_seqs,
+                        conv_channels, conv_kernel_size - 1, n_seqs,
                         conv_input->nb[1], conv_input->nb[2],
-                        ggml_row_size(conv_input->type, s_idx));
+                        s_idx * conv_input->nb[1]);
 
             ggml_tensor * conv_state_update =
                 ggml_view_2d(ctx0,
diff --git a/src/models/qwen35.cpp b/src/models/qwen35.cpp
index d8ffe43ae76c..5a5092372efa 100644
--- a/src/models/qwen35.cpp
+++ b/src/models/qwen35.cpp
@@ -391,7 +391,7 @@ ggml_tensor * llama_model_qwen35::graph::build_layer_attn_linear(
     state = ggml_reshape_4d(ctx0, state, head_v_dim, head_v_dim, num_v_heads, n_seqs);
     cb(state, "state_predelta", il);
 
-    ggml_tensor * conv_output_proper = ggml_ssm_conv(ctx0, conv_input, conv_kernel);
+    ggml_tensor * conv_output_proper = ggml_ssm_conv_channels_major(ctx0, conv_input, conv_kernel);
     cb(conv_output_proper, "conv_output_raw", il);
 
     ggml_tensor * conv_output_silu = ggml_silu(ctx0, conv_output_proper);
diff --git a/src/models/qwen35moe.cpp b/src/models/qwen35moe.cpp
index 7b0876cbb04b..27a79b31cfca 100644
--- a/src/models/qwen35moe.cpp
+++ b/src/models/qwen35moe.cpp
@@ -415,7 +415,7 @@ ggml_tensor * llama_model_qwen35moe::graph::build_layer_attn_linear(
     state = ggml_reshape_4d(ctx0, state, head_v_dim, head_v_dim, num_v_heads, n_seqs);
     cb(state, "state_predelta", il);
 
-    ggml_tensor * conv_output_proper = ggml_ssm_conv(ctx0, conv_input, conv_kernel);
+    ggml_tensor * conv_output_proper = ggml_ssm_conv_channels_major(ctx0, conv_input, conv_kernel);
     cb(conv_output_proper, "conv_output_raw", il);
 
     ggml_tensor * conv_output_silu = ggml_silu(ctx0, conv_output_proper);
diff --git a/src/models/qwen3next.cpp b/src/models/qwen3next.cpp
index 09b66423d5a8..7abadad7dcb1 100644
--- a/src/models/qwen3next.cpp
+++ b/src/models/qwen3next.cpp
@@ -442,7 +442,7 @@ ggml_tensor * llama_model_qwen3next::graph::build_layer_attn_linear(
     state = ggml_reshape_4d(ctx0, state, head_v_dim, head_v_dim, num_v_heads, n_seqs);
     cb(state, "state_predelta", il);
 
-    ggml_tensor * conv_output_proper = ggml_ssm_conv(ctx0, conv_input, conv_kernel);
+    ggml_tensor * conv_output_proper = ggml_ssm_conv_channels_major(ctx0, conv_input, conv_kernel);
     cb(conv_output_proper, "conv_output_raw", il);
 
     ggml_tensor * conv_output_silu = ggml_silu(ctx0, conv_output_proper);
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 1fae3f5176c3..4edaba125079 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -3754,20 +3754,23 @@ struct test_ssm_conv : public test_case {
     const ggml_type type;
     const std::array<int64_t, 4> ne_a;
     const std::array<int64_t, 4> ne_b;
+    const bool channels_major;  // if true, ne_a is [d_inner, d_conv-1+n_t, n_s, 1]
 
     std::string vars() override {
-        return VARS_TO_STR3(type, ne_a, ne_b);
+        return VARS_TO_STR4(type, ne_a, ne_b, channels_major);
     }
 
     test_ssm_conv(ggml_type type = GGML_TYPE_F32,
             std::array<int64_t, 4> ne_a = {10, 10, 10, 1},
-            std::array<int64_t, 4> ne_b = {3, 3, 1, 1})
-        : type(type), ne_a(ne_a), ne_b(ne_b) {}
+            std::array<int64_t, 4> ne_b = {3, 3, 1, 1},
+            bool channels_major = false)
+        : type(type), ne_a(ne_a), ne_b(ne_b), channels_major(channels_major) {}
 
     ggml_tensor * build_graph(ggml_context * ctx) override {
         ggml_tensor * a   = ggml_new_tensor(ctx, type, 4, ne_a.data());
         ggml_tensor * b   = ggml_new_tensor(ctx, type, 4, ne_b.data());
-        ggml_tensor * out = ggml_ssm_conv(ctx, a, b);
+        ggml_tensor * out = channels_major ? ggml_ssm_conv_channels_major(ctx, a, b)
+                                           : ggml_ssm_conv(ctx, a, b);
         return out;
     }
 };
@@ -8480,6 +8483,14 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
             // long token (n_t > 32, exercises the long_token kernel path)
             test_cases.emplace_back(new test_ssm_conv(GGML_TYPE_F32, {d_conv - 1 + 64, d_inner, 1, 1}, {d_conv, d_inner, 1, 1}));
             test_cases.emplace_back(new test_ssm_conv(GGML_TYPE_F32, {d_conv - 1 + 64, d_inner, 4, 1}, {d_conv, d_inner, 1, 1}));
+
+            // channels-major input: ne_a is [d_inner, d_conv-1+n_t, n_s, 1]
+            test_cases.emplace_back(new test_ssm_conv(GGML_TYPE_F32, {d_inner, d_conv, 1, 1}, {d_conv, d_inner, 1, 1}, true));
+            test_cases.emplace_back(new test_ssm_conv(GGML_TYPE_F32, {d_inner, 2 * d_conv, 1, 1}, {d_conv, d_inner, 1, 1}, true));
+            test_cases.emplace_back(new test_ssm_conv(GGML_TYPE_F32, {d_inner, d_conv, 4, 1}, {d_conv, d_inner, 1, 1}, true));
+            // long token (n_t > 32, exercises the long_token kernel path)
+            test_cases.emplace_back(new test_ssm_conv(GGML_TYPE_F32, {d_inner, d_conv - 1 + 64, 1, 1}, {d_conv, d_inner, 1, 1}, true));
+            test_cases.emplace_back(new test_ssm_conv(GGML_TYPE_F32, {d_inner, d_conv - 1 + 64, 4, 1}, {d_conv, d_inner, 1, 1}, true));
         }
     }
 
@@ -8555,6 +8566,14 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_Q8_0, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1}));
     test_cases.emplace_back(new test_mul_mat(GGML_TYPE_MXFP4, GGML_TYPE_F32, 2880, 32, 2880, {1, 1}, {1, 1}));
 
+    // single-row f32 weight (m == 1) against a wide activation (n): the second column count sweeps
+    // across MMVF_MAX_BATCH_SIZE (8) so both the direct mmvf path (n <= 8) and the CUDA/HIP
+    // operand-swap GEMV path (n > 8, which would otherwise fall back to an f32 cuBLAS GEMM) are
+    // exercised and checked against the CPU reference. Regression guard for the mul_mat_vec_f swap.
+    for (int64_t n : {1, 7, 8, 9, 16, 128, 512}) {
+        test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F32, GGML_TYPE_F32, 1, n, 2048, {1, 1}, {1, 1}));
+    }
+
 
 #if 0
     {
diff --git a/tools/llama-bench/CMakeLists.txt b/tools/llama-bench/CMakeLists.txt
index b1c35ee88a5f..e5119c24380d 100644
--- a/tools/llama-bench/CMakeLists.txt
+++ b/tools/llama-bench/CMakeLists.txt
@@ -8,6 +8,13 @@ set_target_properties(${TARGET} PROPERTIES WINDOWS_EXPORT_ALL_SYMBOLS ON)
 target_include_directories(${TARGET} PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
 target_link_libraries(${TARGET} PUBLIC llama-common llama ${CMAKE_THREAD_LIBS_INIT})
 
+# Compile in the post-warmup roofline reset only when the HIP profiler is built
+# (ggml-cuda-roofline.cpp, linked via ggml-hip). No-op in every other build.
+if (GGML_HIP_ROOFLINE)
+    target_compile_definitions(${TARGET} PRIVATE GGML_HIP_ROOFLINE)
+    target_include_directories(${TARGET} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../ggml/src/ggml-cuda)
+endif()
+
 if(LLAMA_TOOLS_INSTALL)
     install(TARGETS ${TARGET} LIBRARY)
 endif()
diff --git a/tools/llama-bench/llama-bench.cpp b/tools/llama-bench/llama-bench.cpp
index 0756893881ff..c96508854566 100644
--- a/tools/llama-bench/llama-bench.cpp
+++ b/tools/llama-bench/llama-bench.cpp
@@ -35,6 +35,13 @@
 #    include <windows.h>
 #endif
 
+// Optional per-op roofline profiler (HIP build with GGML_HIP_ROOFLINE). Used to drop
+// warmup from the report. Compiled out entirely otherwise. See
+// ggml/src/ggml-cuda/ggml-cuda-roofline.cpp.
+#ifdef GGML_HIP_ROOFLINE
+#include "ggml-cuda-roofline.h"
+#endif
+
 // utils
 static uint64_t get_time_ns() {
     using clock = std::chrono::high_resolution_clock;
@@ -2356,6 +2363,12 @@ int llama_bench(int argc, char ** argv) {
             }
         }
 
+        // Discard warmup from the per-op roofline profiler so the report covers only
+        // the measured runs below. Compiled out unless GGML_HIP_ROOFLINE is set.
+#ifdef GGML_HIP_ROOFLINE
+        ggml_cuda_roofline_reset();
+#endif
+
         for (int i = 0; i < params.reps; i++) {
             llama_memory_clear(llama_get_memory(ctx), false);
 
diff --git a/tools/mtmd/CMakeLists.txt b/tools/mtmd/CMakeLists.txt
index ea684d9f156d..6ef1fee20baf 100644
--- a/tools/mtmd/CMakeLists.txt
+++ b/tools/mtmd/CMakeLists.txt
@@ -134,6 +134,13 @@ if (LLAMA_BUILD_TOOLS)
     target_link_libraries  (${TARGET} PRIVATE llama-common mtmd Threads::Threads)
     target_compile_features(${TARGET} PRIVATE cxx_std_17)
 
+    # Compile in the post-warmup roofline reset only when the HIP profiler is built
+    # (ggml-cuda-roofline.cpp, linked via ggml-hip). No-op in every other build.
+    if (GGML_HIP_ROOFLINE)
+        target_compile_definitions(${TARGET} PRIVATE GGML_HIP_ROOFLINE)
+        target_include_directories(${TARGET} PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../../ggml/src/ggml-cuda)
+    endif()
+
     # mtmd-debug tool
     add_executable(llama-mtmd-debug debug/mtmd-debug.cpp)
     set_target_properties(llama-mtmd-debug PROPERTIES OUTPUT_NAME llama-mtmd-debug)
diff --git a/tools/mtmd/mtmd-cli.cpp b/tools/mtmd/mtmd-cli.cpp
index 8704ea79d7a0..f50a037a0b34 100644
--- a/tools/mtmd/mtmd-cli.cpp
+++ b/tools/mtmd/mtmd-cli.cpp
@@ -10,6 +10,13 @@
 #include "mtmd.h"
 #include "mtmd-helper.h"
 
+// Optional per-op roofline profiler (HIP build with GGML_HIP_ROOFLINE). Used to
+// discard warmup so the report covers only the measured prefill. Compiled out
+// otherwise. See ggml/src/ggml-cuda/ggml-cuda-roofline.cpp.
+#ifdef GGML_HIP_ROOFLINE
+#include "ggml-cuda-roofline.h"
+#endif
+
 #include <vector>
 #include <limits.h>
 #include <cinttypes>
@@ -453,6 +460,11 @@ int main(int argc, char ** argv) {
                 return 1; // error is already printed by libmtmd
             }
         }
+        // Discard warmup from the per-op roofline profiler so the report covers
+        // only the measured prefill below. Compiled out unless GGML_HIP_ROOFLINE.
+#ifdef GGML_HIP_ROOFLINE
+        ggml_cuda_roofline_reset();
+#endif
         if (eval_message(ctx, msg)) {
             return 1;
         }