aparler/c-llm
1
0
Fork 0
mirror of https://github.com/karpathy/llm.c.git synced 2026-07-27 20:25:09 -04:00
Code Issues Projects Releases Packages Wiki Activity Actions

fixed numerics for online kernel, and an existing race condition in kernel 2

Browse source
This commit is contained in:
Erik Schultheis 2024-04-15 03:41:09 +03:00
parent 8fb966bc31
commit c73fa8842a
1 changed files with 2 additions and 2 deletions
Download patch file Download diff file Expand all files Collapse all files

4
dev/cuda/softmax_forward.cu
View file

@ -143,8 +143,8 @@ __global__ void softmax_forward_kernel2(float* out, const float* inp, int N, int
shared[tid] = fmaxf(shared[tid], shared[tid + stride]);
}
}
float offset = shared[0];
__syncthreads();
float offset = shared[0];
// compute expf and write the result to global memory
for (int i = tid; i < C; i += block_size) {
out[idx * C + i] = expf(x[i] - offset);
@ -327,7 +327,7 @@ __global__ void softmax_forward_online_kernel1(float* out, const float* inp, int
float* out_row = out + i * C;
float maxval = -INFINITY;
float sum = 0.0f;
double sum = 0.0f;
for (int j = 0; j < C; j++) {
float maxval_prev = maxval;
if (inp_row[j] > maxval) {