Skip to content

[SYCL] fix error: Error OP FLASH_ATTN_EXT on arc770 #26441

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ggerganov merged 1 commit into from
Aug 7, 2026
+18 −16
Merged

[SYCL] fix error: Error OP FLASH_ATTN_EXT on arc770 #26441

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
34 changes: 18 additions & 16 deletions ggml/src/ggml-sycl/fattn-vec.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:

Expand Down Expand Up @@ -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
Expand Down Expand Up @@ -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<D, cols_per_block, 1,
flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
use_logit_softcap, warp_size, nthreads_hw>, 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<D, cols_per_block, 1,
flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
use_logit_softcap, warp_size, nthreads_hw>, 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<D, cols_per_block, 1,
flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
use_logit_softcap, warp_size, nthreads_hw>, 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<D, cols_per_block, 1,
flash_attn_ext_vec<D, cols_per_block, type_K, type_V,
use_logit_softcap, warp_size, nthreads_hw>, warp_size>(
ctx, dst, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false);
}

template <int D, int type_K, int type_V>
Expand Down
Loading