Revert "Fix race in the CUDA DeepSeek FA kernel (#406)"
This reverts commit 36e6e888b75ae93fb5aac212bb0e147d8379ae23. I should have tested. We get NaNs.
This commit is contained in:
parent
0abcf0749e
commit
8f7bd74afb
|
|
@ -898,8 +898,6 @@ static __device__ __forceinline__ void flash_attn_ext_f16_process_tile(
|
||||||
KQ_crs += __shfl_xor_sync(0xFFFFFFFF, KQ_crs, offset, WARP_SIZE);
|
KQ_crs += __shfl_xor_sync(0xFFFFFFFF, KQ_crs, offset, WARP_SIZE);
|
||||||
}
|
}
|
||||||
|
|
||||||
__syncthreads();
|
|
||||||
|
|
||||||
// Write back combined meta data:
|
// Write back combined meta data:
|
||||||
#pragma unroll
|
#pragma unroll
|
||||||
for (int imeta = 0; imeta < nmeta; ++imeta) {
|
for (int imeta = 0; imeta < nmeta; ++imeta) {
|
||||||
|
|
|
||||||
Loading…
Reference in New Issue