From 568615fef11beee6bd3391be407c2d6e0b3dcfcc Mon Sep 17 00:00:00 2001 From: Christopher Dryden Date: Tue, 30 Apr 2024 03:21:42 +0000 Subject: [PATCH 01/16] Added packing for gelu forwards kernel --- dev/cuda/gelu_forward.cu | 72 +++++++++++++++++++++++++++++----------- train_gpt2.cu | 21 ++++++++---- 2 files changed, 67 insertions(+), 26 deletions(-) diff --git a/dev/cuda/gelu_forward.cu b/dev/cuda/gelu_forward.cu index e69abdc..c9fba9e 100644 --- a/dev/cuda/gelu_forward.cu +++ b/dev/cuda/gelu_forward.cu @@ -21,6 +21,22 @@ version 2 uses the Packed128 data structure #include #include "common.h" +// turn on bf16 as default, done up here for now +//#define ENABLE_BF16 + +#if defined(ENABLE_BF16) +typedef __nv_bfloat16 floatX; +typedef __nv_bfloat16 floatN; +#elif defined(ENABLE_FP16) +typedef half floatX; +typedef half floatN; +#else +typedef float floatX; +typedef float floatN; +#endif + +typedef Packed128 x128; + // ---------------------------------------------------------------------------- // CPU code reference @@ -48,19 +64,19 @@ __global__ void gelu_kernel(float* out, const float* inp, int N) { } // elementwise ops are nice and ez -__global__ void gelu_kernel2(float* out, const float* inp, int N) { - int i = (blockIdx.x * blockDim.x + threadIdx.x) * f128::size; +__global__ void gelu_kernel2(floatX* out, const floatX* inp, int N) { + int i = (blockIdx.x * blockDim.x + threadIdx.x) * x128::size; if (i < N) { - f128 packet_out; - f128 packet_in = load128cs(inp + i); // load and do not keep in cache - for(int k = 0; k < packet_in.size; ++k) { - float xi = packet_in[k]; + x128 packed_out; + x128 packed_inp = load128cs(inp + i); // load and do not keep in cache + for(int k = 0; k < packed_inp.size; ++k) { + float xi = (float)packed_inp[k]; float cube = 0.044715f * xi * xi * xi; - packet_out[k] = 0.5f * xi * (1.0f + tanhf(GELU_SCALING_FACTOR * (xi + cube))); + packed_out[k] = (float)(0.5f * xi * (1.0f + tanhf(GELU_SCALING_FACTOR * (xi + cube)))); } // store instead of storecs (without cache streaming) in case it is useful for the // data to be in the cache for the next operation after this GeLU - store128(out + i, packet_out); + store128(out + i, packed_out); } } @@ -73,25 +89,29 @@ void gelu_forward1(float* out, const float* inp, int N, const int block_size) { cudaCheck(cudaGetLastError()); } -void gelu_forward2(float* out, const float* inp, int N, const int block_size) { - const int grid_size = ceil_div(N, block_size) / 4; +void gelu_forward2(floatX* out, const floatX* inp, int N, const int block_size) { + const int grid_size = ceil_div(N, block_size)/x128::size; gelu_kernel2<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } // kernel version dispatch void gelu_forward(int kernel_num, - float* out, - const float* inp, + floatX* out, + const floatX* inp, int B, int T, int C, int block_size) { switch (kernel_num) { +#if !defined(ENABLE_BF16) && !defined(ENABLE_FP16) case 1: gelu_forward1(out, inp, B * T * C, block_size); break; +#endif +#if defined(ENABLE_BF16) case 2: gelu_forward2(out, inp, B * T * C, block_size); break; +#endif default: printf("Invalid kernel number\n"); exit(1); @@ -114,13 +134,6 @@ int main(int argc, char **argv) { float* out = (float*)malloc(B * T * C * sizeof(float)); float* inp = make_random_float(B * T * C); - // move to GPU - float* d_out; - float* d_inp; - cudaCheck(cudaMalloc(&d_out, B * T * C * sizeof(float))); - cudaCheck(cudaMalloc(&d_inp, B * T * C * sizeof(float))); - cudaCheck(cudaMemcpy(d_inp, inp, B * T * C * sizeof(float), cudaMemcpyHostToDevice)); - // read kernel_num from command line int kernel_num = 1; if (argc > 1) { @@ -131,6 +144,19 @@ int main(int argc, char **argv) { // first check the correctness of the kernel gelu_forward_cpu(out, inp, B * T * C); + // move to GPU + floatX* d_out; + floatX* d_inp; + cudaCheck(cudaMalloc(&d_out, B * T * C * sizeof(floatX))); + cudaCheck(cudaMalloc(&d_inp, B * T * C * sizeof(floatX))); + + floatX* inpX = (floatX*)malloc(B * T * C * sizeof(floatX)); + + for (int i = 0; i < B * T * C; i++) { + inpX[i] = (floatX)inp[i]; + } + + cudaCheck(cudaMemcpy(d_inp, inpX, B * T * C * sizeof(floatX), cudaMemcpyHostToDevice)); // time the kernel at different block sizes int block_sizes[] = {32, 64, 128, 256, 512, 1024}; @@ -138,7 +164,12 @@ int main(int argc, char **argv) { int block_size = block_sizes[j]; printf("Checking block size %d.\n", block_size); gelu_forward(kernel_num, d_out, d_inp, B, T, C, block_size); +#if !defined(ENABLE_BF16) && !defined(ENABLE_FP16) validate_result(d_out, out, "out", B * T * C, 1e-5f); +#endif +#if defined(ENABLE_BF16) +#endif + validate_result(d_out, out, "out", B * T * C, 1e-2f); } printf("All results match. Starting benchmarks.\n\n"); @@ -164,8 +195,9 @@ int main(int argc, char **argv) { // free memory free(out); free(inp); + free(inpX); + cudaCheck(cudaFree(d_out)); cudaCheck(cudaFree(d_inp)); - return 0; } \ No newline at end of file diff --git a/train_gpt2.cu b/train_gpt2.cu index 3913bc5..ce3a8f5 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -227,6 +227,7 @@ struct alignas(16) Packed128 { // short-form typedef typedef Packed128 f128; +typedef Packed128 x128; // load a Packed128 from an aligned memory address template @@ -566,7 +567,8 @@ __global__ void permute_kernel_backward(floatX* dinp, __global__ void unpermute_kernel(floatX* inp, floatX *out, int B, int N, int NH, int d) { // out has shape (B, nh, N, d) but we need to unpermute it to (B, N, nh, d) - int idx = blockIdx.x * blockDim.x + threadIdx.x; + + int idx = (blockIdx.x * blockDim.x + threadIdx.x); // out[b][n][nh_][d_] <- inp[b][nh_][n][d_] if (idx < B * NH * N * d) { int b = idx / (NH * N * d); @@ -668,11 +670,18 @@ __global__ void residual_forward_kernel(floatX* out, floatX* inp1, floatX* inp2, #define GELU_SCALING_FACTOR sqrtf(2.0f / M_PI) __global__ void gelu_forward_kernel(floatX* out, const floatX* inp, int N) { - int i = blockIdx.x * blockDim.x + threadIdx.x; + int i = (blockIdx.x * blockDim.x + threadIdx.x) * x128::size; if (i < N) { - float xi = (float)inp[i]; - float cube = 0.044715f * xi * xi * xi; - out[i] = (floatX)(0.5f * xi * (1.0f + tanhf(GELU_SCALING_FACTOR * (xi + cube)))); + x128 packed_out; + x128 packed_inp = load128cs(inp + i); // load and do not keep in cache + for(int k = 0; k < packed_inp.size; ++k) { + float xi =(float)packed_inp[k]; + float cube = 0.044715f * xi * xi * xi; + packed_out[k] = (floatX)(0.5f * xi * (1.0f + tanhf(GELU_SCALING_FACTOR * (xi + cube)))); + } + // store instead of storecs (without cache streaming) in case it is useful for the + // data to be in the cache for the next operation after this GeLU + store128(out + i, packed_out); } } @@ -1192,7 +1201,7 @@ void residual_forward(floatX* out, floatX* inp1, floatX* inp2, int N) { void gelu_forward(floatX* out, const floatX* inp, int N) { const int block_size = 128; - const int grid_size = CEIL_DIV(N, block_size); + const int grid_size = CEIL_DIV(N/x128::size, block_size); gelu_forward_kernel<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } From d7813d281d790d1061997aceed6c64ad9207747c Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Tue, 30 Apr 2024 12:45:37 +0300 Subject: [PATCH 02/16] clear the L2 cache between consecutive invokations of our microbenchmarks to get reliable results --- dev/cuda/common.h | 32 +++++++++++++++++++++++++------- 1 file changed, 25 insertions(+), 7 deletions(-) diff --git a/dev/cuda/common.h b/dev/cuda/common.h index 7edb1e7..169dac3 100644 --- a/dev/cuda/common.h +++ b/dev/cuda/common.h @@ -148,7 +148,7 @@ void validate_result(D* device_result, const T* cpu_reference, const char* name, printf("%f %f\n", cpu_reference[i], (T)out_gpu[i]); } // ensure correctness for all elements. We can set an "ignore" mask by writing NaN - if (fabs(cpu_reference[i] - (T)out_gpu[i]) > tolerance && !isnan(cpu_reference[i])) { + if (fabs(cpu_reference[i] - (T)out_gpu[i]) > tolerance && isfinite(cpu_reference[i])) { printf("Mismatch of %s at %d: CPU_ref: %f vs GPU: %f\n", name, i, cpu_reference[i], (T)out_gpu[i]); nfaults ++; if (nfaults >= 10) { @@ -169,17 +169,35 @@ void validate_result(D* device_result, const T* cpu_reference, const char* name, template float benchmark_kernel(int repeats, Kernel kernel, KernelArgs&&... kernel_args) { cudaEvent_t start, stop; + // prepare buffer to scrub L2 cache between benchmarks + // just memset a large dummy array, recommended by + // https://stackoverflow.com/questions/31429377/how-can-i-clear-flush-the-l2-cache-and-the-tlb-of-a-gpu + // and apparently used in nvbench. + int deviceIdx = 0; + cudaCheck(cudaSetDevice(deviceIdx)); + cudaDeviceProp deviceProp; + cudaCheck(cudaGetDeviceProperties(&deviceProp, deviceIdx)); + void* flush_buffer; + cudaCheck(cudaMalloc(&flush_buffer, deviceProp.l2CacheSize)); + cudaCheck(cudaEventCreate(&start)); cudaCheck(cudaEventCreate(&stop)); - cudaCheck(cudaEventRecord(start, nullptr)); + float elapsed_time = 0.f; for (int i = 0; i < repeats; i++) { + // clear L2 + cudaCheck(cudaMemset(flush_buffer, 0, deviceProp.l2CacheSize)); + // now we can start recording the timing of the kernel + cudaCheck(cudaEventRecord(start, nullptr)); kernel(std::forward(kernel_args)...); + cudaCheck(cudaEventRecord(stop, nullptr)); + cudaCheck(cudaEventSynchronize(start)); + cudaCheck(cudaEventSynchronize(stop)); + float single_call; + cudaCheck(cudaEventElapsedTime(&single_call, start, stop)); + elapsed_time += single_call; } - cudaCheck(cudaEventRecord(stop, nullptr)); - cudaCheck(cudaEventSynchronize(start)); - cudaCheck(cudaEventSynchronize(stop)); - float elapsed_time; - cudaCheck(cudaEventElapsedTime(&elapsed_time, start, stop)); + + cudaCheck(cudaFree(flush_buffer)); return elapsed_time / repeats; } \ No newline at end of file From 51face88d9f651ce05efde81cebca04c9576575e Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 15:48:40 +0000 Subject: [PATCH 03/16] fix bug where backward/step must be outside of amp context --- train_gpt2.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/train_gpt2.py b/train_gpt2.py index 62ad913..7cf3166 100644 --- a/train_gpt2.py +++ b/train_gpt2.py @@ -544,10 +544,10 @@ if __name__ == "__main__": with ctx: logits, loss = model(x, y) del logits - if not args.inference_only: - optimizer.zero_grad(set_to_none=True) - loss.backward() - optimizer.step() + if not args.inference_only: + optimizer.zero_grad(set_to_none=True) + loss.backward() + optimizer.step() # wait on the CPU for all device work to end so we get accurate per-iteration timings below if device == "mps": torch.mps.synchronize() From 9d8a6d13f0309328f4ee5474ac993dcf9a38dd10 Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 16:03:33 +0000 Subject: [PATCH 04/16] add flag overfit_single_batch useful debugging into train_gpt2.cu --- train_gpt2.cu | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index 8451480..e6ccf31 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2087,6 +2087,7 @@ void error_usage() { fprintf(stderr, " -m val_max_batches, up to how many val batches to estimate val loss? (default = 20)\n"); fprintf(stderr, " -s sample_every, how often we inference the model (default = 20)\n"); fprintf(stderr, " -g genT, how many steps of inference we do (default = 64)\n"); + fprintf(stderr, " -a overfit a single batch? 0/1. useful for debugging\n"); exit(EXIT_FAILURE); } @@ -2105,6 +2106,7 @@ int main(int argc, char *argv[]) { int val_max_batches = 20; // how many batches max do we eval for validation loss? int sample_every = 20; // every how many steps to do inference? int genT = 64; // number of steps of inference we will do + int overfit_single_batch = 0; // useful for debugging, 1 = only load a single data batch once for (int i = 1; i < argc; i+=2) { if (i + 1 >= argc) { error_usage(); } // must have arg after flag if (argv[i][0] != '-') { error_usage(); } // must start with dash @@ -2119,6 +2121,7 @@ int main(int argc, char *argv[]) { else if (argv[i][1] == 'm') { val_max_batches = atoi(argv[i+1]); } else if (argv[i][1] == 's') { sample_every = atoi(argv[i+1]); } else if (argv[i][1] == 'g') { genT = atoi(argv[i+1]); } + else if (argv[i][1] == 'a') { overfit_single_batch = atoi(argv[i+1]); } else { error_usage(); } } printf0("+-----------------------+----------------------------------------------------+\n"); @@ -2133,6 +2136,7 @@ int main(int argc, char *argv[]) { printf0("| val_max_batches | %-50d |\n", val_max_batches); printf0("| sample_every | %-50d |\n", sample_every); printf0("| genT | %-50d |\n", genT); + printf0("| overfit_single_batch | %-50d |\n", overfit_single_batch); printf0("+-----------------------+----------------------------------------------------+\n"); // set up the device @@ -2283,7 +2287,9 @@ int main(int argc, char *argv[]) { // do a training step clock_gettime(CLOCK_MONOTONIC, &start); - dataloader_next_batch(&train_loader); + if (overfit_single_batch == 0 || (step == 0 && overfit_single_batch == 1)) { + dataloader_next_batch(&train_loader); + } gpt2_forward(&model, train_loader.inputs, train_loader.targets, B, T); gpt2_zero_grad(&model); gpt2_backward(&model); From b239b67ae1a2f764412ed6f3090a5a0449df0964 Mon Sep 17 00:00:00 2001 From: Franz Louis Cesista Date: Wed, 1 May 2024 00:15:08 +0800 Subject: [PATCH 05/16] add modal script --- dev/cuda/README.md | 6 +++ dev/cuda/benchmark_on_modal.py | 82 ++++++++++++++++++++++++++++++++++ 2 files changed, 88 insertions(+) create mode 100644 dev/cuda/benchmark_on_modal.py diff --git a/dev/cuda/README.md b/dev/cuda/README.md index 041ff00..d020cf6 100644 --- a/dev/cuda/README.md +++ b/dev/cuda/README.md @@ -31,3 +31,9 @@ You'll see that this first forwards the reference code on the CPU, then it runs You'll see that this matches all the CPU results but runs much much faster. The typical process from here on is we copy paste the kernel that ran fastest, adjust it manually (e.g. to hardcode the best block size) and drop it into the training code file, e.g. `train_gpt2.cu`. To add a new version of a kernel, add the kernel to the corresponding file and adjust the docs. To add a new kernel, add the new file and adjust the Makefile. Run `make clean` to clean up binaries from your directory. + +If you do not have a GPU or is having trouble with CUDA dependencies, you can run the benchmarks on the [Modal platform](http://modal.com). For example, to run the benchmark for the attention forward pass on an A100 GPU with 80GB of memory, you can run the following command: + +```bash +GPU_MEM=80 modal run benchmark_on_modal.py --compile-command "nvcc -O3 --use_fast_math attention_forward.cu -o attention_forward -lcublas" --run-command "./attention_forward 1" +``` diff --git a/dev/cuda/benchmark_on_modal.py b/dev/cuda/benchmark_on_modal.py new file mode 100644 index 0000000..e597538 --- /dev/null +++ b/dev/cuda/benchmark_on_modal.py @@ -0,0 +1,82 @@ +""" +Script for running benchmarks on the Modal platform. +This is useful for folks who do not have access to expensive GPUs locally. + +Example usage: +GPU_MEM=80 modal run benchmark_on_modal.py \ + --compile-command "nvcc -O3 --use_fast_math attention_forward.cu -o attention_forward -lcublas" \ + --run-command "./attention_forward 1" + +This will mount the contents of the current directory to the remote container on modal, +compile the `attention_forward.cu` file with `nvcc`, and run the resulting binary on a A100 GPU with 80GB of memory. +""" + +import subprocess +import os +import sys + +import modal +from modal import Image, Stub + +GPU_NAME_TO_MODAL_CLASS_MAP = { + "H100": modal.gpu.H100, + "A100": modal.gpu.A100, + "A10G": modal.gpu.A10G, +} + +N_GPUS = int(os.environ.get("N_GPUS", 1)) +GPU_MEM = int(os.environ.get("GPU_MEM", 40)) +GPU_NAME = os.environ.get("GPU_NAME", "A100") +GPU_CONFIG = GPU_NAME_TO_MODAL_CLASS_MAP[GPU_NAME](count=N_GPUS, memory=GPU_MEM) + +APP_NAME = "llm.c benchmark run" + +# We don't actually need to use the Axolotl image here, but it's reliable +AXOLOTL_REGISTRY_SHA = ( + "d5b941ba2293534c01c23202c8fc459fd2a169871fa5e6c45cb00f363d474b6a" +) +axolotl_image = ( + Image.from_registry(f"winglian/axolotl@sha256:{AXOLOTL_REGISTRY_SHA}") + .run_commands( + "git clone https://github.com/OpenAccess-AI-Collective/axolotl /root/axolotl", + "cd /root/axolotl && git checkout v0.4.0", + ) + .pip_install("huggingface_hub==0.20.3", "hf-transfer==0.1.5") + .env( + dict( + HUGGINGFACE_HUB_CACHE="/pretrained", + HF_HUB_ENABLE_HF_TRANSFER="1", + TQDM_DISABLE="true", + ) + ) +) + +stub = Stub(APP_NAME) + + +def execute_command(command: str): + command_args = command.split(" ") + print(f"{command_args = }") + subprocess.run(command_args, stdout=sys.stdout, stderr=subprocess.STDOUT) + + +@stub.function( + gpu=GPU_CONFIG, + image=axolotl_image, + allow_concurrent_inputs=4, + container_idle_timeout=900, + # This copies everything in this folder to the remote root folder + mounts=[modal.Mount.from_local_dir("./", remote_path="/root/")] +) +def run_benchmark(compile_command: str, run_command: str): + execute_command("pwd") + execute_command("ls") + execute_command(compile_command) + execute_command(run_command) + return None + + +@stub.local_entrypoint() +def inference_main(compile_command: str, run_command: str): + results = run_benchmark.remote(compile_command, run_command) + return results From 654d6f55c5037e07ccff84984f838ed98f71536a Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 17:46:36 +0000 Subject: [PATCH 06/16] add arg to cap the number of steps, and offset all prints to start steps at 1, which i think looks better to the eye --- test_gpt2.cu | 6 +++--- train_gpt2.cu | 6 +++++- train_gpt2.py | 2 +- 3 files changed, 9 insertions(+), 5 deletions(-) diff --git a/test_gpt2.cu b/test_gpt2.cu index 32bea71..e2b540c 100644 --- a/test_gpt2.cu +++ b/test_gpt2.cu @@ -280,7 +280,7 @@ int main(int argc, char *argv[]) { gpt2_update(&model, 1e-4f, 0.9f, 0.999f, 1e-8f, 0.01f, step+1); // print the timing information at the end - printf("step %d: loss %f (took %f ms)\n", step, model.mean_loss, time_elapsed_s * 1000); + printf("step %d: loss %f (took %f ms)\n", step+1, model.mean_loss, time_elapsed_s * 1000); losses[step] = model.mean_loss; } @@ -301,10 +301,10 @@ int main(int argc, char *argv[]) { // compare for (int i = 0; i < 10; i++) { if (fabsf(losses[i] - expected_losses[i]) >= loss_diff_threshold) { - printf("LOSS MISMATCH AT STEP %d: %f %f\n", i, losses[i], expected_losses[i]); + printf("LOSS MISMATCH AT STEP %d: %f %f\n", i+1, losses[i], expected_losses[i]); allok = 0; } else { - printf("loss ok at step %d: %f %f\n", i, losses[i], expected_losses[i]); + printf("loss ok at step %d: %f %f\n", i+1, losses[i], expected_losses[i]); } } diff --git a/train_gpt2.cu b/train_gpt2.cu index e6ccf31..5e6fc80 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2083,6 +2083,7 @@ void error_usage() { fprintf(stderr, " -b batch size B (default = 4)\n"); fprintf(stderr, " -t sequence length T (default = 1024)\n"); fprintf(stderr, " -l learning rate (default = 3e-4f)\n"); + fprintf(stderr, " -x max_steps of optimization to run (-1 (default) = disable, run 1 epoch)\n"); fprintf(stderr, " -v val_loss_every, how often we evaluate val loss (default = 20)\n"); fprintf(stderr, " -m val_max_batches, up to how many val batches to estimate val loss? (default = 20)\n"); fprintf(stderr, " -s sample_every, how often we inference the model (default = 20)\n"); @@ -2107,6 +2108,7 @@ int main(int argc, char *argv[]) { int sample_every = 20; // every how many steps to do inference? int genT = 64; // number of steps of inference we will do int overfit_single_batch = 0; // useful for debugging, 1 = only load a single data batch once + int max_steps = -1; for (int i = 1; i < argc; i+=2) { if (i + 1 >= argc) { error_usage(); } // must have arg after flag if (argv[i][0] != '-') { error_usage(); } // must start with dash @@ -2117,6 +2119,7 @@ int main(int argc, char *argv[]) { else if (argv[i][1] == 'b') { B = atoi(argv[i+1]); } // Per-GPU batch size else if (argv[i][1] == 't') { T = atoi(argv[i+1]); } else if (argv[i][1] == 'l') { learning_rate = atof(argv[i+1]); } + else if (argv[i][1] == 'x') { max_steps = atoi(argv[i+1]); } else if (argv[i][1] == 'v') { val_loss_every = atoi(argv[i+1]); } else if (argv[i][1] == 'm') { val_max_batches = atoi(argv[i+1]); } else if (argv[i][1] == 's') { sample_every = atoi(argv[i+1]); } @@ -2132,6 +2135,7 @@ int main(int argc, char *argv[]) { printf0("| batch size B | %-50d |\n", B); printf0("| sequence length T | %-50d |\n", T); printf0("| learning rate | %-50f |\n", learning_rate); + printf0("| max_steps | %-50d |\n", max_steps); printf0("| val_loss_every | %-50d |\n", val_loss_every); printf0("| val_max_batches | %-50d |\n", val_max_batches); printf0("| sample_every | %-50d |\n", sample_every); @@ -2187,7 +2191,7 @@ int main(int argc, char *argv[]) { dataloader_init(&train_loader, &multi_gpu_config, train_tokens_filename, B, T); DataLoader val_loader; dataloader_init(&val_loader, &multi_gpu_config, val_tokens_filename, B, T); - int train_num_batches = train_loader.num_batches; // let's do 1 epoch by default for now + int train_num_batches = (max_steps == -1) ? train_loader.num_batches : max_steps; // default = 1 epoch int val_num_batches = train_loader.num_batches < val_max_batches ? train_loader.num_batches : val_max_batches; printf0("| train_num_batches | %-50d |\n", train_num_batches); printf0("| val_num_batches | %-50d |\n", val_num_batches); diff --git a/train_gpt2.py b/train_gpt2.py index 7cf3166..abc1eab 100644 --- a/train_gpt2.py +++ b/train_gpt2.py @@ -557,7 +557,7 @@ if __name__ == "__main__": t1 = time.time() # the 0th iteration is often an outlier (much slower) => skip logging it tokens_per_second = ddp_world_size * B * T / (t1-t0) - print0(f"iteration {i}, loss: {loss.item():.4f}, time: {(t1-t0)*1000:.3f}ms, tok/s: {tokens_per_second:.2f}") + print0(f"iteration {i+1}, loss: {loss.item():.4f}, time: {(t1-t0)*1000:.3f}ms, tok/s: {tokens_per_second:.2f}") if i > 0 and i > args.num_iterations - 20: timings.append(t1-t0) From b84571f7459f1c98d90b30b680c4e4c0506987e3 Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 17:54:15 +0000 Subject: [PATCH 07/16] also add argparse to force tf32 to zero --- train_gpt2.cu | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/train_gpt2.cu b/train_gpt2.cu index 5e6fc80..4ede51a 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2089,6 +2089,7 @@ void error_usage() { fprintf(stderr, " -s sample_every, how often we inference the model (default = 20)\n"); fprintf(stderr, " -g genT, how many steps of inference we do (default = 64)\n"); fprintf(stderr, " -a overfit a single batch? 0/1. useful for debugging\n"); + fprintf(stderr, " -f enable_tf32 override (default: 1, set to 0 to disable tf32)\n"); exit(EXIT_FAILURE); } @@ -2109,6 +2110,7 @@ int main(int argc, char *argv[]) { int genT = 64; // number of steps of inference we will do int overfit_single_batch = 0; // useful for debugging, 1 = only load a single data batch once int max_steps = -1; + int override_enable_tf32 = 1; for (int i = 1; i < argc; i+=2) { if (i + 1 >= argc) { error_usage(); } // must have arg after flag if (argv[i][0] != '-') { error_usage(); } // must start with dash @@ -2125,6 +2127,7 @@ int main(int argc, char *argv[]) { else if (argv[i][1] == 's') { sample_every = atoi(argv[i+1]); } else if (argv[i][1] == 'g') { genT = atoi(argv[i+1]); } else if (argv[i][1] == 'a') { overfit_single_batch = atoi(argv[i+1]); } + else if (argv[i][1] == 'f') { override_enable_tf32 = atoi(argv[i+1]); } else { error_usage(); } } printf0("+-----------------------+----------------------------------------------------+\n"); @@ -2158,6 +2161,7 @@ int main(int argc, char *argv[]) { // TF32 precision is equivalent to torch.set_float32_matmul_precision('high') int enable_tf32 = cuda_arch_major >= 8 ? 1 : 0; + if (override_enable_tf32 == 0) { enable_tf32 = 0; } // force to zero via arg cublas_compute_type = enable_tf32 ? CUBLAS_COMPUTE_32F_FAST_TF32 : CUBLAS_COMPUTE_32F; cublasMath_t cublas_math_mode = enable_tf32 ? CUBLAS_TF32_TENSOR_OP_MATH : CUBLAS_DEFAULT_MATH; cublasCheck(cublasSetMathMode(cublas_handle, cublas_math_mode)); From 44656c385085380f50b22b50649a448bd7e2a326 Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 18:38:38 +0000 Subject: [PATCH 08/16] allow fp32 precision in the test script as well --- test_gpt2.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test_gpt2.cu b/test_gpt2.cu index e2b540c..0d1f57d 100644 --- a/test_gpt2.cu +++ b/test_gpt2.cu @@ -109,7 +109,7 @@ int main(int argc, char *argv[]) { // build the GPT-2 model from a checkpoint GPT2 model; - gpt2_build_from_checkpoint(&model, "gpt2_124M_bf16.bin"); + gpt2_build_from_checkpoint(&model, load_filename); size_t V = model.config.vocab_size; size_t Vp = model.config.padded_vocab_size; size_t maxT = model.config.max_seq_len; From 050cbfa42c37fd52de5b8eeb577ef4a333a9494e Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 18:48:12 +0000 Subject: [PATCH 09/16] override train split to val split if we are debugging and trying to overfit a single batch of data, following the python script reference behavior --- train_gpt2.cu | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/train_gpt2.cu b/train_gpt2.cu index 4ede51a..1ef1963 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -2137,7 +2137,7 @@ int main(int argc, char *argv[]) { printf0("| output log file | %-50s |\n", output_log_file == NULL ? "NULL" : output_log_file); printf0("| batch size B | %-50d |\n", B); printf0("| sequence length T | %-50d |\n", T); - printf0("| learning rate | %-50f |\n", learning_rate); + printf0("| learning rate | %-50e |\n", learning_rate); printf0("| max_steps | %-50d |\n", max_steps); printf0("| val_loss_every | %-50d |\n", val_loss_every); printf0("| val_max_batches | %-50d |\n", val_max_batches); @@ -2189,7 +2189,10 @@ int main(int argc, char *argv[]) { char train_tokens_filename[128]; char val_tokens_filename[128]; assert(strlen(input_dataset_prefix) < 100); // being bit lazy here, make sure we don't overflow - sprintf(train_tokens_filename, "%s_train.bin", input_dataset_prefix); + // if we're only overfitting a single batch for debugging, let's overfit the first batch + // from val instead of train split, because val is smaller and a bit faster + const char* train_split = (overfit_single_batch == 1) ? "val" : "train"; + sprintf(train_tokens_filename, "%s_%s.bin", input_dataset_prefix, train_split); sprintf(val_tokens_filename, "%s_val.bin", input_dataset_prefix); DataLoader train_loader; dataloader_init(&train_loader, &multi_gpu_config, train_tokens_filename, B, T); @@ -2296,6 +2299,7 @@ int main(int argc, char *argv[]) { // do a training step clock_gettime(CLOCK_MONOTONIC, &start); if (overfit_single_batch == 0 || (step == 0 && overfit_single_batch == 1)) { + // if we're overfitting a single batch, we'll only call this at step = 0 dataloader_next_batch(&train_loader); } gpt2_forward(&model, train_loader.inputs, train_loader.targets, B, T); From 70b4de8700efca73fd39ce4bb421118a910fbaeb Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 18:51:45 +0000 Subject: [PATCH 10/16] add comment documenting how to reproduce python reference exactly --- train_gpt2.cu | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/train_gpt2.cu b/train_gpt2.cu index 1ef1963..72de93c 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -24,6 +24,13 @@ Also we're using TinyStories here for example as it is a bigger dataset Example launch using bfloat16 on 4 GPUs, same as above: mpirun -np 4 ./train_gpt2cu -b 8 -v 200 -s 200 -i data/TinyStories + +If you'd like to see train_gpt2.cu produce identical results to +`python train_gpt2.py`, you can run it like this: +make train_gpt2cu PRECISION=FP32 +./train_gpt2cu -b 4 -t 64 -l 1e-4 -v 200 -s 200 -a 1 -x 10 -f 0 +This reads & runs in fp32, B=4, T=64, LR=1e-4, val/sample never (200), +-a 1 is "overfit single batch", -x 10 is 10 iterations, and -f 0 disables tf32 */ #include From 9141e04693d53243058c613048b0c97c05ce95db Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 19:35:05 +0000 Subject: [PATCH 11/16] fixes to the PR, careful with float/floatX etc --- dev/cuda/gelu_forward.cu | 16 ++++++++-------- train_gpt2.cu | 2 +- 2 files changed, 9 insertions(+), 9 deletions(-) diff --git a/dev/cuda/gelu_forward.cu b/dev/cuda/gelu_forward.cu index c9fba9e..9bb3ff6 100644 --- a/dev/cuda/gelu_forward.cu +++ b/dev/cuda/gelu_forward.cu @@ -9,10 +9,10 @@ If encountering "error: identifier "M_PI" is undefined", add the following lines #define _USE_MATH_DEFINES #include OR #include -version 1 is naive port from CPU code to kernel +version 1 is naive CPU port, for use in float ./gelu_forward 1 -version 2 uses the Packed128 data structure +version 2 is bfloat16 with the Packed128 data structure ./gelu_forward 2 */ @@ -22,7 +22,7 @@ version 2 uses the Packed128 data structure #include "common.h" // turn on bf16 as default, done up here for now -//#define ENABLE_BF16 +#define ENABLE_BF16 #if defined(ENABLE_BF16) typedef __nv_bfloat16 floatX; @@ -54,7 +54,7 @@ void gelu_forward_cpu(float* out, const float* inp, int N) { // GPU kernels // elementwise ops are nice and ez -__global__ void gelu_kernel(float* out, const float* inp, int N) { +__global__ void gelu_forward_kernel1(float* out, const float* inp, int N) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < N) { float xi = inp[i]; @@ -64,7 +64,7 @@ __global__ void gelu_kernel(float* out, const float* inp, int N) { } // elementwise ops are nice and ez -__global__ void gelu_kernel2(floatX* out, const floatX* inp, int N) { +__global__ void gelu_forward_kernel2(floatX* out, const floatX* inp, int N) { int i = (blockIdx.x * blockDim.x + threadIdx.x) * x128::size; if (i < N) { x128 packed_out; @@ -72,7 +72,7 @@ __global__ void gelu_kernel2(floatX* out, const floatX* inp, int N) { for(int k = 0; k < packed_inp.size; ++k) { float xi = (float)packed_inp[k]; float cube = 0.044715f * xi * xi * xi; - packed_out[k] = (float)(0.5f * xi * (1.0f + tanhf(GELU_SCALING_FACTOR * (xi + cube)))); + packed_out[k] = (floatX)(0.5f * xi * (1.0f + tanhf(GELU_SCALING_FACTOR * (xi + cube)))); } // store instead of storecs (without cache streaming) in case it is useful for the // data to be in the cache for the next operation after this GeLU @@ -85,13 +85,13 @@ __global__ void gelu_kernel2(floatX* out, const floatX* inp, int N) { void gelu_forward1(float* out, const float* inp, int N, const int block_size) { const int grid_size = ceil_div(N, block_size); - gelu_kernel<<>>(out, inp, N); + gelu_forward_kernel1<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } void gelu_forward2(floatX* out, const floatX* inp, int N, const int block_size) { const int grid_size = ceil_div(N, block_size)/x128::size; - gelu_kernel2<<>>(out, inp, N); + gelu_forward_kernel2<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } diff --git a/train_gpt2.cu b/train_gpt2.cu index 37c157c..b141f45 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -574,7 +574,7 @@ __global__ void permute_kernel_backward(floatX* dinp, __global__ void unpermute_kernel(floatX* inp, floatX *out, int B, int N, int NH, int d) { // out has shape (B, nh, N, d) but we need to unpermute it to (B, N, nh, d) - + int idx = (blockIdx.x * blockDim.x + threadIdx.x); // out[b][n][nh_][d_] <- inp[b][nh_][n][d_] if (idx < B * NH * N * d) { From 7ba1ed8a1a48b39012e82b7f94bc93af2158efb8 Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 19:39:20 +0000 Subject: [PATCH 12/16] delete packed128 .fp32, use (float) ok --- dev/cuda/common.h | 3 --- train_gpt2.cu | 3 --- 2 files changed, 6 deletions(-) diff --git a/dev/cuda/common.h b/dev/cuda/common.h index 169dac3..0be5c0d 100644 --- a/dev/cuda/common.h +++ b/dev/cuda/common.h @@ -52,9 +52,6 @@ struct alignas(16) Packed128 { __device__ const ElementType& operator[](int index) const { return payload[index]; } - __device__ float fp32(int index) { - return static_cast(payload[index]); - } __device__ int4 get_bits() const { int4 bits; static_assert(sizeof(bits) == sizeof(payload), "Size mismatch."); diff --git a/train_gpt2.cu b/train_gpt2.cu index b141f45..bfc8d16 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -218,9 +218,6 @@ struct alignas(16) Packed128 { __device__ const ElementType& operator[](int index) const { return payload[index]; } - __device__ float fp32(int index) { - return static_cast(payload[index]); - } __device__ int4 get_bits() const { int4 bits; static_assert(sizeof(bits) == sizeof(payload), "Size mismatch."); From c3515cfe47a721e9cff011a035e92e46de8b6f5d Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Tue, 30 Apr 2024 23:45:34 +0300 Subject: [PATCH 13/16] fixed potential error and generalized gelu forward --- dev/cuda/common.h | 17 +++++++++++++++++ dev/cuda/gelu_forward.cu | 22 +++++++++------------- 2 files changed, 26 insertions(+), 13 deletions(-) diff --git a/dev/cuda/common.h b/dev/cuda/common.h index 0be5c0d..b5b8596 100644 --- a/dev/cuda/common.h +++ b/dev/cuda/common.h @@ -32,6 +32,23 @@ void cublasCheck(cublasStatus_t status, const char *file, int line) } #define cublasCheck(status) { cublasCheck((status), __FILE__, __LINE__); } +// ---------------------------------------------------------------------------- +// converting size_t typed but small values to integer. + +int int_check(size_t small_but_size_t, const char *file, int line) { + if(small_but_size_t > INT32_MAX) { + fprintf(stderr, "Error: Integer conversion failed at %s:%d\n", file, line); + fprintf(stderr, "Error details:\n"); + fprintf(stderr, " File: %s\n", file); + fprintf(stderr, " Line: %d\n", line); + fprintf(stderr, " Value: %zu\n", small_but_size_t); + exit(EXIT_FAILURE); + } + return (int)small_but_size_t; +} + +#define toIntCheck(value) int_check(value, __FILE__, __LINE__) + // ---------------------------------------------------------------------------- // Packed128 data structure, which forces the compiler to use 128-bit loads/stores // in GPUs that support (the LDG.128 and STS.128 instructions) diff --git a/dev/cuda/gelu_forward.cu b/dev/cuda/gelu_forward.cu index 9bb3ff6..a8edf31 100644 --- a/dev/cuda/gelu_forward.cu +++ b/dev/cuda/gelu_forward.cu @@ -9,7 +9,7 @@ If encountering "error: identifier "M_PI" is undefined", add the following lines #define _USE_MATH_DEFINES #include OR #include -version 1 is naive CPU port, for use in float +version 1 is naive CPU port ./gelu_forward 1 version 2 is bfloat16 with the Packed128 data structure @@ -54,7 +54,7 @@ void gelu_forward_cpu(float* out, const float* inp, int N) { // GPU kernels // elementwise ops are nice and ez -__global__ void gelu_forward_kernel1(float* out, const float* inp, int N) { +__global__ void gelu_forward_kernel1(floatX* out, const floatX* inp, int N) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < N) { float xi = inp[i]; @@ -83,14 +83,14 @@ __global__ void gelu_forward_kernel2(floatX* out, const floatX* inp, int N) { // ---------------------------------------------------------------------------- // kernel launcher -void gelu_forward1(float* out, const float* inp, int N, const int block_size) { +void gelu_forward1(floatX* out, const floatX* inp, int N, const int block_size) { const int grid_size = ceil_div(N, block_size); gelu_forward_kernel1<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } void gelu_forward2(floatX* out, const floatX* inp, int N, const int block_size) { - const int grid_size = ceil_div(N, block_size)/x128::size; + const int grid_size = ceil_div(N, toIntCheck(block_size * x128::size)); gelu_forward_kernel2<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } @@ -102,16 +102,12 @@ void gelu_forward(int kernel_num, int B, int T, int C, int block_size) { switch (kernel_num) { -#if !defined(ENABLE_BF16) && !defined(ENABLE_FP16) case 1: gelu_forward1(out, inp, B * T * C, block_size); break; -#endif -#if defined(ENABLE_BF16) case 2: gelu_forward2(out, inp, B * T * C, block_size); break; -#endif default: printf("Invalid kernel number\n"); exit(1); @@ -120,7 +116,7 @@ void gelu_forward(int kernel_num, // ---------------------------------------------------------------------------- -int main(int argc, char **argv) { +int main(int argc, const char **argv) { srand(0); int B = 8; @@ -165,11 +161,11 @@ int main(int argc, char **argv) { printf("Checking block size %d.\n", block_size); gelu_forward(kernel_num, d_out, d_inp, B, T, C, block_size); #if !defined(ENABLE_BF16) && !defined(ENABLE_FP16) - validate_result(d_out, out, "out", B * T * C, 1e-5f); + float tol = 1e-5; +#else + float tol = 1e-2f; #endif -#if defined(ENABLE_BF16) -#endif - validate_result(d_out, out, "out", B * T * C, 1e-2f); + validate_result(d_out, out, "out", B * T * C, tol); } printf("All results match. Starting benchmarks.\n\n"); From 603d862830b6d2e2ea4f31da7465fc548b3381ff Mon Sep 17 00:00:00 2001 From: Erik Schultheis Date: Tue, 30 Apr 2024 23:52:08 +0300 Subject: [PATCH 14/16] fix up bandwidth calculation --- dev/cuda/gelu_forward.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dev/cuda/gelu_forward.cu b/dev/cuda/gelu_forward.cu index a8edf31..464f034 100644 --- a/dev/cuda/gelu_forward.cu +++ b/dev/cuda/gelu_forward.cu @@ -182,7 +182,7 @@ int main(int argc, const char **argv) { // napkin math: estimate the memory bandwidth achieved // for each (B,T,C) output element, we do 1 read and 1 write, 4 bytes each // and e.g. A100 40GB PCIe is advertised at 1,555GB/s - long memory_ops = B * T * C * 2 * 4; + long memory_ops = B * T * C * 2 * (int)sizeof(floatX); float memory_bandwidth = memory_ops / elapsed_time / 1e6; printf("block_size %4d | time %.4f ms | bandwidth %.2f GB/s\n", block_size, elapsed_time, memory_bandwidth); From 769c911ab678c775b777f0d9505d8a19c0eff6c0 Mon Sep 17 00:00:00 2001 From: Andrej Karpathy Date: Tue, 30 Apr 2024 23:48:06 +0000 Subject: [PATCH 15/16] small tweaks --- dev/cuda/common.h | 22 +++------------------- dev/cuda/gelu_forward.cu | 2 +- train_gpt2.cu | 13 ++++++++----- 3 files changed, 12 insertions(+), 25 deletions(-) diff --git a/dev/cuda/common.h b/dev/cuda/common.h index b5b8596..ae352a5 100644 --- a/dev/cuda/common.h +++ b/dev/cuda/common.h @@ -32,23 +32,6 @@ void cublasCheck(cublasStatus_t status, const char *file, int line) } #define cublasCheck(status) { cublasCheck((status), __FILE__, __LINE__); } -// ---------------------------------------------------------------------------- -// converting size_t typed but small values to integer. - -int int_check(size_t small_but_size_t, const char *file, int line) { - if(small_but_size_t > INT32_MAX) { - fprintf(stderr, "Error: Integer conversion failed at %s:%d\n", file, line); - fprintf(stderr, "Error details:\n"); - fprintf(stderr, " File: %s\n", file); - fprintf(stderr, " Line: %d\n", line); - fprintf(stderr, " Value: %zu\n", small_but_size_t); - exit(EXIT_FAILURE); - } - return (int)small_but_size_t; -} - -#define toIntCheck(value) int_check(value, __FILE__, __LINE__) - // ---------------------------------------------------------------------------- // Packed128 data structure, which forces the compiler to use 128-bit loads/stores // in GPUs that support (the LDG.128 and STS.128 instructions) @@ -75,8 +58,9 @@ struct alignas(16) Packed128 { memcpy(&bits, &payload, sizeof(bits)); return bits; } - - static constexpr const size_t size = sizeof(int4) / sizeof(ElementType); + // e.g. sizeof(int4) is 16 (4 X 4 bytes), sizeof(bfloat16) = 2, so size = 8 + // so in the case where ElementType = bfloat16, we store 8 elements in one Packed128 + static constexpr const int size = sizeof(int4) / sizeof(ElementType); ElementType payload[size]; }; diff --git a/dev/cuda/gelu_forward.cu b/dev/cuda/gelu_forward.cu index 464f034..6a582d5 100644 --- a/dev/cuda/gelu_forward.cu +++ b/dev/cuda/gelu_forward.cu @@ -90,7 +90,7 @@ void gelu_forward1(floatX* out, const floatX* inp, int N, const int block_size) } void gelu_forward2(floatX* out, const floatX* inp, int N, const int block_size) { - const int grid_size = ceil_div(N, toIntCheck(block_size * x128::size)); + const int grid_size = ceil_div(N, block_size * x128::size); gelu_forward_kernel2<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } diff --git a/train_gpt2.cu b/train_gpt2.cu index bfc8d16..ec0c125 100644 --- a/train_gpt2.cu +++ b/train_gpt2.cu @@ -673,13 +673,13 @@ __global__ void residual_forward_kernel(floatX* out, floatX* inp1, floatX* inp2, } #define GELU_SCALING_FACTOR sqrtf(2.0f / M_PI) -__global__ void gelu_forward_kernel(floatX* out, const floatX* inp, int N) { +__global__ void gelu_forward_kernel2(floatX* out, const floatX* inp, int N) { int i = (blockIdx.x * blockDim.x + threadIdx.x) * x128::size; if (i < N) { x128 packed_out; x128 packed_inp = load128cs(inp + i); // load and do not keep in cache for(int k = 0; k < packed_inp.size; ++k) { - float xi =(float)packed_inp[k]; + float xi = (float)packed_inp[k]; float cube = 0.044715f * xi * xi * xi; packed_out[k] = (floatX)(0.5f * xi * (1.0f + tanhf(GELU_SCALING_FACTOR * (xi + cube)))); } @@ -1204,9 +1204,9 @@ void residual_forward(floatX* out, floatX* inp1, floatX* inp2, int N) { } void gelu_forward(floatX* out, const floatX* inp, int N) { - const int block_size = 128; - const int grid_size = CEIL_DIV(N/x128::size, block_size); - gelu_forward_kernel<<>>(out, inp, N); + const int block_size = 512; + const int grid_size = CEIL_DIV(N, block_size * x128::size); + gelu_forward_kernel2<<>>(out, inp, N); cudaCheck(cudaGetLastError()); } @@ -2003,6 +2003,8 @@ void dataloader_init(DataLoader *loader, const MultiGpuConfig* multi_gpu_config, cudaMallocHost((void**)&loader->batch, (B * T + 1) * sizeof(int)); loader->inputs = loader->batch; loader->targets = loader->batch + 1; // targets are shifted by one + // note: we definitely want to advance by B * T; That is the "stride" by which we move + // the window of tokens. We only load B * T + 1 tokens because our targets are offset by 1 loader->num_batches = loader->file_size / (loader->num_processes * B * T * sizeof(int)); } @@ -2021,6 +2023,7 @@ void dataloader_next_batch(DataLoader *loader) { fseekCheck(loader->tokens_file, loader->current_position, SEEK_SET); freadCheck(loader->batch, sizeof(int), B*T+1, loader->tokens_file); // advance the current position by B*T*num_processes integers + // note: the "stride" of tokens by which we move each time is definitely B * T loader->current_position += loader->num_processes * B * T * sizeof(int); } From 78eba5fff39b2c3a2caf4e6b920988095686ffa1 Mon Sep 17 00:00:00 2001 From: Jake Hemstad Date: Tue, 30 Apr 2024 18:52:42 -0500 Subject: [PATCH 16/16] Add llm.cpp fork to README --- README.md | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/README.md b/README.md index 0eb8dd3..84f8269 100644 --- a/README.md +++ b/README.md @@ -290,6 +290,10 @@ Lastly, I will be a lot more sensitive to complexity in the root folder of the p ## notable forks +- CUDA C++ + - [llm.cpp](https://github.com/gevtushenko/llm.c) by @[gevtushenko](https://github.com/gevtushenko): a port of this project using the [CUDA C++ Core Libraries](https://github.com/NVIDIA/cccl) + - A presentation this fork was covered in [this lecture](https://www.youtube.com/watch?v=WiB_3Csfj_Q) in the [CUDA MODE Discord Server](https://discord.gg/cudamode) + - Mojo - [llm.🔥](https://github.com/dorjeduck/llm.mojo) by @[dorjeduck](https://github.com/dorjeduck): a Mojo port of this project