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

fix the compiler warnings and errors

Browse source
This commit is contained in:
lancerts 2024-06-06 11:46:40 -07:00
parent d931fc0507
commit 4b41087c64
1 changed files with 4 additions and 4 deletions
Download patch file Download diff file Expand all files Collapse all files

8
dev/cuda/matmul_backward_bias.cu
View file

@ -2,7 +2,7 @@
Kernels for matmul backward pass bias only.
Compile example:
nvcc -O3 -lcublas -lcublasLt matmul_backward_bias.cu -lineinfo -o matmul_backward_bias
nvcc -O3 -lcublas -lcublasLt -std=c++17 matmul_backward_bias.cu -lineinfo -o matmul_backward_bias
./matmul_backward_bias 1
./matmul_backward_bias 2
@ -116,7 +116,7 @@ __global__ void matmul_backward_bias_kernel2(floatX* dbias, const floatX* dout,
sum = cg::reduce(warp, sum, cg::plus<float>{});
// write the result to output (global memory)
if(warp.thread_rank() == 0) {
dbias[idx] += sum;
dbias[idx] = (float)dbias[idx] + sum;
}
}
@ -148,7 +148,7 @@ __global__ void matmul_backward_bias_kernel3(floatX* dbias, const floatX* dout,
float block_sum = cg::reduce(warp, warp_sum, cg::plus<float>{}); // sum(x)
// write the result to output (global memory)
if(threadIdx.x == 0) {
dbias[idx] += block_sum;
dbias[idx] = (float)dbias[idx] + block_sum;
}
}
@ -188,7 +188,7 @@ __global__ void matmul_backward_bias_kernel4(floatX* dbias, const floatX* dout,
for (int j = 0; j < vstep; j++) {
dout_sum += smem[lane_id + j * warpSize];
}
dbias[tl + lane_id] += dout_sum;
dbias[tl + lane_id] = (float)dbias[tl + lane_id] + dout_sum;
}
}