diff --git a/.github/workflows/ci_gpu.yml b/.github/workflows/ci_gpu.yml index d151506..9162d3b 100644 --- a/.github/workflows/ci_gpu.yml +++ b/.github/workflows/ci_gpu.yml @@ -31,19 +31,19 @@ jobs: run: python train_gpt2.py - name: Compile training and testing program - run: make test_gpt2cu train_gpt2cu test_gpt2fp32cu train_gpt2fp32cu + run: make test_gpt2cu train_gpt2cu test_gpt2fp32cu train_gpt2fp32cu - name: Train model (With OpenMP) run: OMP_NUM_THREADS=8 ./train_gpt2cu - name: Train model (FP32) with gpt2_124M.bin - run: | + run: | PRECISION=FP32 make train_gpt2cu ./train_gpt2cu -b 4 -t 64 -l 1e-4 -v 200 -s 200 -a 1 -x 10 -e gpt2_124M.bin - + - name: Build FP32 precision run: PRECISION=FP32 make test_gpt2cu profile_gpt2cu - + - name: Run default run: ./test_gpt2cu @@ -52,7 +52,7 @@ jobs: - name: Run recompute LN run: ./test_gpt2cu -r 2 - + - name: Build BF16 precision run: PRECISION=BF16 make train_gpt2cu test_gpt2cu profile_gpt2cu @@ -67,18 +67,18 @@ jobs: - name: Run recompute LN run: ./test_gpt2cu -r 2 - + - name: Train model fp32 (With OpenMP) run: OMP_NUM_THREADS=8 ./train_gpt2fp32cu - name: Execute testing program (With OpenMP) run: OMP_NUM_THREADS=8 ./test_gpt2cu - + - name: Execute testing program fp32 (With OpenMP) run: OMP_NUM_THREADS=8 ./test_gpt2fp32cu - name: Compile training and testing program without OpenMP - run: NO_OMP=1 make test_gpt2cu train_gpt2cu test_gpt2fp32cu train_gpt2fp32cu + run: NO_OMP=1 make test_gpt2cu train_gpt2cu test_gpt2fp32cu train_gpt2fp32cu - name: Train model (No OpenMP) run: NO_OMP=1 ./train_gpt2cu @@ -88,14 +88,14 @@ jobs: - name: Execute testing program (No OpenMP) run: ./test_gpt2cu -b 32 - + - name: Execute testing program fp32 (No OpenMP) run: ./test_gpt2fp32cu - name: Install cuDNN-frontend - run: + run: git clone https://github.com/NVIDIA/cudnn-frontend.git - + - name: Build with cuDNN run: USE_CUDNN=1 make test_gpt2cu train_gpt2cu test_gpt2fp32cu train_gpt2fp32cu diff --git a/dev/cuda/adamw.cu b/dev/cuda/adamw.cu index 20a6560..74dfc2e 100644 --- a/dev/cuda/adamw.cu +++ b/dev/cuda/adamw.cu @@ -159,7 +159,7 @@ int main(int argc, char **argv) { // create random data on host (to be used for the CPU reference implementation) float* params_memory = make_random_float(num_parameters); float* grads_memory = make_random_float(num_parameters); - float* m_memory = make_random_float_01(num_parameters); + float* m_memory = make_random_float(num_parameters); float* v_memory = make_random_float_01(num_parameters); // move to GPU diff --git a/dev/cuda/classifier_fused.cu b/dev/cuda/classifier_fused.cu index d58f101..9d93f9d 100644 --- a/dev/cuda/classifier_fused.cu +++ b/dev/cuda/classifier_fused.cu @@ -114,7 +114,7 @@ __device__ SoftmaxParams prepare_softmax(cg::thread_block_tile<32>& warp, int64_t idx, const float* inp, int V, int P) { // this warp (of 32) threads processes one row of inp, i.e. inp[idx, :] of shape (V,) // note that inp is actually (B * T, P) but we only use the first V elements - // this function tehen calculates: + // this function then calculates: // 1) the max value to subtract for numerical stability and // 2) the sum normalization factor const float* x = inp + idx * P; @@ -680,8 +680,8 @@ int main(int argc, char **argv) { cudaCheck(cudaSetDevice(deviceIdx)); // create host memory of random numbers - float* logits = make_random_float_01(B * T * V); - float* probs = (float*)malloc(B * T * V * sizeof(float)); + float* logits = make_random_float(B * T * V); + float* probs = make_random_float_01(B * T * V); float* dlogits = (float*)malloc(B * T * V * sizeof(float)); float* losses = (float*)malloc(B * T * sizeof(float)); float* dlosses = make_random_float(B * T); @@ -760,6 +760,7 @@ int main(int argc, char **argv) { free(losses); free(dlosses); free(targets); + free(outliers); cudaCheck(cudaFree(d_dlogits)); cudaCheck(cudaFree(d_losses)); cudaCheck(cudaFree(d_logits)); diff --git a/dev/cuda/crossentropy_softmax_backward.cu b/dev/cuda/crossentropy_softmax_backward.cu index 27521bf..65c72a2 100644 --- a/dev/cuda/crossentropy_softmax_backward.cu +++ b/dev/cuda/crossentropy_softmax_backward.cu @@ -99,7 +99,7 @@ int main(int argc, char **argv) { cudaCheck(cudaSetDevice(deviceIdx)); // create host memory of random numbers - float* probs = make_random_float(B * T * V); + float* probs = make_random_float_01(B * T * V); int* targets = make_random_int(B * T, V); float* dlosses = make_random_float(B * T); float* dlogits = make_zeros_float(B * T * V); diff --git a/llmc/gelu.cuh b/llmc/gelu.cuh index ce9aa12..cd5c297 100644 --- a/llmc/gelu.cuh +++ b/llmc/gelu.cuh @@ -50,7 +50,7 @@ __global__ void gelu_backward_inplace_kernel(floatX* d_in_out, const floatX* inp void gelu_forward(floatX* out, const floatX* inp, int N, cudaStream_t stream) { NVTX_RANGE_FN(); const int block_size = 512; - assert(N % block_size == 0); + assert(N % (block_size * x128::size) == 0); const int grid_size = CEIL_DIV(N, block_size * x128::size); gelu_forward_kernel2<<>>(out, inp); cudaCheck(cudaGetLastError()); @@ -59,7 +59,7 @@ void gelu_forward(floatX* out, const floatX* inp, int N, cudaStream_t stream) { void gelu_backward_inplace(floatX* d_in_out, const floatX* inp, const int N, cudaStream_t stream) { NVTX_RANGE_FN(); const int block_size = 128; - assert(N % block_size == 0); + assert(N % (block_size * x128::size) == 0); const int grid_size = CEIL_DIV(N, block_size * x128::size); gelu_backward_inplace_kernel<<>>(d_in_out, inp); cudaCheck(cudaGetLastError()); diff --git a/llmc/layernorm.cuh b/llmc/layernorm.cuh index a49f0bf..9c9afa3 100644 --- a/llmc/layernorm.cuh +++ b/llmc/layernorm.cuh @@ -368,7 +368,7 @@ void layernorm_forward(floatX* out, floatX* mean, floatX* rstd, void residual_forward(floatX* out, const floatX* inp1, const floatX* inp2, int N, cudaStream_t stream) { NVTX_RANGE_FN(); const int block_size = 256; - assert(N % block_size == 0); + assert(N % (block_size * x128::size) == 0); const int grid_size = CEIL_DIV(N, block_size * x128::size); residual_forward_kernel<<>>(out, inp1, inp2); cudaCheck(cudaGetLastError());