Daniel Bevenius
3913f8730e
* ggml : remove adding extra dim timestep embedding This commit updates the ggml_timestep_embedding function to no longer add an extra dimension when the specified dimension is odd. The motivation for this change is that this introduces an unnecessary dimension when the dimension is odd, which caused an issue in the kernels which were not expecting this extra dimension and it resulted in uninitialized memory for the second to last dimension. * ggml-cuda : fix padding in timestep embedding kernel This commit removes the zeroing out of the last dimension now that we are not adding the extra padding dimension. * ggml-metal : fix padding in timestep embedding kernel This commit fixes the zero padding for odd dimensions in the timestep embedding kernel * ggml-opencl : fix padding in timestep embedding kernel This commit fixes the zero padding for odd dimensions in the timestep embedding kernel. * ggml-sycl : fix padding in timestep embedding kernel This commit fixes the zero padding for odd dimensions in the timestep embedding kernel. * ggml-vulkan : fix padding in timestep embedding kernel This commit fixes the zero padding for odd dimensions in the timestep embedding kernel. * ggml-cpu : fix padding in timestep embedding function This commit removes the zeroing out of the last dimension now that we are not adding the extra padding dimension.
49 lines
1.4 KiB
Common Lisp
49 lines
1.4 KiB
Common Lisp
kernel void kernel_timestep_embedding(
|
|
global const void * p_timesteps,
|
|
ulong off_timesteps,
|
|
global void * p_dst,
|
|
ulong off_dst,
|
|
int dst_nb1_bytes,
|
|
int logical_dim,
|
|
int max_period
|
|
) {
|
|
int local_i;
|
|
int local_j;
|
|
int local_half_dim;
|
|
float local_timestep_val;
|
|
float local_freq;
|
|
float local_arg;
|
|
global float * local_embed_data_ptr;
|
|
global const float * local_timesteps_input_ptr;
|
|
global float * local_dst_output_base_ptr;
|
|
|
|
local_timesteps_input_ptr = (global const float *)((global char *)p_timesteps + off_timesteps);
|
|
local_dst_output_base_ptr = (global float *)((global char *)p_dst + off_dst);
|
|
|
|
local_i = get_global_id(1);
|
|
local_j = get_global_id(0);
|
|
|
|
local_half_dim = logical_dim / 2;
|
|
local_embed_data_ptr = (global float *)((global char *)local_dst_output_base_ptr + local_i * dst_nb1_bytes);
|
|
|
|
if (logical_dim % 2 != 0 && local_j == local_half_dim) {
|
|
local_embed_data_ptr[2 * local_half_dim] = 0.0f;
|
|
}
|
|
|
|
if (local_j >= local_half_dim) {
|
|
return;
|
|
}
|
|
|
|
local_timestep_val = local_timesteps_input_ptr[local_i];
|
|
|
|
if (local_half_dim == 0) {
|
|
local_freq = 1.0f;
|
|
} else {
|
|
local_freq = exp(-log((float)max_period) * (float)local_j / (float)local_half_dim);
|
|
}
|
|
|
|
local_arg = local_timestep_val * local_freq;
|
|
local_embed_data_ptr[local_j] = cos(local_arg);
|
|
local_embed_data_ptr[local_j + local_half_dim] = sin(local_arg);
|
|
}
|