CUDA: refactor swizzling code by JohannesGaessler · Pull Request #29612 · ggml-org/llama.cpp

JohannesGaessler · 2026-09-28T20:14:46Z

I refactored and generalized the CUDA swizzling code a bit with support for strides that are not a multiple of 128 bytes as well as Volta and AMD. Unfortunately the performance for those cases was worse than with padding so the corresponding code paths are disabled. Still, I think we should keep them since they will be potentially useful in the future due to the reduced VRAM requirements.

The functions load_ldmatrix_swizzled and load_ldmatrix_trans_swizzled could become replacements for the non-_swizzled variants in the future since they are a superset of the functionality if the stride is known at compile time. However, I am not yet 100% sure about this so I'm avoiding mass-changing the kernels that use them for now.

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: No

ynankani · 2026-09-30T06:44:01Z

+            : "l"(xs0 + offset_ij));
+#elif defined(VOLTA_MMA_AVAILABLE)
+#pragma unroll
+        for (int o = 0; o < 8; o += 4) {


the array bound should be below?

Suggested change

for (int o = 0; o < 8; o += 4) {

for (int o = 0; o < t.ne; o += 4) {

ynankani · 2026-09-30T07:01:06Z

looks like bool swz is no longer needed in the template declaration.

ynankani · 2026-09-30T07:02:46Z

same swz arg is not needed in this call

ynankani · 2026-09-30T07:03:06Z

same swz arg is not needed in this call

ynankani · 2026-09-30T07:03:29Z

same swz arg is not needed in this call

ynankani · 2026-09-30T07:03:40Z

same swz arg is not needed in this call

ynankani · 2026-09-30T07:04:10Z

same swz arg is not needed in this call

ynankani · 2026-10-02T06:14:46Z

@IMbackK would you like to take look at AMD side changes and approve?

IMbackK

Its fine from a hip perspective, It looks correct statically to me and passes tests when enabled on cdna.
I dident bother trying to determine performance or if the layout is ideal since it will remain disabled on hip for now

* CUDA: refactor swizzling code * fix templates/loop bounds (cherry picked from commit 2ca15f5)

ServeurpersoCom · 2026-10-09T07:29:38Z

Hey @JohannesGaessler, since this refactor the HIP quality check is red on master: flash_attn_ext_f16<320, 256, 2, 32> now uses 257 VGPRs on gfx908, one spill over the 256 limit. Would you rather recover that register in the kernel, or should I add it to the ignore list in gcn-cdna-vgpr-check.py next to the other big head sizes?

JohannesGaessler · 2026-10-09T07:40:03Z

I don't remember where it was but I think in some slop fork someone claimed that the original swizzling PR was problematic for compiler optimizations on AMD because of the cast to char * and back again. I was unable to measure any performance differences on my AMD GPUs. So I would say to ignore it.

ServeurpersoCom · 2026-10-09T07:51:14Z

I don't remember where it was but I think in some slop fork someone claimed that the original swizzling PR was problematic for compiler optimizations on AMD because of the cast to char * and back again. I was unable to measure any performance differences on my AMD GPUs. So I would say to ignore it.

Perfect, thanks! Small fix CI here then: #30212

…-10-10) Conflicts: - ggml-backend.cpp: upstream moved input copies to ggml_backend_sched_copy_input with host weights last (ggml-org#29943); the GGML_SCHED_BATCH_INPUTS fast path now runs in the first loop before falling back to it. - ggml-cuda/fattn-common.cuh: launch_fattn takes upstream async_kv_preload, then our parallel_blocks_fixed (GQA-6 tile launcher updated). - ggml-cuda/fattn-mma-f16.cuh: upstream swizzle refactor (ggml-org#29612); its helpers take ncols2 like the rest of the fork's config getters. - ggml-cuda/gated_delta_net.cu: keep the fork's kernel and launch geometry (its columns_per_block template argument is not upstream's cols_per_warp, ggml-org#30087). - ggml-cuda/mmq.cu: upstream per-tile src1 padding (ggml-org#29953); ggml_cuda_mul_mat_q_src1_nbytes and the MMQ input cache use the worst-case padding so producers and projections with other tiles can share the buffer. - ggml-cuda/mmq-vec-dot.cuh: fork vec_dot kernels pass GGML_PREC_Q8 to the config helpers (ggml-org#30168). - ggml-cuda/mmvf.cu/.cuh: keep both includes and the HIP declarations, upstream warp_size argument. - ggml-cuda/rope.cu: upstream grid-stride rms_norm+mul+rope (ggml-org#28175) with the fork's M-RoPE theta (nchannels instead of gridDim.y) and rope_store_cast. - ggml-cuda/top-k.cu: upstream shape-based selection (ggml-org#28713), RDNA4 tournament kept for few long rows. - src/llama-batch.*: mixed token/embd batches (ggml-org#29622) copy embd rows per entry; other batches keep pointing at the ext rows. - src/llama-context.cpp: expert copy callback also on the second MTP scheduler. - src/llama-graph.cpp: build_rs keeps the GET_ROWS of the ubatch states (GDN fusions, deferred state) and gathers the extra states from the tail like upstream's custom getter path (ggml-org#29856).

* CUDA: refactor swizzling code * fix templates/loop bounds

JohannesGaessler requested a review from a team as a code owner September 28, 2026 20:14

JohannesGaessler mentioned this pull request Sep 28, 2026

cuda, hip: GDN chunked kernel #29353

Merged

github-actions Bot added ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Sep 28, 2026

ynankani self-assigned this Sep 29, 2026

ynankani reviewed Sep 30, 2026

View reviewed changes