Gfx11

.github/workflows/code-style.yml

@@ -5,17 +5,21 @@ 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 }}
   cancel-in-progress: true
 
 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

.github/workflows/editorconfig.yml

@@ -5,17 +5,21 @@ 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 }}
   cancel-in-progress: true
 
 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

.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

.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

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)

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,

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};

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:

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;
             }

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() {

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);