Skip to content

Commit 27e8f08

Browse files
committed
ggml-sycl : fix padding in timestep embedding kernel
This commit fixes the zero padding for odd dimensions in the timestep embedding kernel similar to the fix that was applied to the cpu backend in Commit 9de447d ("ggml-cpu : fix padding in ggml_timestep_embedding (#15917)"). The motivation for this is that currently if an odd dimension is used, the padding check incorrectly uses the dimension value for indexing. For example, with dim=15: Elements 0-6 are set to cosine values Elements 7-13 are set to sine values Element 14 is left uninitialized (contains garbage) Element 15 is correctly set to zero This fix changes embed_data[dim] to embed_data[2 * half] so that element 14 (the first unused element) is properly set to zero as well as the last element.
1 parent 8263ffa commit 27e8f08

File tree

1 file changed

+3
-1
lines changed

1 file changed

+3
-1
lines changed

ggml/src/ggml-sycl/tsembd.cpp

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,11 +21,13 @@ static void timestep_embedding_f32(
2121
int j = item_ct1.get_local_id(2) + item_ct1.get_group(2) * item_ct1.get_local_range(2);
2222
float * embed_data = (float *)((char *)dst + i*nb1);
2323

24+
int half = dim / 2;
25+
2426
if (dim % 2 != 0 && j == ((dim + 1) / 2)) {
27+
embed_data[2 * half] = 0.f;
2528
embed_data[dim] = 0.f;
2629
}
2730

28-
int half = dim / 2;
2931
if (j >= half) {
3032
return;
3133
}

0 commit comments

Comments
 (0)