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

Add a local convenience Makefile for dev/cuda/, fix 2 minor warnings observed when building all targets.

Browse source
This commit is contained in:
austinvhuang 2024-04-26 21:19:44 -04:00
parent 311d6d8441
commit d937f93aaa
4 changed files with 92 additions and 3 deletions
Download patch file Download diff file Expand all files Collapse all files

1
.gitignore vendored
View file

@ -20,6 +20,7 @@ dev/cuda/*_backward
dev/cuda/classifier_fused
dev/cuda/adamw
dev/cuda/matmul_backward_bias
dev/cuda/nccl_all_reduce
*.obj
*.exe

88
dev/cuda/Makefile Normal file
View file

@ -0,0 +1,88 @@
NVCC := $(shell which nvcc 2>/dev/null)
ifeq ($(NVCC),)
$(error nvcc not found.)
endif
TARGETS = adamw attention_backward attention_forward classifier_fused crossentropy_forward crossentropy_softmax_backward encoder_backward encoder_forward gelu_forward layernorm_backward layernorm_forward matmul_backward matmul_backward_bias matmul_forward nccl_all_reduce residual_forward softmax_forward trimat_forward
CFLAGS = -O3 --use_fast_math
all: $(TARGETS)
adamw: adamw.cu
$(NVCC) $(CFLAGS) adamw.cu -o adamw
attention_backward: attention_backward.cu
$(NVCC) $(CFLAGS) attention_backward.cu -o attention_backward -lcublas
attention_forward: attention_forward.cu
$(NVCC) $(CFLAGS) attention_forward.cu -o attention_forward -lcublas
classifier_fused: classifier_fused.cu
$(NVCC) $(CFLAGS) classifier_fused.cu -o classifier_fused
crossentropy_forward: crossentropy_forward.cu
$(NVCC) $(CFLAGS) crossentropy_forward.cu -o crossentropy_forward
crossentropy_softmax_backward: crossentropy_softmax_backward.cu
$(NVCC) $(CFLAGS) crossentropy_softmax_backward.cu -o crossentropy_softmax_backward
encoder_backward: encoder_backward.cu
$(NVCC) $(CFLAGS) encoder_backward.cu -o encoder_backward
encoder_forward: encoder_forward.cu
$(NVCC) $(CFLAGS) encoder_forward.cu -o encoder_forward
gelu_forward: gelu_forward.cu
$(NVCC) $(CFLAGS) gelu_forward.cu -o gelu_forward
layernorm_backward: layernorm_backward.cu
$(NVCC) $(CFLAGS) layernorm_backward.cu -o layernorm_backward
layernorm_forward: layernorm_forward.cu
$(NVCC) $(CFLAGS) layernorm_forward.cu -o layernorm_forward
matmul_backward: matmul_backward.cu
$(NVCC) $(CFLAGS) -Xcompiler -fopenmp matmul_backward.cu -o matmul_backward -lcublas
matmul_backward_bias: matmul_backward_bias.cu
$(NVCC) $(CFLAGS) matmul_backward_bias.cu -o matmul_backward_bias
matmul_forward: matmul_forward.cu
$(NVCC) $(CFLAGS) -Xcompiler -fopenmp matmul_forward.cu -o matmul_forward -lcublas -lcublasLt
nccl_all_reduce: nccl_all_reduce.cu
$(NVCC) -lmpi -lnccl -I/usr/lib/x86_64-linux-gnu/openmpi/include -L/usr/lib/x86_64-linux-gnu/openmpi/lib/ nccl_all_reduce.cu -o nccl_all_reduce
residual_forward: residual_forward.cu
$(NVCC) $(CFLAGS) residual_forward.cu -o residual_forward
softmax_forward: softmax_forward.cu
$(NVCC) $(CFLAGS) softmax_forward.cu -o softmax_forward
trimat_forward: trimat_forward.cu
$(NVCC) $(CFLAGS) trimat_forward.cu -o trimat_forward -lcublas
clean:
rm -f $(TARGETS)
run_all: all
./adamw
./attention_backward
./attention_forward
./classifier_fused
./crossentropy_forward
./crossentropy_softmax_backward
./encoder_backward
./encoder_forward
./gelu_forward
./layernorm_backward
./layernorm_forward
./matmul_backward
./matmul_backward_bias
./matmul_forward
./nccl_all_reduce
./residual_forward
./softmax_forward
./trimat_forward

4
dev/cuda/encoder_backward.cu
View file

@ -75,7 +75,7 @@ __global__ void encoder_backward_kernel2(float* dwte, float* dwpe,
int c = blockIdx.x * blockDim.x + threadIdx.x;
if (c >= C) { return; } // guard
int BT = B * T;
for (int i = 0; i < B * T; i++) {
for (int i = 0; i < BT; i++) {
int t = i % T;
int ix = inp[i];
float dout_btc = dout[i * C + c];
@ -196,4 +196,4 @@ int main(int argc, char **argv) {
cudaFree(d_dwpe);
return 0;
}
}

2
dev/cuda/layernorm_backward.cu
View file

@ -122,7 +122,7 @@ __global__ void layernorm_backward_kernel1(float* dinp, float* dweight, float* d
const float* dout_bt = dout + b * T * C + t * C;
const float* inp_bt = inp + b * T * C + t * C;
float* dinp_bt = dinp + b * T * C + t * C;
const const float mean_bt = mean[b * T + t];
const float mean_bt = mean[b * T + t];
const float rstd_bt = rstd[b * T + t];
// first: two reduce operations