* Removes __restrict__ from PDL kernel headers due to incompatibility with
PDL. Adds preprocessor directives based on arch in kernel body to add
__restrict__ to retain performance on older architectures.
* Simplifies new __restrict__ usage via macro
* Add hopper to PDL __restrict__ fix.
Co-authored-by: Oliver Simons <osimons@nvidia.com>
---------
Co-authored-by: Oliver Simons <osimons@nvidia.com>
* Adds initial PDL setup.
* Adds PDL barriers based on simple heuristic: place "sync" before first input pointer access, and "launch" after last write, e.g. to tensors like dst.
* Further optimization pass of the first half of kernels
* Optimized PDL barriers for the second batch of kernels
* Further refinements after rebase.
* Moves pdl logic to separate function, removes some whitespace
* Strips post-hoc PDL logic
* Adds stream capture PDL setup. Enrolls quantize_q8_1 to leverage pdl to
overlap execution with previous kernels
* Enrolls mul_mat_vec_q, rms_norm_f32 and k_bin_bcast (partly) into PDL
* Enrolls mmvf, rope, set-rows and topk kernels for gpt-oss into PDL
* Introduce ggml_cuda_kernel_launch, to abstract away cudaLaunchKernelEx,
to enable hip/musa compatibility
* Enrolls cpy_scalar_contiguous, k_get_rows_float and rms_norm_f32
* Enrolls flash_attn_combine_results
* Fix: Drops needless and broken check of CUDA arch for PDL. PDL either
works or is without effect.
* Enrolls flash-attention kernels to pdl
* Fix: inlines ggml_cuda_kernel_launch, and uses perfect forwarding for
kernels args. This fixes PDL.
* Perf: Enrolls k_bin_bcast variadic template invocation into PDL, via
and template alias and template expansion
* Enrolls all remaining kernels for qwen3-coder-next into PDL
* Remove all PDL LC calls to create a baseline
* Added LC according to internal guidance and tested kernel performance.
* Enrols missing qwen3-5 kernels passively into PDL.
* Kernel optimizations (LC signals) for qwen3.5
* Enrolls ssm-scan kernels into PDL
* Adds GGML_CUDA_PDL command line option to toggle PDL.
* Fix: Ada and lower compilation by guarding PDL calls correctly
* Cleanup: Removes commented out GGML_CUDA_PDL_LC
* Cleanup: Removes experimental comments
* Adds 90-virtual to build script so that Hopper GPUs can leverage PDL.
* Adds stricter checks to enable PDL, adds env-check to disable it, and removes now superfluous compile option to enable PDL.
* Fix: Correct PDL en/disablement based on device-side arch check. Host
side check is UB. Required moving from macros to inlined functions
* Fix: default-disable PDL. Enable by setting GGML_CUDA_ENABLE_PDL=1
* Enable PDL by default for Hopper+ devices
* Enrolls softcap_f32 and two flash_attn kernels into PDL.
* Improves flash attn PDL barrier placement
* Fix: Perf regression on ada; excludes ada and below from PDL launches
* Improves some sync barrier placements
* Drops superfluous constructor
* Adds #endif guard comments
* Reverts experimental change to top-k-moe.cu, which moved expensive allocations
in front of the PDL barrier. It did not have a meaningful impact.
* Exchanges GGML_CUDA_DISABLE_PDL with GGML_CUDA_PDL. IFF GGML_CUDA_PDL=0
PDL is disabled
* Revert "Drops superfluous constructor". Adds const to remaining
arguments
This reverts commit 12b1d250da0089ae02a9bb71bbb3fd6d70f6f2f1.
* Cleanup: Removes and fixes some comments and whitespace
* Clarifies comment of sync-barrier position
* Relocates and refactors PDL launch functions and accessories
* Adds error checking to the regular kernel launch path
* Drops "auto" in favor of "ggml_cuda_kernel_params"
* Adds "const" to ggml_cuda_kernel_launch_params
* [Whitespace] Adds final newline to common.cuh to make editorconfig CI job happy
Adds RDNA3 support to the CUDA mma FA kernel. To make the RDNA3 tensor cores work with the FP16 accumulation for VKQ the tiles they need to be 32 logical units long in direction of the attention head; for head sizes 80 and 112 that are not exactly divided by 32 the regular length of 16 with FP32 accumulation is used instead. The longer tiles also enable more efficient transposition for a warp size of 32 which is why it's also used for RDNA4. However, this scrambles the data layout of the accumulators along the attention head dimension. To prevent accidental misuse I added another entry to ggml_cuda_mma::data_layout.
I also tuned the kernel parameters for RDNA3, RDNA4, and CDNA1 in general, during which I discovered that the kernel can be made to work for head sizes up to 256 for CDNA. For RDNA3/4 I was not able to get better performance that the tile kernel for head sizes > 128.
* mimo-v2.5: add flash attention mma/tiles for for d_kq=192 d_v=128
* mimo-v2.5: follow (256, 256) fattn templates
* mimo-v2.5: cleanup comments
* mimo-v2.5: further comment cleanup
* mimo-v2.5: address PR feedback
fix GQA handling
check for other dangling 320/576 carveouts and mirror them for 192
Add to backend ops test so new paths are covered
* ggml-cuda: add flash-attn support for DKQ=320/DV=256 with ncols2=32 (GQA=32)
Adds MMA-f16 and tile kernel configs, dispatch logic, template instances,
and tile .cu file for Mistral Small 4 (head sizes 320/256), restricting to
ncols2=32 to support GQA ratio 32 only.
* Adding check to return BEST_FATTN_KERNEL_NONE in case GQA!=32
* Apply suggestions from code review
Address review comments
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Address review comments and making kernel config default to DQK=512, DV=512 instead of DQK=256,DV=256
* Fixed bug with sinks=1, with ncols=32, there are two warp-groups created but sinks index is same(0,...,15) for both the groups hence with sinks=1, output is not matching with CPU output. Added sink_base which will be base index for each warp_group (threadIdx.y / np)
* Apply suggestions from code review
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Update ggml/src/ggml-cuda/template-instances/generate_cu_files.py
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* flash attention support for head dimension 512 added
* FA D=512 - match 576 configs, limit ncols2, revert vec cap
* fix HIP tile kernel build for D=512
* fix HIP tile kernel occupancy for D=512 on AMD
* Apply suggestions from code review
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* fix tile FA compilation
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* mla : pass V as a view of K to the FA op
* cuda : adjust mla logic to new layout
* kv-cache : fix rope shift
* tests : remove comment
* cuda : fix reusable_cutoff
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* musa: fix build warnings
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
* fix warning: comparison of integers of different signs: 'const int' and 'unsigned int' [-Wsign-compare]
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
---------
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
* musa: fix build warning (unused parameter)
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
* musa: upgrade MUSA SDK version to rc4.0.1
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
* musa: use mudnn::Unary::IDENTITY op to accelerate D2D memory copy
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
* Update ggml/src/ggml-cuda/cpy.cu
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* musa: remove MUDNN_CHECK_GEN and use CUDA_CHECK_GEN instead in MUDNN_CHECK
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
---------
Signed-off-by: Xiaodong Ye <xiaodong.ye@mthreads.com>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
- Find out active blocks per SM using cudaOccupancyMaxActiveBlocksPerMultiprocessor API. Use this value to determine the optimal parallel_blocks value.
- Prefer vector flash attention kernels over MMA kernel for BS=1
Fixes Issue: #12182
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* CUDA: use mma PTX instructions for FlashAttention
* __shfl_sync workaround for movmatrix
* add __shfl_sync to HIP
Co-authored-by: Diego Devesa <slarengh@gmail.com>
Johannes Gäßler