From 9374d751de4da08f6b7aade24e766e2e4ec1e34d Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Wed, 29 May 2024 00:48:37 +0300 Subject: [PATCH 01/20] MFU for other GPUs --- train_gpt2.cu | 76 +++++++++++++++++++++++++++++++++++++++++++++++++-- 1 file changed, 74 insertions(+), 2 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index c511b49..070f6d0 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2872,6 +2872,74 @@ void gpt2_multi_gpu_gather(GPT2 *model, MultiGpuConfig* multi_gpu_config) #endif } +float get_flops_promised(const char* device) { + // for the non-top models, actual performance numbers aren't that easy to find, e.g., + // here https://www.techpowerup.com/gpu-specs/rtx-a4000.c3756, does "Theoretical Performance" + // seems to be without tensor cores. + // So, instead we use that all these cards just use the same types of tensor cores in different + // numbers and at different frequencies. Then we just need to look up these two easily accesible + // numbers for all the other GPUs. + // linear scaling seems to work: comparing spec sheet and calculation: + // 4080: 304TCs, 2505 GHz; 97.5TFlops = 165.2/512*304 /2520 * 2505 + // + // Original numbers for the top GPUS are from. + // https://resources.nvidia.com/en-us-tensor-core + // https://images.nvidia.com/aem-dam/Solutions/geforce/ada/nvidia-ada-gpu-architecture.pdf + + struct PerfData { + float TF_32; // tensor-core performance 32 bit + float BF_16_32; // bf16 with 32 bit accumulate + float FP_16_32; // fp16 with 32 bit accumulate + float FP_16_16; // fp16 with 16 bit accumulate + float FP_8_32; // and so on + float FP_8_16; + float CLOCK; // clock frequency from the spec sheet + float CORES; // #TCs from the spec sheet + }; + + // basic data from the nvidia whitepapers + static const PerfData AMPERE = {156.f, 312.f, 312.f, 312.f, -1, -1, 1410.f, 432.f}; + static const PerfData HOPPER = {378.f, 756.f, 756.f, 756.f, 1513.f, 1513.f, 1620.f, 456.f}; + static const PerfData ADA = {82.6f, 165.2f, 165.2f, 330.3f, 330.3f, 660.6f, 2520.f, 512.f}; + + // helper to get the value for a particular GPU with different #TCs and clock frequency. + auto adjust = [](const PerfData& orig, float new_cores, float new_mhz){ + float value; + switch (PRECISION_MODE) { + case PRECISION_BF16: + value = orig.BF_16_32; + break; + case PRECISION_FP32: + value = orig.TF_32; + break; + case PRECISION_FP16: + value = orig.FP_16_32; + break; + default: + fprintf(stderr, "Invalid precision mode, how did this even compile?"); + exit(1); + } + return value / orig.CORES * new_cores / orig.CLOCK * new_mhz; + }; + + // TODO fill in more GPUs + static std::pair gpu_db[] = { + {"NVIDIA A100-PCIE-40GB", adjust(AMPERE, 432, 1410)}, + {"NVIDIA GeForce RTX 4060 Ti", adjust(ADA, 136, 2535)}, + {"NVIDIA RTX A4000", adjust(AMPERE, 192, 1560)}, + }; + + // just do a linear search until you find our GPU + for(const auto& entry : gpu_db) { + if(entry.first == device) { + return entry.second; + } + } + + // signal that we don't know + return -1; +} + float gpt2_estimate_mfu(GPT2 *model, int num_tokens, float dt) { // estimate model flops utilization (MFU) in units of A100 bfloat16 peak FLOPS // see PaLM paper Appendix B as ref: https://arxiv.org/abs/2204.02311 @@ -2885,7 +2953,10 @@ float gpt2_estimate_mfu(GPT2 *model, int num_tokens, float dt) { size_t flops_per_step = flops_per_token * num_tokens; // express our flops throughput as ratio of A100 bfloat16 peak flops float flops_achieved = (float)flops_per_step * (1.0f / dt); // per second - float flops_promised = 312e12f; // A100 GPU bfloat16 peak flops is 312 TFLOPS + float flops_promised = get_flops_promised(deviceProp.name) * 1e12f; + if(flops_promised < 0) { + return -1.f; // don't know + } float mfu = flops_achieved / flops_promised; return mfu; } @@ -3258,6 +3329,7 @@ int main(int argc, char *argv[]) { ? (cublas_compute == CUBLAS_COMPUTE_32F_FAST_TF32 ? "TF32" : "FP32") : (PRECISION_MODE == PRECISION_FP16 ? "FP16" : "BF16"); printf0("| device | %-50s |\n", deviceProp.name); + printf0("| TFlops | %-50.1f |\n", get_flops_promised(deviceProp.name)); printf0("| precision | %-50s |\n", precision_str); printf0("+-----------------------+----------------------------------------------------+\n"); @@ -3563,7 +3635,7 @@ int main(int argc, char *argv[]) { } float accumulated_loss = multi_gpu_config.num_processes == 1 ? model.mean_loss : model.accumulated_mean_loss; float mfu = gpt2_estimate_mfu(&model, B * T * grad_accum_steps, time_elapsed_ms / 1000.0f); - printf0("step %4d/%d | train loss %7.6f | norm %6.4f | lr %.2e | %.2f ms | %.1f%% A100 fp16 MFU | %.0f tok/s\n", + printf0("step %4d/%d | train loss %7.6f | norm %6.4f | lr %.2e | %.2f ms | %.1f%% bf16 MFU | %.0f tok/s\n", step + 1, train_num_batches, accumulated_loss, grad_norm, step_learning_rate, time_elapsed_ms, 100*mfu, bias_corrected_ema_tokens_per_second); logger_log_train(&logger, step, model.mean_loss); From e2b411362efd95a910edc390c182b11b1826ad52 Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Wed, 29 May 2024 00:54:23 +0300 Subject: [PATCH 02/20] include string view --- train_gpt2.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/train_gpt2.cu b/train_gpt2.cu index 070f6d0..e5df83b 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -39,6 +39,7 @@ This reads & runs in fp32, B=4, T=64, LR=1e-4, val/sample never (200), #include #include #include +#include #include #include #include From 184cc9d10a196cdaf29de1b631f046dd288d61ba Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Wed, 29 May 2024 01:31:16 +0300 Subject: [PATCH 03/20] added H100 SMX version --- train_gpt2.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/train_gpt2.cu b/train_gpt2.cu index e5df83b..24f6c4f 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2928,6 +2928,7 @@ float get_flops_promised(const char* device) { {"NVIDIA A100-PCIE-40GB", adjust(AMPERE, 432, 1410)}, {"NVIDIA GeForce RTX 4060 Ti", adjust(ADA, 136, 2535)}, {"NVIDIA RTX A4000", adjust(AMPERE, 192, 1560)}, + {"NVIDIA H100 80GB HBM3", adjust(HOPPER, 528, 1830)}, // HBM3 = SXM5 }; // just do a linear search until you find our GPU From de7a9b83cd2b86fc492c69f78b57f96c0d20cae2 Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Wed, 29 May 2024 02:09:20 +0300 Subject: [PATCH 04/20] more GPUs --- train_gpt2.cu | 36 +++++++++++++++++++++++++++++++++++- 1 file changed, 35 insertions(+), 1 deletion(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index 24f6c4f..e2da136 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2899,6 +2899,7 @@ float get_flops_promised(const char* device) { }; // basic data from the nvidia whitepapers + static const PerfData VOLTA = {125.0f, -1, 125.f, -1, -1, -1, 1530.f, 640.f}; static const PerfData AMPERE = {156.f, 312.f, 312.f, 312.f, -1, -1, 1410.f, 432.f}; static const PerfData HOPPER = {378.f, 756.f, 756.f, 756.f, 1513.f, 1513.f, 1620.f, 456.f}; static const PerfData ADA = {82.6f, 165.2f, 165.2f, 330.3f, 330.3f, 660.6f, 2520.f, 512.f}; @@ -2925,9 +2926,42 @@ float get_flops_promised(const char* device) { // TODO fill in more GPUs static std::pair gpu_db[] = { + {"Tesla V100-SXM2-16GB", adjust(VOLTA, 640, 1530)}, + {"Tesla V100-PCIE-32GB", adjust(VOLTA, 640, 1530)}, {"NVIDIA A100-PCIE-40GB", adjust(AMPERE, 432, 1410)}, - {"NVIDIA GeForce RTX 4060 Ti", adjust(ADA, 136, 2535)}, + {"NVIDIA A100-PCIE-80GB", adjust(AMPERE, 432, 1410)}, + {"NVIDIA A100-SXM4-40GB", adjust(AMPERE, 432, 1410)}, + {"NVIDIA A100-SXM4-80GB", adjust(AMPERE, 432, 1410)}, + {"NVIDIA RTX A2000", adjust(AMPERE, 104, 1200)}, {"NVIDIA RTX A4000", adjust(AMPERE, 192, 1560)}, + {"NVIDIA RTX A4500", adjust(AMPERE, 224, 1650)}, + {"NVIDIA RTX A5000", adjust(AMPERE, 256, 1695)}, + {"NVIDIA RTX A5500", adjust(AMPERE, 320, 1770)}, + {"NVIDIA RTX A6000", adjust(AMPERE, 336, 1800)}, + {"NVIDIA GeForce RTX 3090 Ti", adjust(AMPERE, 336, 1860)}, + {"NVIDIA GeForce RTX 3090", adjust(AMPERE, 328, 1695)}, + {"NVIDIA GeForce RTX 3080 Ti", adjust(AMPERE, 320, 1665)}, + {"NVIDIA GeForce RTX 3080", adjust(AMPERE, 272, 1710)}, + {"NVIDIA GeForce RTX 3070 Ti", adjust(AMPERE, 192, 1770)}, + {"NVIDIA GeForce RTX 3070", adjust(AMPERE, 184, 1725)}, + {"NVIDIA GeForce RTX 3060 Ti", adjust(AMPERE, 152, 1665)}, + {"NVIDIA GeForce RTX 3060", adjust(AMPERE, 112, 1777)}, + {"NVIDIA RTX A2000 ADA", adjust(ADA, 88, 2130)}, + {"NVIDIA RTX A4000 ADA", adjust(ADA, 192, 2175)}, + {"NVIDIA RTX A4500 ADA", adjust(ADA, 224, 2580)}, + {"NVIDIA RTX A5000 ADA", adjust(ADA, 400, 2550)}, + {"NVIDIA RTX A5880 ADA", adjust(ADA, 440, 2460)}, + {"NVIDIA RTX A6000 ADA", adjust(ADA, 568, 2505)}, + {"NVIDIA GeForce RTX 4090", adjust(ADA, 512, 2520)}, + {"NVIDIA GeForce RTX 4080 SUPER", adjust(ADA, 320, 2550)}, + {"NVIDIA GeForce RTX 4080", adjust(ADA, 304, 2505)}, + {"NVIDIA GeForce RTX 4070 Ti SUPER", adjust(ADA, 264, 2610)}, + {"NVIDIA GeForce RTX 4070 Ti", adjust(ADA, 240, 2610)}, + {"NVIDIA GeForce RTX 4070 SUPER", adjust(ADA, 224, 2475)}, + {"NVIDIA GeForce RTX 4070", adjust(ADA, 184, 2475)}, + {"NVIDIA GeForce RTX 4070", adjust(ADA, 184, 2475)}, + {"NVIDIA GeForce RTX 4060 Ti", adjust(ADA, 136, 2535)}, + {"NVIDIA GeForce RTX 4060", adjust(ADA, 96, 2460)}, {"NVIDIA H100 80GB HBM3", adjust(HOPPER, 528, 1830)}, // HBM3 = SXM5 }; From daa95e04ae44bd77cbb69660756eb0e9284543b8 Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Wed, 29 May 2024 02:27:11 +0300 Subject: [PATCH 05/20] ampere-consumer --- train_gpt2.cu | 39 ++++++++++++++++++++------------------- 1 file changed, 20 insertions(+), 19 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index e2da136..97509e7 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2900,7 +2900,8 @@ float get_flops_promised(const char* device) { // basic data from the nvidia whitepapers static const PerfData VOLTA = {125.0f, -1, 125.f, -1, -1, -1, 1530.f, 640.f}; - static const PerfData AMPERE = {156.f, 312.f, 312.f, 312.f, -1, -1, 1410.f, 432.f}; + static const PerfData AMPERE_DATACENTER = {156.f, 312.f, 312.f, 312.f, -1, -1, 1410.f, 432.f}; + static const PerfData AMPERE_CONSUMER = {40.f, 80.f, 80.f, 160.f, -1, -1, 1860.f, 336.f}; static const PerfData HOPPER = {378.f, 756.f, 756.f, 756.f, 1513.f, 1513.f, 1620.f, 456.f}; static const PerfData ADA = {82.6f, 165.2f, 165.2f, 330.3f, 330.3f, 660.6f, 2520.f, 512.f}; @@ -2928,24 +2929,24 @@ float get_flops_promised(const char* device) { static std::pair gpu_db[] = { {"Tesla V100-SXM2-16GB", adjust(VOLTA, 640, 1530)}, {"Tesla V100-PCIE-32GB", adjust(VOLTA, 640, 1530)}, - {"NVIDIA A100-PCIE-40GB", adjust(AMPERE, 432, 1410)}, - {"NVIDIA A100-PCIE-80GB", adjust(AMPERE, 432, 1410)}, - {"NVIDIA A100-SXM4-40GB", adjust(AMPERE, 432, 1410)}, - {"NVIDIA A100-SXM4-80GB", adjust(AMPERE, 432, 1410)}, - {"NVIDIA RTX A2000", adjust(AMPERE, 104, 1200)}, - {"NVIDIA RTX A4000", adjust(AMPERE, 192, 1560)}, - {"NVIDIA RTX A4500", adjust(AMPERE, 224, 1650)}, - {"NVIDIA RTX A5000", adjust(AMPERE, 256, 1695)}, - {"NVIDIA RTX A5500", adjust(AMPERE, 320, 1770)}, - {"NVIDIA RTX A6000", adjust(AMPERE, 336, 1800)}, - {"NVIDIA GeForce RTX 3090 Ti", adjust(AMPERE, 336, 1860)}, - {"NVIDIA GeForce RTX 3090", adjust(AMPERE, 328, 1695)}, - {"NVIDIA GeForce RTX 3080 Ti", adjust(AMPERE, 320, 1665)}, - {"NVIDIA GeForce RTX 3080", adjust(AMPERE, 272, 1710)}, - {"NVIDIA GeForce RTX 3070 Ti", adjust(AMPERE, 192, 1770)}, - {"NVIDIA GeForce RTX 3070", adjust(AMPERE, 184, 1725)}, - {"NVIDIA GeForce RTX 3060 Ti", adjust(AMPERE, 152, 1665)}, - {"NVIDIA GeForce RTX 3060", adjust(AMPERE, 112, 1777)}, + {"NVIDIA A100-PCIE-40GB", adjust(AMPERE_DATACENTER, 432, 1410)}, + {"NVIDIA A100-PCIE-80GB", adjust(AMPERE_DATACENTER, 432, 1410)}, + {"NVIDIA A100-SXM4-40GB", adjust(AMPERE_DATACENTER, 432, 1410)}, + {"NVIDIA A100-SXM4-80GB", adjust(AMPERE_DATACENTER, 432, 1410)}, + {"NVIDIA RTX A2000", adjust(AMPERE_CONSUMER, 104, 1200)}, + {"NVIDIA RTX A4000", adjust(AMPERE_CONSUMER, 192, 1560)}, + {"NVIDIA RTX A4500", adjust(AMPERE_CONSUMER, 224, 1650)}, + {"NVIDIA RTX A5000", adjust(AMPERE_CONSUMER, 256, 1695)}, + {"NVIDIA RTX A5500", adjust(AMPERE_CONSUMER, 320, 1770)}, + {"NVIDIA RTX A6000", adjust(AMPERE_CONSUMER, 336, 1800)}, + {"NVIDIA GeForce RTX 3090 Ti", adjust(AMPERE_CONSUMER, 336, 1860)}, + {"NVIDIA GeForce RTX 3090", adjust(AMPERE_CONSUMER, 328, 1695)}, + {"NVIDIA GeForce RTX 3080 Ti", adjust(AMPERE_CONSUMER, 320, 1665)}, + {"NVIDIA GeForce RTX 3080", adjust(AMPERE_CONSUMER, 272, 1710)}, + {"NVIDIA GeForce RTX 3070 Ti", adjust(AMPERE_CONSUMER, 192, 1770)}, + {"NVIDIA GeForce RTX 3070", adjust(AMPERE_CONSUMER, 184, 1725)}, + {"NVIDIA GeForce RTX 3060 Ti", adjust(AMPERE_CONSUMER, 152, 1665)}, + {"NVIDIA GeForce RTX 3060", adjust(AMPERE_CONSUMER, 112, 1777)}, {"NVIDIA RTX A2000 ADA", adjust(ADA, 88, 2130)}, {"NVIDIA RTX A4000 ADA", adjust(ADA, 192, 2175)}, {"NVIDIA RTX A4500 ADA", adjust(ADA, 224, 2580)}, From 545063223720505e4da4ef1bdea176d8006bf4f6 Mon Sep 17 00:00:00 2001 From: Christopher Date: Thu, 30 May 2024 03:01:52 +0000 Subject: [PATCH 06/20] Removed unnecesary shared memory due to blockreduce using static defined shared memory --- dev/cuda/classifier_fused.cu | 2 +- train_gpt2.cu | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/dev/cuda/classifier_fused.cu b/dev/cuda/classifier_fused.cu index 2125b87..5b6c986 100644 --- a/dev/cuda/classifier_fused.cu +++ b/dev/cuda/classifier_fused.cu @@ -664,7 +664,7 @@ void fused_classifier5(float* dlogits, float* losses, int B, int T, int V, int P, int block_size) { const int N = B * T; const int grid_size = N; - fused_classifier_kernel5<<>>((floatX*)dlogits, (floatX*)losses, NULL, (floatX*)logits, (floatX*)dlosses, targets, B, T, V, P); + fused_classifier_kernel5<<>>((floatX*)dlogits, (floatX*)losses, NULL, (floatX*)logits, (floatX*)dlosses, targets, B, T, V, P); cudaCheck(cudaGetLastError()); } diff --git a/train_gpt2.cu b/train_gpt2.cu index 5918e21..053c2e3 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -1899,7 +1899,7 @@ void fused_classifier(Type* logits, Type* losses, const int block_size = 1024; const int N = B * T; const int grid_size = N; - fused_classifier_kernel5<<>>(logits, losses, (floatX*)NULL, dloss, targets, B, T, V, P); + fused_classifier_kernel5<<>>(logits, losses, (floatX*)NULL, dloss, targets, B, T, V, P); cudaCheck(cudaGetLastError()); } From cc4bed08067ef2d2413e3c3e28c54636d6e90f8e Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Sat, 1 Jun 2024 14:30:40 +0300 Subject: [PATCH 07/20] added missing checks --- train_gpt2.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index ff4d9c2..6919dbf 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2885,8 +2885,8 @@ void gpt2_free(GPT2 *model) { cudaCheck(cudaFree(model->grads_acts_memory)); cudaCheck(cudaFree(model->inputs)); cudaCheck(cudaFree(model->targets)); - cudaFreeHost(model->cpu_losses); - cudaFreeHost(model->cpu_losses_fp32); + cudaCheck(cudaFreeHost(model->cpu_losses)); + cudaCheck(cudaFreeHost(model->cpu_losses_fp32)); free(model->workload_indices); free(model->bucket_info); } From 761fb2cf2b287b15f6810da4e92273463dc53851 Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Sat, 1 Jun 2024 14:31:26 +0300 Subject: [PATCH 08/20] rename gelu_backward to gelu_backward_inplace, and made it obvious in the code that this is an inplace operation --- train_gpt2.cu | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index 6919dbf..3c21491 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -942,12 +942,12 @@ __global__ void gelu_forward_kernel2(floatX* out, const floatX* inp) { store128(out + idx, packed_out); } -__global__ void gelu_backward_kernel(floatX* dinp, const floatX* inp, const floatX* dout) { +__global__ void gelu_backward_inplace_kernel(floatX* d_in_out, const floatX* inp) { int idx = (blockIdx.x * blockDim.x + threadIdx.x) * x128::size; x128 packed_dinp; x128 packed_inp = load128cs(inp + idx); - x128 packed_dout = load128cs(dout + idx); + x128 packed_dout = load128(d_in_out + idx); for (int k = 0; k < packed_inp.size; ++k) { float x = (float)packed_inp[k]; float cube = 0.044715f * x * x * x; @@ -958,7 +958,7 @@ __global__ void gelu_backward_kernel(floatX* dinp, const floatX* inp, const floa float local_grad = 0.5f * (1.0f + tanh_out) + x * 0.5f * sech_out * GELU_SCALING_FACTOR * (1.0f + 3.0f * 0.044715f * x * x); packed_dinp[k] = (floatX)(local_grad * (float)packed_dout[k]); } - store128(dinp + idx, packed_dinp); + store128(d_in_out + idx, packed_dinp); } template @@ -1747,12 +1747,12 @@ void gelu_forward(floatX* out, const floatX* inp, int N) { cudaCheck(cudaGetLastError()); } -void gelu_backward(floatX* dinp, const floatX* inp, const floatX* dout, const int N) { +void gelu_backward_inplace(floatX* d_in_out, const floatX* inp, const int N) { NVTX_RANGE_FN(); const int block_size = 128; assert(N % block_size == 0); const int grid_size = CEIL_DIV(N, block_size * x128::size); - gelu_backward_kernel<<>>(dinp, inp, dout); + gelu_backward_inplace_kernel<<>>(d_in_out, inp); cudaCheck(cudaGetLastError()); } @@ -2660,7 +2660,7 @@ void gpt2_backward(GPT2 *model, int* inputs) { gelu_forward(l_fch_gelu, l_fch, B*T*4*C); } matmul_backward(dl_bt4c, dl_fcprojw, dl_fcprojb, dresidual, l_fch_gelu, l_fcprojw, scratchF, B, T, 4*C, C); - gelu_backward(dl_bt4c, l_fch, dl_bt4c, B*T*4*C); + gelu_backward_inplace(dl_bt4c, l_fch, B*T*4*C); if(model->recompute >= 2) { // same as gelu above, l_ln1 and l_ln2 are just buffers if recompute >= 2, recompute them here on demand layernorm_forward(l_ln2, l_ln2_mean, l_ln2_rstd, l_residual2, l_ln2w, l_ln2b, B, T, C); From 8eafd40615ce2c8a9575b84803a37161f31802cd Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sat, 1 Jun 2024 16:46:32 +0200 Subject: [PATCH 09/20] Update compile cmd in the dev/cuda README --- dev/cuda/README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dev/cuda/README.md b/dev/cuda/README.md index d020cf6..22ad4d0 100644 --- a/dev/cuda/README.md +++ b/dev/cuda/README.md @@ -7,7 +7,7 @@ See the top of each file for how to compile and run the kernel. Alternatively, t For example, we can look at the top of `layernorm_forward.cu` to build the forward pass kernels for the LayerNorm: ```bash -nvcc -O3 --use_fast_math layernorm_forward.cu -o layernorm_forward +nvcc -O3 --use_fast_math -lcublas -lcublasLt layernorm_forward.cu -o layernorm_forward ``` or simply From 290c00a3626f6ffe60deebb4191d9b9edc371efd Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sat, 1 Jun 2024 18:27:24 +0200 Subject: [PATCH 10/20] Remove redundant CPU computation --- dev/cuda/encoder_backward.cu | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/dev/cuda/encoder_backward.cu b/dev/cuda/encoder_backward.cu index 5322187..7c14d0e 100644 --- a/dev/cuda/encoder_backward.cu +++ b/dev/cuda/encoder_backward.cu @@ -163,14 +163,17 @@ int main(int argc, char **argv) { } printf("Using kernel %d\n", kernel_num); - // set up block sizes + // first check the correctness of the kernel + encoder_backward_cpu(dwte, dwpe, dout, inp, B, T, C); + + // time the kernel at different block sizes int block_sizes[] = {32, 64, 128, 256, 512, 1024}; - // first check the correctness of the kernel for (int j = 0; j < sizeof(block_sizes) / sizeof(int); j++) { int block_size = block_sizes[j]; + cudaCheck(cudaMemset(d_dwte, 0, V * C * sizeof(float))); + cudaCheck(cudaMemset(d_dwpe, 0, T * C * sizeof(float))); printf("Checking block size %d.\n", block_size); - encoder_backward_cpu(dwte, dwpe, dout, inp, B, T, C); encoder_backward(kernel_num, d_dwte, d_dwpe, d_dout, d_inp, B, T, C, block_size); validate_result(d_dwte, dwte, "dwte", V * C, 1e-5f); validate_result(d_dwpe, dwpe, "dwpe", T * C, 1e-5f); From a902323f365c0254f3166adeaeeb52985c718cc5 Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Sun, 2 Jun 2024 12:40:28 +0300 Subject: [PATCH 11/20] fix compilation with older nvcc --- train_gpt2.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index bf6acf1..a57d801 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -228,10 +228,10 @@ struct alignas(16) Packed128 { return result; } __device__ static Packed128 zeros() { - return constant(0); + return constant(0.f); } __device__ static Packed128 ones() { - return constant(1); + return constant(1.f); } __device__ ElementType& operator[](int index) { From ae9c95727369eb1cb4ea4a526f7ced840149642f Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sun, 2 Jun 2024 15:04:45 +0200 Subject: [PATCH 12/20] Add clarification on Box-Muller --- llmc/rand.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/llmc/rand.h b/llmc/rand.h index ba13de9..60ed393 100644 --- a/llmc/rand.h +++ b/llmc/rand.h @@ -163,8 +163,8 @@ void uniform_(float* data, unsigned int numel, float from, float to, mt19937_sta } } -// Box�Muller transform - +// Box-Muller transform: maps uniform random numbers to Gaussian distributed numbers +// https://en.wikipedia.org/wiki/Box%E2%80%93Muller_transform void normal_fill_16(float* data, float mean, float std, mt19937_state* state) { #define EPSILONE 1e-12 for (unsigned int t = 0; t < 8; t++) { From 79cbc03ccf23b411c37f603665a6eb911d83abf3 Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sun, 2 Jun 2024 15:06:06 +0200 Subject: [PATCH 13/20] Use local params for num blocks --- train_gpt2.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index bf6acf1..4f762cb 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2830,7 +2830,7 @@ float gpt2_update(GPT2 *model, float learning_rate, float beta1, float beta2, fl // - the position embeddings actively participate at every forward/backward pass float wd = (i == 0 || i == 1 || i == 4 || i == 6 || i == 10 || i == 12) ? weight_decay : 0.0f; // ok finally call the kernel - size_t num_blocks = CEIL_DIV(num_parameters, block_size); + size_t num_blocks = CEIL_DIV(local_params, block_size); adamw_kernel3<<>>(params_ptr, master_ptr, grad_ptr, m_ptr, v_ptr, local_params, learning_rate, beta1, beta2, beta1_correction, beta2_correction, From 30d81cce153bf0143398a7a813c4eac5bbca0aa1 Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sun, 2 Jun 2024 21:28:42 +0200 Subject: [PATCH 14/20] Fix mem leak --- train_gpt2.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index bf6acf1..2ee029f 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2554,7 +2554,7 @@ void gpt2_backward(GPT2 *model, int* inputs) { model->grads_memory = malloc_and_point_parameters(&model->grads, model->param_elements, model->param_sizeof); // we're going to be clever for the activations backward pass. we don't need to exactly // mirror the forward pass activations and we will save memory. - size_t bw_act_sizes[NUM_ACTIVATION_TENSORS]; + size_t bw_act_sizes[NUM_BACKWARD_TENSORS]; fill_in_grad_act_sizes(bw_act_sizes, model->batch_size, model->seq_len, model->config); // count up and allocate the space model->num_grad_acts = 0; @@ -2985,6 +2985,7 @@ void load_state(int* step, GPT2* model, DataLoader* loader, const char* filename cudaCheck(cudaMemcpy(model->m_memory, cpu_buffer, shard_num_parameters * sizeof(float), cudaMemcpyHostToDevice)); freadCheck(cpu_buffer, sizeof(float), shard_num_parameters, state_file); cudaCheck(cudaMemcpy(model->v_memory, cpu_buffer, shard_num_parameters * sizeof(float), cudaMemcpyHostToDevice)); + free(cpu_buffer); } // ---------------------------------------------------------------------------- From df0a826ef3e846aaaa8efe21079c0c361e4df10b Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sun, 2 Jun 2024 21:37:52 +0200 Subject: [PATCH 15/20] Close file --- train_gpt2.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/train_gpt2.cu b/train_gpt2.cu index 2ee029f..7050c67 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2986,6 +2986,7 @@ void load_state(int* step, GPT2* model, DataLoader* loader, const char* filename freadCheck(cpu_buffer, sizeof(float), shard_num_parameters, state_file); cudaCheck(cudaMemcpy(model->v_memory, cpu_buffer, shard_num_parameters * sizeof(float), cudaMemcpyHostToDevice)); free(cpu_buffer); + fclose(state_file); } // ---------------------------------------------------------------------------- From b4e8bd7f661f31508b3de29f05dc912046c73a7b Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sun, 2 Jun 2024 22:45:26 +0200 Subject: [PATCH 16/20] Fix zero stage 1 norm computation - WIP --- train_gpt2.cu | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index bf6acf1..940c9a6 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2763,10 +2763,21 @@ float gpt2_update(GPT2 *model, float learning_rate, float beta1, float beta2, fl // gradient clipping // repurposing this buffer (which isn't needed now) to write grad norm into it float* grad_norm_squared = (float*)model->acts.output; - global_norm_squared(grad_norm_squared, (floatX*)model->grads_memory, model->num_parameters); + if (multi_gpu_config->zero_stage == 1) { + // we only have a local shard of the gradients, so we need to compute the norm of the local shard + global_norm_squared(grad_norm_squared, (floatX*)model->grads_memory + shard_offset, shard_num_parameters); + } else { + // we have all the gradients, so we can compute the norm of all the gradients + global_norm_squared(grad_norm_squared, (floatX*)model->grads_memory, model->num_parameters); + } // transfer the gradient norm to CPU float grad_norm_squared_cpu = 0.0f; cudaCheck(cudaMemcpy(&grad_norm_squared_cpu, grad_norm_squared, sizeof(float), cudaMemcpyDeviceToHost)); + if (multi_gpu_config->zero_stage == 1) { + // we have to sum the gradient norm across all GPUs as they're only partial + grad_norm_squared_cpu = multi_gpu_cpu_float_sum(grad_norm_squared_cpu) + } + if(!isfinite(grad_norm_squared_cpu)) { // may happen due to some issue (e.g. overflow?) // TODO: later may want to keep a global counter of instabilities like this From 9a455edc2599dcadb34e0548b60814a9216b5518 Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sun, 2 Jun 2024 23:03:41 +0200 Subject: [PATCH 17/20] Improve comments --- train_gpt2.cu | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index 940c9a6..f610bf4 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2764,17 +2764,17 @@ float gpt2_update(GPT2 *model, float learning_rate, float beta1, float beta2, fl // repurposing this buffer (which isn't needed now) to write grad norm into it float* grad_norm_squared = (float*)model->acts.output; if (multi_gpu_config->zero_stage == 1) { - // we only have a local shard of the gradients, so we need to compute the norm of the local shard - global_norm_squared(grad_norm_squared, (floatX*)model->grads_memory + shard_offset, shard_num_parameters); + // only the local grad shard has correct values -> compute the squared norm of the local grad shard + global_norm_squared(grad_norm_squared, grads_memory + shard_offset, shard_num_parameters); } else { - // we have all the gradients, so we can compute the norm of all the gradients - global_norm_squared(grad_norm_squared, (floatX*)model->grads_memory, model->num_parameters); + // all grad shards have correct values -> compute the squared norm of the whole grad vector + global_norm_squared(grad_norm_squared, grads_memory, model->num_parameters); } // transfer the gradient norm to CPU float grad_norm_squared_cpu = 0.0f; cudaCheck(cudaMemcpy(&grad_norm_squared_cpu, grad_norm_squared, sizeof(float), cudaMemcpyDeviceToHost)); if (multi_gpu_config->zero_stage == 1) { - // we have to sum the gradient norm across all GPUs as they're only partial + // sum the squared norm across all GPUs as they're only partial (computed per shard) grad_norm_squared_cpu = multi_gpu_cpu_float_sum(grad_norm_squared_cpu) } From a8a5a444bc0da711899010c867cb3691f235f02b Mon Sep 17 00:00:00 2001 From: Aleksa Gordic Date: Sun, 2 Jun 2024 23:04:47 +0200 Subject: [PATCH 18/20] Add ; - ah c my old friend :) --- train_gpt2.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index f610bf4..c554ba8 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2775,7 +2775,7 @@ float gpt2_update(GPT2 *model, float learning_rate, float beta1, float beta2, fl cudaCheck(cudaMemcpy(&grad_norm_squared_cpu, grad_norm_squared, sizeof(float), cudaMemcpyDeviceToHost)); if (multi_gpu_config->zero_stage == 1) { // sum the squared norm across all GPUs as they're only partial (computed per shard) - grad_norm_squared_cpu = multi_gpu_cpu_float_sum(grad_norm_squared_cpu) + grad_norm_squared_cpu = multi_gpu_cpu_float_sum(grad_norm_squared_cpu); } if(!isfinite(grad_norm_squared_cpu)) { From 4944ac574381fd40e0ca13f0841988eade7b7238 Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Sun, 2 Jun 2024 23:28:13 +0000 Subject: [PATCH 19/20] adjust comments --- train_gpt2.cu | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index c554ba8..113333e 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2764,17 +2764,20 @@ float gpt2_update(GPT2 *model, float learning_rate, float beta1, float beta2, fl // repurposing this buffer (which isn't needed now) to write grad norm into it float* grad_norm_squared = (float*)model->acts.output; if (multi_gpu_config->zero_stage == 1) { - // only the local grad shard has correct values -> compute the squared norm of the local grad shard + // ^1 because of the ncclReduceScatter() in gpt2_multi_gpu_accumulate, + // grads_memory only contains the averaged gradients at the local shard + // so we only calculate the grad norm at the grads_memory belonging to the local shard global_norm_squared(grad_norm_squared, grads_memory + shard_offset, shard_num_parameters); } else { - // all grad shards have correct values -> compute the squared norm of the whole grad vector + // the ncclAllReduce() in gpt2_multi_gpu_accumulate has averaged the gradients across all GPUs + // so each GPU can compute the squared norm over the whole grad vector, with no added comms needed global_norm_squared(grad_norm_squared, grads_memory, model->num_parameters); } // transfer the gradient norm to CPU float grad_norm_squared_cpu = 0.0f; cudaCheck(cudaMemcpy(&grad_norm_squared_cpu, grad_norm_squared, sizeof(float), cudaMemcpyDeviceToHost)); if (multi_gpu_config->zero_stage == 1) { - // sum the squared norm across all GPUs as they're only partial (computed per shard) + // further sum the (partial) squared norm across all GPUs (see comment ^1 above) grad_norm_squared_cpu = multi_gpu_cpu_float_sum(grad_norm_squared_cpu); } From 08fc45b4926dbe5c5970307cc4ae5ab064e4dad6 Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Mon, 3 Jun 2024 02:08:29 +0000 Subject: [PATCH 20/20] delete copy paste code for gpt2 model init --- train_gpt2.cu | 72 ++++++++++++++++++++++++--------------------------- 1 file changed, 34 insertions(+), 38 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index 1e6dfa3..06da2f5 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2151,13 +2151,43 @@ typedef struct { floatX* cpu_losses; // CPU buffer to copy the losses to, allocated with cudaMallocHost float* cpu_losses_fp32; // same but fp32 unsigned long long rng_state; // the RNG state for seeding stochastic rounding etc. - int use_master_weights; - int recompute; + int use_master_weights; // keep master weights copy in float for optim update? 0|1 + int recompute; // recompute gelu | layernorm forward during model backward? 0|1|2 // todo - if other functions need cpu scratch buffers in the future, reuse as generic scratch? int* workload_indices; // encoder_backward, B*T*num_c_groups (int) int4* bucket_info; // encoder_backward, B*T*num_c_groups (int4) - size for worst case } GPT2; +void gpt2_init_common(GPT2 *model) { + // common inits outside of the model weights + // the weights are initialized either in: + // - gpt2_build_from_checkpoint() if loading from a checkpoint + // - gpt2_build_from_random() if starting from scratch + // memory lazily initialized in forward() + model->acts_memory = NULL; + model->inputs = NULL; + model->targets = NULL; + model->cpu_losses = NULL; + model->cpu_losses_fp32 = NULL; + // the B,T params are determined and set, fixed on first batch in forward() + model->batch_size = 0; + model->seq_len = 0; + model->mean_loss = -1.0f; // -1.0f designates no loss, set at end of forward() + // memory lazily initialized in backward() + model->grads_memory = NULL; + model->grads_acts_memory = NULL; + model->workload_indices = NULL; // on cpu, for encoder_backward + model->bucket_info = NULL; // on cpu, for encoder_backward + // memory lazily initialized in update() + model->m_memory = NULL; + model->v_memory = NULL; + model->master_weights = NULL; + // other default settings + model->rng_state = 13371337; // used in stochastic rounding + model->use_master_weights = 1; // safe default: do keep master weights in fp32 + model->recompute = 1; // good default: recompute gelu but not layernorm +} + void gpt2_write_to_checkpoint(GPT2 *model, const char* checkpoint_path) { // write the model to a checkpoint file printf0("Writing model to %s\n", checkpoint_path); @@ -2247,25 +2277,7 @@ void gpt2_build_from_checkpoint(GPT2 *model, const char* checkpoint_path) { free(params_memory_cpu); fcloseCheck(model_file); - // other inits - model->acts_memory = NULL; - model->grads_memory = NULL; - model->m_memory = NULL; - model->v_memory = NULL; - model->master_weights = NULL; - model->grads_acts_memory = NULL; - model->inputs = NULL; - model->targets = NULL; - model->cpu_losses = NULL; - model->cpu_losses_fp32 = NULL; - model->workload_indices = NULL; - model->bucket_info = NULL; - model->batch_size = 0; - model->seq_len = 0; - model->mean_loss = -1.0f; // -1.0f will designate no loss - model->rng_state = 13371337; - model->use_master_weights = 1; // keep master weights copy in float for optim update? - model->recompute = 1; // default to recompute gelu during backward + gpt2_init_common(model); } void gpt2_build_from_random(GPT2 *model, int depth) { @@ -2354,23 +2366,7 @@ void gpt2_build_from_random(GPT2 *model, int depth) { cudaCheck(cudaMemcpy(model->params_memory, params_memory_cpu, model->num_parameters_bytes, cudaMemcpyHostToDevice)); free(params_memory_cpu); - // other inits and defaults - model->acts_memory = NULL; - model->grads_memory = NULL; - model->m_memory = NULL; - model->v_memory = NULL; - model->master_weights = NULL; - model->grads_acts_memory = NULL; - model->inputs = NULL; - model->targets = NULL; - model->cpu_losses = NULL; - model->cpu_losses_fp32 = NULL; - model->batch_size = 0; - model->seq_len = 0; - model->mean_loss = -1.0f; // -1.0f designates no loss - model->rng_state = 13371337; - model->use_master_weights = 1; // keep master weights copy in float for optim update? - model->recompute = 1; // default to recompute gelu during backward + gpt2_init_common(model); } void gpt2_forward(GPT2 *model, int* inputs, int* targets, size_t B, size_t T, int grad_accum_steps=1) {