ver4a/llama.cpp
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity

CUDA: fix FlashAttention on Turing (#13415)

Browse source
This commit is contained in:
Johannes Gäßler 2025-05-10 09:16:52 +02:00 • committed by GitHub
parent 7fef11766c
commit d8919424f1
No known key found for this signature in database
GPG key ID: B5690EEEBB952194
1 changed files with 1 additions and 1 deletions
Download patch file Download diff file Expand all files Collapse all files

2
ggml/src/ggml-cuda/fattn-mma-f16.cuh
View file

@ -546,7 +546,7 @@ static __device__ __forceinline__ void flash_attn_ext_f16_iter(
const int i0_stop = i0_start + 2*c::nbatch_V2 < DV ? i0_start + 2*c::nbatch_V2 : DV;
const int i0_diff = i0_stop - i0_start;
if (nstages == 1) {
if (nstages <= 1) {
constexpr bool use_cp_async = nstages == 1;
flash_attn_ext_f16_load_tile<stride_tile_V, nwarps, c::nbatch_fa, use_cp_async>
(V_h2 + k_VKQ_0*stride_V + i0_start/2, tile_V, i0_diff/2, stride_V);