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

gla: Put the barrier inside the main logic loop

Browse source
This commit is contained in:
Akarshan Biswas 2025-01-13 17:44:07 +05:30
parent 81d852902e
commit db3ded2c8a
No known key found for this signature in database
GPG key ID: 52A578A14B32134D
1 changed files with 3 additions and 1 deletions
Download patch file Download diff file Expand all files Collapse all files

4
ggml/src/ggml-sycl/gla.cpp
View file

@ -30,9 +30,11 @@ static void gated_linear_attn_f32_kernel(const dpct::queue_ptr stream, u_int B,
for (u_int i = 0; i < head_size; i++) { for (u_int i = 0; i < head_size; i++) {
state[i] = s[batch_i * state_size + head_i * head_size * head_size + i * head_size + tid]; state[i] = s[batch_i * state_size + head_i * head_size * head_size + i * head_size + tid];
} }
item.barrier(sycl::access::fence_space::local_space); //sync threads
for (u_int t = batch_i * n_seq_tokens * C + head_i * head_size + tid; for (u_int t = batch_i * n_seq_tokens * C + head_i * head_size + tid;
t < (batch_i + 1) * n_seq_tokens * C + head_i * head_size + tid; t += C) { t < (batch_i + 1) * n_seq_tokens * C + head_i * head_size + tid; t += C) {
item.barrier(sycl::access::fence_space::local_space); //sync threads
_k[tid] = k[t]; _k[tid] = k[t];
_r[tid] = r[t]; _r[tid] = r[t];
_td[tid] = td[t]; _td[tid] = td[t];