From 2ad093af75d1905d933e5005024b38d770f482ba Mon Sep 17 00:00:00 2001 From: Concedo <39025047+LostRuins@users.noreply.github.com> Date: Fri, 3 Jul 2026 15:27:30 +0800 Subject: [PATCH] Revert "CUDA: consistent use of __restrict__ + PDL for FA (#25185)" This reverts commit b820cc8e6f94843d32f92c8ebd7db837dae2bd8b. --- ggml/src/ggml-cuda/fattn-common.cuh | 10 +++------- 1 file changed, 3 insertions(+), 7 deletions(-) diff --git a/ggml/src/ggml-cuda/fattn-common.cuh b/ggml/src/ggml-cuda/fattn-common.cuh index e67cc7fdf..b76121fef 100644 --- a/ggml/src/ggml-cuda/fattn-common.cuh +++ b/ggml/src/ggml-cuda/fattn-common.cuh @@ -664,10 +664,7 @@ constexpr __device__ dequantize_V_t get_dequantize_V() { template __launch_bounds__(FATTN_KQ_STRIDE/2, 1) static __global__ void flash_attn_mask_to_KV_max( - const half2 * mask_ptr, int * KV_max_ptr, const int ne30, const int64_t s31, const int64_t s33) { - const half2 * GGML_CUDA_RESTRICT mask = mask_ptr; - int * GGML_CUDA_RESTRICT KV_max = KV_max_ptr; - + const half2 * __restrict__ mask, int * __restrict__ KV_max, const int ne30, const int64_t s31, const int64_t s33) { const int ne31 = gridDim.x; const int tid = threadIdx.x; const int sequence = blockIdx.y; @@ -1102,9 +1099,8 @@ void launch_fattn( const int iter_k = K->ne[1] / FATTN_KQ_STRIDE; KV_max.alloc(ne_KV_max); - ggml_cuda_kernel_launch_params launch_params = ggml_cuda_kernel_launch_params(blocks_num_KV_max, block_dim_KV_max, 0, main_stream); - ggml_cuda_kernel_launch(flash_attn_mask_to_KV_max, launch_params, - (const half2 *) mask->data, KV_max.ptr, iter_k, s31, s33); + flash_attn_mask_to_KV_max<<>> + ((const half2 *) mask->data, KV_max.ptr, iter_k, s31, s33); CUDA_CHECK(cudaGetLastError()); }