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

Initialize nreduce as size_t

Browse source
This commit is contained in:
Akarshan Biswas 2024-12-10 17:28:48 +05:30
parent fb2e66e825
commit 32164aa48d
No known key found for this signature in database
GPG key ID: 52A578A14B32134D
1 changed files with 6 additions and 6 deletions
Download patch file Download diff file Expand all files Collapse all files

12
ggml/src/ggml-sycl/norm.cpp
View file

@ -31,8 +31,8 @@ static void norm_f32(const float* x, float* dst, const int ncols, const float ep
*/
item_ct1.barrier(sycl::access::fence_space::local_space);
mean_var = 0.f;
int nreduce = nwarps / WARP_SIZE;
for (size_t i = 0; i < (size_t) nreduce; i += 1)
size_t nreduce = nwarps / WARP_SIZE;
for (size_t i = 0; i < nreduce; i += 1)
{
mean_var += s_sum[lane_id + i * WARP_SIZE];
}
@ -55,7 +55,7 @@ static void group_norm_f32(const float* x, float* dst, const int group_size, con
const int nthreads = item_ct1.get_local_range(2);
const int nwarps = nthreads / WARP_SIZE;
start += item_ct1.get_local_id(2);
int nreduce = nwarps / WARP_SIZE;
size_t nreduce = nwarps / WARP_SIZE;
if (end >= ne_elements) {
end = ne_elements;
@ -86,7 +86,7 @@ static void group_norm_f32(const float* x, float* dst, const int group_size, con
*/
item_ct1.barrier();
tmp = 0.f;
for (size_t i = 0; i < (size_t) nreduce; i += 1)
for (size_t i = 0; i < nreduce; i += 1)
{
tmp += s_sum[lane_id + i * WARP_SIZE];
}
@ -121,7 +121,7 @@ static void group_norm_f32(const float* x, float* dst, const int group_size, con
*/
item_ct1.barrier();
tmp = 0.f;
for (size_t i = 0; i < (size_t) nreduce; i += 1)
for (size_t i = 0; i < nreduce; i += 1)
{
tmp += s_sum[lane_id + i * WARP_SIZE];
}
@ -163,7 +163,7 @@ static void rms_norm_f32(const float* x, float* dst, const int ncols, const floa
converged control flow. You may need to adjust the code.
*/
item_ct1.barrier(sycl::access::fence_space::local_space);
int nreduce = nwarps / WARP_SIZE;
size_t nreduce = nwarps / WARP_SIZE;
tmp = 0.f;
for (size_t i = 0; i < nreduce; i += 1)
{