diff --git a/ggml/src/ggml-sycl/fattn-vec.hpp b/ggml/src/ggml-sycl/fattn-vec.hpp index 04baac4414..53ad0eaee4 100644 --- a/ggml/src/ggml-sycl/fattn-vec.hpp +++ b/ggml/src/ggml-sycl/fattn-vec.hpp @@ -73,6 +73,7 @@ static void flash_attn_ext_vec(const char* __restrict__ Q, const int32_t nb31, const int32_t nb32, const int64_t nb33) { + #ifdef SYCL_FLASH_ATTN // Skip unused kernel variants for faster compilation: @@ -469,7 +470,6 @@ static void flash_attn_ext_vec(const char* __restrict__ Q, } } - item_ct1.barrier(sycl::access::fence_space::local_space); #pragma unroll @@ -591,22 +591,24 @@ void ggml_sycl_flash_attn_ext_vec_case_impl(ggml_backend_sycl_context & ctx, ggm const auto arch = ggml_sycl_info().devices[ctx.device].hw_info.arch; const int nthreads = ggml_sycl_fattn_vec_get_nthreads_device(arch); - // 256 threads would overflow the 64 KB work-group local memory at D == 512, so keep 128 there. - if (D <= 256 && nthreads == 256) { - constexpr int nthreads_hw = 256; - constexpr int nwarps = nthreads_hw / warp_size; - launch_fattn, warp_size>( - ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false); - } else { - constexpr int nthreads_hw = 128; - constexpr int nwarps = nthreads_hw / warp_size; - launch_fattn, warp_size>( - ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false); + if constexpr (D <= 256) { + if (nthreads == 256) { + constexpr int nthreads_hw = 256; + constexpr int nwarps = nthreads_hw / warp_size; + launch_fattn, warp_size>( + ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false); + return; + } } + + constexpr int nthreads_hw = 128; + constexpr int nwarps = nthreads_hw / warp_size; + launch_fattn, warp_size>( + ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false); } template