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

careful to guard the overloaded warpReduceSum and small fixes

Browse source
This commit is contained in:
Andrej Karpathy 2024-05-02 19:24:04 +00:00
parent 6879c36cf3
commit 01df41fbba
1 changed files with 9 additions and 7 deletions
Download patch file Download diff file Expand all files Collapse all files

16
train_gpt2.cu
View file

@ -213,13 +213,6 @@ __device__ void atomicAddX(float* addr, float val) {
atomicAdd(addr, val);
}
__device__ floatX warpReduceSum(floatX val) {
for (int offset = 16; offset > 0; offset /= 2) {
val += __shfl_xor_sync(0xFFFFFFFF, val, offset);
}
return val;
}
// warp-level reduction for summing values
__device__ float warpReduceSum(float val) {
for (int offset = 16; offset > 0; offset /= 2) {
@ -236,6 +229,15 @@ __device__ float warpReduceMax(float val) {
return val;
}
#if defined(ENABLE_BF16) || defined(ENABLE_FP16)
__device__ floatX warpReduceSum(floatX val) {
for (int offset = 16; offset > 0; offset /= 2) {
val += __shfl_xor_sync(0xFFFFFFFF, val, offset);
}
return val;
}
#endif
// ----------------------------------------------------------------------------
// Packed128 data structure, which forces the compiler to use 128-bit loads/stores
// in GPUs that support (the LDG.128 and STS.128 instructions)