Public/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-09-19 09:15:04 +02:00
Code Issues Packages Projects Releases Wiki Activity

Improves flash attn PDL barrier placement

Browse Source
This commit is contained in:
aendk
2026-05-15 14:13:29 +02:00
parent ac3365305d
commit a459f2fb32
2 changed files with 3 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/fattn-mma-f16.cuh
+1
View File
@@ -1559,6 +1559,7 @@ static __global__ void flash_attn_ext_f16(
const int32_t nb21, const int32_t nb22, const int64_t nb23,
const int32_t ne31, const int32_t ne32, const int32_t ne33,
const int32_t nb31, const int32_t nb32, const int64_t nb33) {
ggml_cuda_pdl_sync(); // just for safety, as this kernel is indirectly launched via ggml_cuda_kernel_launch
#if defined(FLASH_ATTN_AVAILABLE) && (defined(VOLTA_MMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || (defined(AMD_WMMA_AVAILABLE) && defined(RDNA4)) || defined(AMD_MFMA_AVAILABLE))
// Skip unused kernel variants for faster compilation:
ggml/src/ggml-cuda/fattn-tile.cuh
+2 -2
View File
@@ -870,8 +870,6 @@ static __global__ void flash_attn_tile(
constexpr int DKQp = (DKQ + 2*warp_size - 1) & ~(2*warp_size - 1); // DKQ padded to multiple of 2*warp_size.
constexpr int DVp = (DV + 2*warp_size - 1) & ~(2*warp_size - 1); // DV padded to multiple of 2*warp_size.
ggml_cuda_pdl_sync(); // needs to guard Q, K, V, mask, sinks, KV_max, dst, dst_meta data accesses. Conservatively placed, not optimal
// Q_tmp == SRAM buffer to hold Q data for the entire lifetime of the kernel.
// KV_tmp == SRAM buffer to hold fragments of K/V data while iterating over ne11.
// KV_tmp is padded to avoid memory conflicts for K (cpy_ne) and OOB accesses for V (DVp-DV).
@@ -896,6 +894,8 @@ static __global__ void flash_attn_tile(
}
float KQ_sum[cpw] = {0.0f};
ggml_cuda_pdl_sync(); // needs to guard Q, K, V, mask, sinks, KV_max, dst, dst_meta data accesses.
// Load Q data, convert to FP16 if fast:
#pragma unroll
for (int jc0 = 0; jc0 < cpw; ++jc0) {