Compare commits

..

1515 Commits

Author SHA1 Message Date
Andrej Karpathy 18298f3a40 i misspelled reduced 2024-08-16 18:16:46 +00:00
Andrej Karpathy 8c586f9158 reduce across GPUs nicer 2024-08-16 18:08:34 +00:00
Andrej Karpathy e6856bc566 fallback to memory allocation of m,v,master_weights on host automatically in case of OOM. will run slower but won't OOM 2024-08-16 17:23:42 +00:00
Andrej Karpathy 2882ec6b9e fix makefile for multigpu setups 2024-08-16 17:22:30 +00:00
Erik Schultheis f72c1f2c64 hint to host 2024-08-16 01:44:08 +03:00
Erik Schultheis c8457572fa just try to allocate on device; fallback if that fails 2024-08-16 01:43:56 +03:00
Erik Schultheis 0d52d2a3d7 fall back to cudaMallocManaged for optimizer states if we're out of memory 2024-08-16 01:42:30 +03:00
Andrej 4c84bc743c Merge pull request #740 from karpathy/gordicaleksa-fix_dataloader2
Gordicaleksa fix dataloader2
2024-08-13 11:18:57 -07:00
Andrej Karpathy 755458d0e1 fix tokenizer omg 2024-08-13 17:24:15 +00:00
Andrej Karpathy 16635d41a2 attempt to fix PR 2024-08-13 02:13:14 +00:00
Andrej Karpathy 9740a652cd Merge branch 'fix_dataloader2' of https://github.com/gordicaleksa/llm.c into gordicaleksa-fix_dataloader2 2024-08-12 23:50:10 +00:00
Andrej 1787210306 Merge pull request #738 from ademeure/faster_compile
Improve compile time (simple makefile changes)
2024-08-12 16:45:15 -07:00
Aleksa Gordic 68624cc6ac Refactor tokenize, delegate instead 2024-08-12 20:26:21 +02:00
Aleksa Gordic e8739c89a0 Refactor magic/version header dict 2024-08-12 20:17:53 +02:00
ademeure fc8834409b Improve compile time via nvidia-smi instead of __nvcc_device_query (& add "--threads=0"). 2024-08-12 17:07:36 +00:00
ademeure 9b1c14bd93 Improve compile time via nvidia-smi instead of __nvcc_device_query (& add "--threads=0"). 2024-08-12 16:51:39 +00:00
Aleksa Gordic eb1b359dae Replace fire with argparse 2024-08-10 22:45:46 +02:00
Aleksa Gordic 311c149a2a Add fire to reqs 2024-08-10 22:39:07 +02:00
Aleksa Gordic 50765638c0 Add fineweb llama 3 support 2024-08-10 22:35:46 +02:00
Aleksa Gordic f8c1ff590f Use fire module, update readme 2024-08-10 22:29:19 +02:00
Aleksa Gordic 15b76bf0fb Add tiny shakespeare 2024-08-10 22:23:33 +02:00
Aleksa Gordic b018b0c85c Change default dataset path to tiny stories 2024-08-10 21:50:39 +02:00
Aleksa Gordic 71419cbd61 Update header description 2024-08-10 21:47:09 +02:00
Aleksa Gordic d3151c4ace Add llama 3 tokenization and data loader 2024-08-10 21:25:57 +02:00
Andrej 6e6a528111 Merge pull request #725 from gordicaleksa/llama
Add LLaMA 3 Python support
2024-08-08 14:08:20 -07:00
Aleksa Gordic d773c88e49 Remove unnecessary comment 2024-08-08 21:06:01 +02:00
Aleksa Gordic c5c87fc4e7 Add comments, fix stop tokens 2024-08-08 20:54:17 +02:00
Aleksa Gordic 8b672ffcb9 Add data loader not implemented exception 2024-08-08 18:55:17 +02:00
Aleksa Gordic f1c91f8ae3 Address Andrej's PR comments 2024-08-08 18:45:04 +02:00
Aleksa Gordic 89addd3e3e Remove pycache 2024-08-08 18:20:51 +02:00
Aleksa Gordic ac01536b99 Remove the whole llmc_py 2024-08-08 18:20:18 +02:00
Aleksa Gordic dfd459bfc4 Remove prompts.json 2024-08-08 18:19:40 +02:00
Aleksa Gordic 624ed3ce30 Rename the new file to train llama3 2024-08-08 18:15:37 +02:00
Aleksa Gordic b7c98c9365 Add llama config error check 2024-08-08 18:12:35 +02:00
Aleksa Gordic b25e325c69 Add explicit external mask 2024-08-08 17:55:01 +02:00
Aleksa Gordic d4ef9c5afd Remove llmc_py, single file 2024-08-08 17:09:41 +02:00
Aleksa Gordic 72dcfeb404 Remove float(-inf) 2024-08-04 22:20:46 +02:00
Aleksa Gordic 8e55d16863 Support HF & Meta models 2024-08-04 16:27:12 +02:00
Aleksa Gordic 0749a4af42 Minor refactor - remove wpe pos array from fwd 2024-08-04 09:28:17 +00:00
Aleksa Gordic 4b386a2a89 Test training loop works 2024-08-04 09:13:29 +00:00
Aleksa Gordic cd902735b2 Fix saving model / state logic 2024-08-04 08:54:21 +00:00
Aleksa Gordic 7768a36f37 Add prompts file 2024-08-04 10:07:48 +02:00
Aleksa Gordic 879cc5f46c Remove init weights as it's gpt-2 specific 2024-08-04 08:00:44 +00:00
Aleksa Gordic bad7857dff Done - need to test train loop and saving model 2024-08-03 22:37:54 +02:00
Aleksa Gordic b1acb5988f Got to llama 3 inference (end) 2024-08-03 22:22:38 +02:00
Aleksa Gordic d855c9695a Refactoring - got to main 2024-08-03 22:16:13 +02:00
Aleksa Gordic c1d2b7fec1 Removed from pretrained for GPT-2 2024-08-03 22:02:02 +02:00
Aleksa Gordic 8a48df7b16 MLP GPT2->LLaMA3 2024-08-03 21:53:26 +02:00
Aleksa Gordic ccdbdfd4ce Going towards pure llama 3 file - fixed attn 2024-08-03 21:49:45 +02:00
Aleksa Gordic 8919b66c37 Fix generate for gpt-2 2024-08-03 14:14:05 +02:00
Aleksa Gordic 180215fd60 Unify generate and generate_llama 2024-08-03 12:38:49 +02:00
Aleksa Gordic fa7bcc3f29 One more minor attn fix 2024-08-03 11:54:10 +02:00
Aleksa Gordic fdd5345931 Refactor attn mechanism 2024-08-03 11:49:26 +02:00
Aleksa Gordic 0199e51a26 Refactor MLP 2024-08-03 10:46:15 +02:00
Aleksa Gordic de879d129b Refactor checkpoint state dict map func 2024-08-03 10:28:32 +02:00
Aleksa Gordic c0c08ba53d Have prompts in a file instead of inline, prompt 4 is different 2024-08-03 10:21:56 +02:00
Aleksa Gordic f50f2de8a2 Refactor attn, change numerics but equivalent 2024-08-03 10:11:26 +02:00
Aleksa Gordic 465aac4aa5 Equivalent to nano llama 3 reference code 2024-08-03 10:03:28 +02:00
Aleksa Gordic c414d0284a Minor refactor 2024-08-02 23:12:49 +02:00
Aleksa Gordic 838cd13c4f Refactor 2024-08-02 22:38:11 +02:00
Aleksa Gordic 41bf8e09d7 Equivalent with nano llama 3 2024-08-02 21:27:05 +02:00
Andrej 29aacba1f5 Merge pull request #705 from gordicaleksa/refactor_c
Refactor C code
2024-07-30 13:25:04 -07:00
Andrej ef12d1b80e Merge pull request #717 from ngc92/nvml
Nvidia management library for more detailed GPU state printing
2024-07-30 13:10:23 -07:00
Andrej Karpathy 3cefe09f7e Merge branch 'master' of github.com:karpathy/llm.c 2024-07-30 20:08:13 +00:00
Andrej Karpathy e48ebc0138 Merge branch 'mspronesti-master' 2024-07-30 20:08:04 +00:00
Andrej Karpathy 1566d6b9f6 revert makefile change 2024-07-30 20:07:54 +00:00
Andrej 9cc357e27d Merge pull request #715 from karpathy/feature/restore_from_master
Feature/restore from master
2024-07-30 13:04:38 -07:00
Andrej Karpathy 51dd102328 remove confusing comment 2024-07-30 20:01:43 +00:00
Erik Schultheis e4b5e5775f disable if header is not found 2024-07-29 10:07:37 +02:00
Erik Schultheis bb9467dcca make gpu log configurable 2024-07-29 09:40:50 +02:00
Erik Schultheis bdff672741 add temperature 2024-07-29 09:40:50 +02:00
Erik Schultheis 294a6bca55 throttle reason 2024-07-29 09:40:50 +02:00
Erik Schultheis 743062142c units + more frequent reporting 2024-07-29 09:40:50 +02:00
Erik Schultheis 7d7954caa7 nvml for more detailed gpu status info 2024-07-29 09:40:50 +02:00
Andrej Karpathy 2b827f1659 bring back state allocation into build_from_checkpoint 2024-07-29 02:39:50 +00:00
mspronesti 574ef9b87d chore(permute): use common.h utils for validation and benchmarking 2024-07-29 00:19:31 +01:00
mspronesti a49da7d917 make: generalize TARGETS definition 2024-07-29 00:19:03 +01:00
Andrej Karpathy a794bcb395 small fixes 2024-07-28 17:10:20 +00:00
Andrej Karpathy b2ae84745d Merge branch 'restore_from_master_weights' of https://github.com/ademeure/llm.c into ademeure-restore_from_master_weights 2024-07-28 16:23:16 +00:00
Andrej b4623bc5e7 Merge pull request #712 from indianspeedster/dev-cuda
Added permute kernel in dev/cuda
2024-07-28 06:08:20 -10:00
Andrej 7379cd10b4 Merge pull request #713 from mspronesti/master
fix(dev/cuda): memory leaks
2024-07-28 06:05:44 -10:00
Andrej 362c6a8d7f Merge pull request #700 from YuchenJin/integer-overflow
Fix integer overflow by using `size_t` for parameter sizes.
2024-07-27 07:26:58 -10:00
mspronesti c6e74ea322 fix(dev/cuda): memory leaks 2024-07-26 00:33:51 +01:00
indianspeedster a80bce65e5 Added compilation steps in makefile for permute.cu 2024-07-25 18:02:06 -04:00
indianspeedster 8952c524af added kernel to permute in dev/cuda 2024-07-25 17:35:29 -04:00
ademeure 52e6e0f80b fix bug from merge (init_state set to false too late) 2024-07-25 21:32:01 +00:00
ademeure 4d77ece6b4 simplified further (don't need non-functional error checking...) 2024-07-25 22:04:14 +01:00
ademeure 2eabc223cc simplify a little bit 2024-07-25 22:04:14 +01:00
ademeure 9781627a8a allow restoring from checkpoint of different precision 2024-07-25 22:04:14 +01:00
ademeure 5cae10f907 make restoring from master weights actually work 2024-07-25 22:04:14 +01:00
ademeure f470fbd56d Allow restoring weights from the master weights of a checkpoint (deterministically by also saving RNG state of last update) 2024-07-25 22:04:14 +01:00
Andrej cb4451137b Merge pull request #703 from rosslwheeler/bugfix/msvc_latest_fix
Fix for upgraded Cuda 12.5.1 and Microsoft latest compiler
2024-07-25 13:31:41 -07:00
Andrej e33402f7f3 Merge pull request #694 from ngc92/model-init-cleanup
Model init cleanup
2024-07-25 13:14:11 -07:00
Aleksa Gordic 16c990fdc0 Refactor C code 2024-07-21 11:20:04 +02:00
Ross Wheeler 9dc920a66f Move WIN32 code into unistd.h 2024-07-20 12:05:20 -07:00
Ross Wheeler b203820232 Including unistd.h in cudnn_att.cpp for latest MSVC compiler 2024-07-20 11:57:10 -07:00
YuchenJin 30f4f223be Fix integer overflow by using size_t for parameter sizes. 2024-07-19 08:55:44 -07:00
Erik Schultheis 9bb7a36730 tiny cleanup 2024-07-18 21:47:26 +02:00
Erik Schultheis f2ac85ea53 estimate maximum batch size 2024-07-18 20:12:29 +02:00
Erik Schultheis 794ed07f3c print memory report after allocating model state 2024-07-18 20:06:13 +02:00
Erik Schultheis 251421102d cleanup; these mallocs are no longer needed 2024-07-18 20:00:51 +02:00
Erik Schultheis bfce9bdbcf bugfixes 2024-07-18 19:43:47 +02:00
Erik Schultheis 0bcb3e5840 allocate_weights utility function
reverted some earlier changes + more error checking
2024-07-18 19:29:49 +02:00
Erik Schultheis cde79ffb49 allocate_state utility function 2024-07-18 19:19:06 +02:00
Erik Schultheis 5c89416adf also allocate gradient eagerly
moved helper buffer allocation to forward, so backward now no longer does any memory allocation
2024-07-18 18:54:51 +02:00
Andrej Karpathy 72698a5ba3 update gpt3 125M script to be better. should exactly reproduce gpt-3 125M, except we don't use banded attention and we use fineweb 2024-07-18 16:43:14 +00:00
Erik Schultheis b20def92d4 consolidate model init 2024-07-18 17:23:06 +02:00
Erik Schultheis cbe04c3f92 moved dtype utilities to cuda_utils 2024-07-18 17:13:38 +02:00
Andrej 85d17f4aa2 Merge pull request #668 from abhilash1910/tpc
Add Habana gaudi2 tpc kernel link
2024-07-15 18:22:15 -07:00
Andrej Karpathy 835bb2e6d7 delete segmentation fault when B < 4 atm 2024-07-16 01:20:16 +00:00
Andrej 88514df231 Merge pull request #689 from karpathy/refactor/code-to-zerocuh
Refactor/code to zerocuh
2024-07-15 12:51:01 -07:00
Andrej Karpathy 5afc57064c document the imports from zero.cuh to be explicit 2024-07-15 19:09:52 +00:00
Andrej Karpathy 882a214ad4 Merge branch 'move-to-zero' of https://github.com/ngc92/llm.c into ngc92-move-to-zero 2024-07-15 18:59:22 +00:00
Andrej a6dc64533a Merge pull request #688 from karpathy/feature/gpt3v1
feature/gpt3v1
2024-07-15 11:53:05 -07:00
Andrej Karpathy 220125f39c ok let's allow T < max_seq_len because our tests use it. make warning very prominent 2024-07-15 18:21:34 +00:00
Andrej Karpathy faf8d159a3 use int instead of bool 2024-07-15 18:01:01 +00:00
Andrej Karpathy 528b1b0e31 fix fat fingers 2024-07-15 17:56:41 +00:00
Andrej Karpathy 039904ee9f revert legacy behavior in .sh scripts for -e option 2024-07-15 17:55:26 +00:00
Andrej Karpathy 64db3ded85 refactor gpt2 gpt3 descriptor to keep legacy behavior of -e, be strict everywhere, remove interpolation for now 2024-07-15 17:51:27 +00:00
Andrej Karpathy 4e629e22a1 Merge branch 'gpt3-series' of https://github.com/ngc92/llm.c into ngc92-gpt3-series 2024-07-15 16:46:30 +00:00
Andrej f45c2191b6 Merge pull request #684 from rosslwheeler/feature/loss_checker_ci
Adding CI check for exceeding loss tolerance
2024-07-15 09:30:59 -07:00
Andrej e71c1565c9 Merge pull request #686 from indianspeedster/cuda
Added  cudaCheck wherever missing.
2024-07-15 09:28:02 -07:00
indianspeedster 6bb562bd88 added cudaCheck where missing for proper error checking. 2024-07-13 13:20:39 -04:00
Andrej a6f6bc74a8 Merge pull request #682 from AnswerDotAI/master
Add a README link under related related projects for gpu.cpp under WebGPU C++
2024-07-13 10:12:59 -07:00
Andrej 0e2ff48440 Merge pull request #685 from planetis-m/patch-1
Add nim port
2024-07-13 10:08:26 -07:00
Antonis Geralis a86a69c135 Add nim port 2024-07-13 18:57:51 +03:00
Ross Wheeler 56621bba0e Adding CI check for exceeding loss tolerance
LLMC-26
2024-07-13 01:14:20 -07:00
austinvhuang ddd8aa0f34 add gpu.cpp link to related projects in readme 2024-07-12 14:42:24 -04:00
Erik Schultheis 37b6e328c1 enable model interpolation 2024-07-12 15:07:05 +02:00
Erik Schultheis b5d5ed7f68 updated example scripts 2024-07-12 14:46:09 +02:00
Erik Schultheis d19a6267a8 add gpt3 model series 2024-07-12 14:43:22 +02:00
Erik Schultheis 39a41fa701 move cpu float reduction to zero 2024-07-12 11:12:58 +02:00
Erik Schultheis 4d1ec6832a rename parameters to avoid name clashes with global variable 2024-07-11 19:04:02 +02:00
Erik Schultheis 09b78bc902 move set_zero_configs into zero.cuh 2024-07-11 09:55:31 +02:00
Andrej bdb0fb5599 Merge pull request #671 from ademeure/less_slow_inference
Faster inference by changing (B,T) to (1,t)
2024-07-10 17:47:04 -07:00
Andrej db2454f88f Merge pull request #620 from ngc92/small-fixes
small fixes based on clang-tidy
2024-07-09 12:53:34 -07:00
Erik Schultheis 2b76155976 small fixes based on clang-tidy 2024-07-09 18:37:09 +02:00
ademeure 2fd6fee971 fix the corner case where T < 256 2024-07-07 12:35:29 +00:00
ademeure efa6767a8a Make inference sanity checks much faster by changing (B,T) to (1,256) 2024-07-07 12:23:26 +00:00
Abhilash Majumder a2737d366a add gaudi2 tpc kernel link 2024-07-04 00:34:13 +05:30
Andrej f0abd09ffa Merge pull request #641 from gordicaleksa/add_close_check_socket
Add check versions of functions
2024-07-02 14:26:03 -07:00
Aleksa Gordic 6b62f3b9ae Fix linker issue 2024-07-02 21:54:50 +02:00
Aleksa Gordic d27331cced Move socket header to utils.h 2024-07-02 21:17:28 +02:00
Aleksa Gordic 5645bd18b1 CI fix - add sockets header 2024-07-02 21:17:28 +02:00
Aleksa Gordic 1ae3243185 Add check versions of socket close 2024-07-02 21:17:28 +02:00
Andrej Karpathy c6efd90a86 upload good 1558M sh script 2024-07-02 16:49:09 +00:00
Andrej 39270cc69a Merge pull request #665 from ngc92/zero-grad
zero-grad is async and part of backward call
2024-07-01 16:08:07 -07:00
Erik Schultheis 486f98e3cc zero-grad is async and part of backward call 2024-07-02 00:47:41 +02:00
Andrej Karpathy 79223b950b Merge branch 'ademeure-optimize_attproj_fcproj' 2024-07-01 22:07:24 +00:00
Andrej Karpathy cee396ab5a resolve merge conflict 2024-07-01 22:07:14 +00:00
Andrej Karpathy 48fee049cb add edu fineweb sh download script, change paths to my HF repo, add hellaswag to starter pack 2024-07-01 21:41:15 +00:00
Andrej a876282eb8 Merge pull request #653 from ademeure/cublaslt_refactor
Matmul refactor using only cuBLASLt + GELU Fusion
2024-07-01 11:53:02 -07:00
ademeure 00247dca9b Disable GELU fusion by deffault & fix potential argv[] out of bounds access 2024-07-01 18:40:21 +00:00
Andrej 2e6a397d5f Merge pull request #663 from karpathy/fix/reshuffle_zero_grad
Fix/reshuffle zero grad
2024-07-01 10:39:23 -07:00
Andrej Karpathy 7b92930021 move things around a bit, a few comments, and also update python script to correspond to C more 2024-07-01 17:20:46 +00:00
Andrej Karpathy 8253823432 Merge branch 'fix_periodic_spikes' of https://github.com/gordicaleksa/llm.c into gordicaleksa-fix_periodic_spikes-2 2024-07-01 17:14:22 +00:00
Andrej 942fed59fe Merge pull request #662 from karpathy/overfit-batch-actually
fix the overfit single batch behavior to actually overfit batch, not …
2024-07-01 10:12:20 -07:00
Andrej Karpathy 7fd7717816 fix the overfit single batch behavior to actually overfit batch, not microbatch, we do this more cleanly by simply resetting the dataloader every step 2024-07-01 17:00:40 +00:00
Aleksa Gordic 74e5ba7260 Add micro step guard 2024-07-01 12:11:02 +02:00
Aleksa Gordic 78e7627859 Remove local args for testing 2024-07-01 11:31:55 +02:00
Aleksa Gordic de8454d2f4 Make the grad cleaning more robust 2024-07-01 11:30:36 +02:00
Aleksa Gordic d2f06aed01 Fix periodic spikes - clean loss after val 2024-07-01 11:24:16 +02:00
Andrej Karpathy 61625d168c only rank 0 creates log dir 2024-07-01 06:31:37 +00:00
ademeure 20e450ca70 merge fix 2024-06-30 19:48:53 +00:00
ademeure 6e29dd000d fix default for test_gpt2cu 2024-06-30 20:45:37 +01:00
ademeure 0f11a4ec49 Make GELU fusion default 2 for H100+ and 0 on older GPUs for now 2024-06-30 20:45:37 +01:00
ademeure fa5a196be0 Refactor matmul to only use cuBLASLt and optionally use fused GELU 2024-06-30 20:45:37 +01:00
Andrej Karpathy ea36d04e39 small cosmetic fixes to fineweb data file 2024-06-30 15:52:51 +00:00
ademeure ba4e2c05ad Remove per-layer attproj and fcproj activation tensors (reuse output as scratch buffer instead). Significant memory saving. 2024-06-30 15:28:05 +00:00
Andrej 2daaf6ab83 Merge pull request #629 from rosslwheeler/feature/tests
CI Dataloader test and ptx/sass file generator
2024-06-29 15:17:31 -07:00
Andrej Karpathy e6e365a1a8 Merge branch 'eliebak-add_edu_fineweb_support' 2024-06-28 19:35:58 +00:00
Andrej Karpathy ff7a90bd8d Merge branch 'add_edu_fineweb_support' of https://github.com/eliebak/llm.c into eliebak-add_edu_fineweb_support 2024-06-28 19:35:44 +00:00
Andrej Karpathy 2cf7c5b7d3 Merge branch 'ngc92-mixed-dtypes-2' 2024-06-27 17:47:22 +00:00
Andrej Karpathy 012915c06c handle unknown datatype to squash compiler warning 2024-06-27 17:46:46 +00:00
Andrej 06cd83eacc Merge pull request #652 from ademeure/cudnn_determinism
Make cuDNN deterministic for Flash Attention backward
2024-06-27 10:22:04 -07:00
ademeure 29a067f916 Make cuDNN deterministic for Flash Attention backward (requires frontend >=1.5) 2024-06-27 13:33:02 +00:00
Ross Wheeler d6d08be3c0 Addressed review comments
Explicit precision set for tests
2024-06-27 00:58:19 -07:00
Ross Wheeler eb0b74e559 Add H100 and A100 ptx/sass files 2024-06-27 00:30:38 -07:00
rosslwheeler 267efc9a0e Merge branch 'karpathy:master' into feature/tests 2024-06-26 22:38:21 -07:00
ngc92 98875dd022 macro-based alternative 2024-06-27 02:31:02 +02:00
ngc92 d38c614eb9 add some comments 2024-06-26 23:58:31 +03:00
Erik Schultheis 7beaa1ed96 keep layernorm stats in fp32 2024-06-26 20:48:11 +03:00
Erik Schultheis f0c3e35de8 switch losses to fp32 2024-06-26 20:48:08 +03:00
Erik Schultheis ca47165d29 allocation function for mixed dtype activation 2024-06-26 20:48:03 +03:00
Andrej 68535cb6ce Merge pull request #649 from karpathy/feature/openmpi_flexible_path
tune openmpi make a bit more convenient
2024-06-26 09:11:22 -07:00
Andrej Karpathy 7c283df9c4 use ?= for OPENMPI_DIR so it can be overriden 2024-06-26 16:06:54 +00:00
Andrej c1ac121c77 Merge pull request #646 from Ricardicus/makefile-fix
Makefile: Respect the NO_MULTI_GPU
2024-06-26 08:42:59 -07:00
Andrej a08e32e474 Merge pull request #648 from krrishnarraj/mention_fork
[readme] mention opencl fork
2024-06-26 08:41:15 -07:00
Andrej Karpathy 2a8f82f14f tune openmpi make a bit more convenient 2024-06-26 15:39:10 +00:00
Krishnaraj Bhat 0aa0eb09a6 [readme] mention opencl fork 2024-06-26 20:32:21 +05:30
Rickard Hallerbäck 66cd8c3f1f Makefile: respect the NO_MULTI_GPU 2024-06-26 14:55:02 +02:00
Andrej 44fde98344 Merge pull request #637 from karpathy/feature/outliers
add outlier detector, test for it, and start tracking z score of loss
2024-06-25 21:16:29 -07:00
Andrej Karpathy e12d090483 resolve merge conflicts 2024-06-26 04:05:46 +00:00
Andrej Karpathy ac018c3302 Merge branch 'ngc92-more-streams' 2024-06-25 23:41:55 +00:00
Andrej Karpathy 9c7f1f9701 minor comments 2024-06-25 23:41:40 +00:00
Andrej Karpathy 4f84e3daa7 Merge branch 'more-streams' of https://github.com/ngc92/llm.c into ngc92-more-streams 2024-06-25 23:25:57 +00:00
Erik Schultheis 7171f78f04 fix check call 2024-06-26 02:21:44 +03:00
Andrej Karpathy c0e84eb859 Merge branch 'more-streams' of https://github.com/ngc92/llm.c into ngc92-more-streams 2024-06-25 23:16:50 +00:00
Erik Schultheis 4aae1e529b fix up profile 2024-06-26 02:11:55 +03:00
Erik Schultheis 6c24f13b55 check token validity 2024-06-26 02:01:24 +03:00
Erik Schultheis 7b13f20343 moved targets to backward/validate 2024-06-26 01:56:23 +03:00
Andrej Karpathy 496d8bf449 Merge branch 'more-streams' of https://github.com/ngc92/llm.c into ngc92-more-streams 2024-06-25 22:44:00 +00:00
Erik Schultheis 04e1c183ec set mean_loss to -1 if not calculated 2024-06-26 01:18:19 +03:00
Andrej 16b5bd53fb Merge pull request #636 from karpathy/feature/rolling_checkpoints
rolling checkpoints
2024-06-25 14:43:21 -07:00
Andrej Karpathy 11e2fd1cf2 separate out write_checkpoint and delete_checkpoint 2024-06-25 20:46:03 +00:00
Erik Schultheis 7018280448 fix zero-1 2024-06-25 21:28:52 +03:00
Erik Schultheis 09503dea4f make windows happy 2024-06-25 21:28:52 +03:00
Erik Schultheis 2413b0a6d0 fix profile 2024-06-25 21:28:52 +03:00
Erik Schultheis 39f976f013 added missing free 2024-06-25 21:28:52 +03:00
Erik Schultheis 8b36c2f60b fix loss accumulation 2024-06-25 21:28:52 +03:00
Erik Schultheis 75d4610442 update norm computation: multi-process reduce on GPU 2024-06-25 21:28:52 +03:00
Erik Schultheis 8d994782cf moved loss calculation to backward 2024-06-25 21:28:52 +03:00
Erik Schultheis 4745f53589 boolean_constant utilities 2024-06-25 21:28:52 +03:00
Andrej 1f49cc1c03 Merge pull request #639 from rosslwheeler/bug_fixes/windows_clean_up
Add unistd.h to fix Windows cuDNN build
2024-06-25 11:23:31 -07:00
Andrej be7a278351 Merge pull request #640 from ahrefs/multigpu-fs-build-fix
Add missing MULTI_GPU compiler flag
2024-06-25 11:22:51 -07:00
Ross Wheeler 53985a77f4 CI for Windows needs cuDNN installed - revert 2024-06-25 10:43:06 -07:00
Ross Wheeler 2e6fd9b150 Add unistd.h for Windows cuDNN build
Small fix - move WIN32 change into dev.unistd.h

Add USE_CUDNN=1 build for Windows for CI
2024-06-25 10:36:59 -07:00
chinthysl 934e6f22c5 Fix missing compiler flag 2024-06-25 09:35:24 +00:00
rosslwheeler 2e32e25e37 Merge branch 'karpathy:master' into feature/tests 2024-06-24 23:05:55 -07:00
Andrej Karpathy b51fe0124d oops forgot to update defaults comment in the argparse, smallfix 2024-06-25 05:07:42 +00:00
Andrej Karpathy 665ff144fb skip update if lossz or gradz are above thresholds determined by new args -sl and -sg 2024-06-25 04:50:38 +00:00
Andrej Karpathy 329aee333e fix training loop in other parts 2024-06-25 04:30:34 +00:00
Andrej Karpathy 59830367f2 refactor grad norm and add tracking of grad norm outliers as well 2024-06-25 03:38:21 +00:00
Andrej Karpathy 7a0d864f37 oops i was only using window size 16 for testing let's make it 128 or something 2024-06-25 01:56:04 +00:00
Andrej Karpathy 569f9f16f9 add outlier detector, test for it, and start tracking z score of the loss, for now only print it 2024-06-25 01:37:31 +00:00
Andrej Karpathy 8952f0235b rolling checkpoints - checkpoints are either MINOR or MAJOR and minor checkpoints get deleted with a rolling window. This is an optimization that will allow us to save state more often, but preserve disk space overall. And if we ever see loss spikes and such, it is easier to reset the state to an earlier checkpoint that is more recent 2024-06-24 23:39:18 +00:00
Andrej 3a986bc81b Merge pull request #619 from gordicaleksa/cast_2d_noise
Cast Get2dNoiseUint computation to uint
2024-06-24 14:43:20 -07:00
Andrej 76b133d3f4 Merge pull request #631 from gordicaleksa/minor_refactor2
Minor refactor
2024-06-24 14:38:27 -07:00
Andrej 69b50adcf5 Merge pull request #632 from gordicaleksa/multi_node_my
MPI/TCP/FS for NCCL-init
2024-06-24 14:36:30 -07:00
Aleksa Gordic 4af6a6ab3d Remove /tmp as default init value for fs_path 2024-06-24 21:26:55 +00:00
Aleksa Gordic 2667aa1e3e Fix windows include issue 2024-06-24 23:18:38 +02:00
Aleksa Gordic ac1e37b659 Remove linux socket funcs when windows 2024-06-24 18:56:48 +00:00
Aleksa Gordic 60c95c82aa Add a multi node readme section 2024-06-24 18:50:56 +00:00
Aleksa Gordic f9d0e8b67a Add socket logic for Windows 2024-06-24 18:26:53 +00:00
Aleksa Gordic 05ef1e7dc0 Fix makefile 2024-06-24 14:53:58 +00:00
Aleksa Gordic c420a522b2 Fix arg permutation 2024-06-24 14:38:49 +00:00
Aleksa Gordic fbf4f59c8f Add MPI under multi gpu guard 2024-06-24 16:26:12 +02:00
Aleksa Gordic a259eeabe5 Fix CI errors 2024-06-24 16:26:12 +02:00
Aleksa Gordic 3abdf9a3e4 Add gpus per node arg 2024-06-24 16:26:12 +02:00
Aleksa Gordic 75a2c8b10f Add usage for new args 2024-06-24 16:26:12 +02:00
Aleksa Gordic bfdea887eb Add launch scripts - tested 2024-06-24 16:26:12 +02:00
Aleksa Gordic 1bc3a5ef6e Add all 3 init methods 2024-06-24 16:26:12 +02:00
Aleksa Gordic 1c6c5ff0b7 TCP working 2024-06-24 16:26:12 +02:00
Aleksa Gordic 5500b61da5 fs working 2024-06-24 16:26:12 +02:00
Aleksa Gordic 08c055cd72 Free up unified buffer mem 2024-06-24 16:26:12 +02:00
Aleksa Gordic 1bd737546c Mini refactor 2024-06-24 16:26:12 +02:00
Aleksa Gordic ac72875c24 Use NCCL stream 2024-06-24 16:26:12 +02:00
Aleksa Gordic 52db158367 Can not call set device after NCCL init 2024-06-24 16:26:12 +02:00
Aleksa Gordic 0addb503ef Use NCCL where possible 2024-06-24 16:26:12 +02:00
Andrej Karpathy 2a4be7f349 Merge branch 'ngc92-streams-io' 2024-06-23 01:14:50 +00:00
Andrej Karpathy 98e928f1d8 minor fixes for disk device io 2024-06-23 01:14:42 +00:00
Andrej Karpathy 2543b62b88 Merge branch 'streams-io' of https://github.com/ngc92/llm.c into ngc92-streams-io 2024-06-23 00:48:05 +00:00
Aleksa Gordic 71f10e7960 Rename gpt2 backward 2024-06-22 14:32:51 +02:00
Ross Wheeler 5b95bff118 Dataloader test and ptx/sass file generator
Setup new test infrastructure
2024-06-21 17:23:03 -07:00
Andrej Karpathy 72a2158812 remove a spurious print that is making the table a bit uglier 2024-06-21 15:08:30 +00:00
Andrej Karpathy c6e21c78b5 Merge branch 'master' of github.com:karpathy/llm.c 2024-06-21 02:40:19 +00:00
Andrej Karpathy 61ed938dfb delete spurious check it's fine 2024-06-21 02:40:11 +00:00
Andrej 5b82822462 Merge pull request #627 from karpathy/feature/lr_schedulers
feature/lr_schedulers
2024-06-20 19:31:42 -07:00
Andrej Karpathy 4b5fecc8b2 small touchups and add the wsd schedule 2024-06-21 00:49:01 +00:00
Andrej Karpathy 96ebb16298 simplify lr schedulers 2024-06-21 00:29:47 +00:00
Andrej Karpathy 9abd9ec8a8 Merge branch 'add_lr_schedules' of https://github.com/gordicaleksa/llm.c into gordicaleksa-add_lr_schedules 2024-06-21 00:12:15 +00:00
Andrej dad58b6800 Merge pull request #625 from gordicaleksa/add_nccl_to_readme
Add NCCL instruction to README
2024-06-20 16:06:31 -07:00
Andrej Karpathy 2677d82dcb add shell script that downloads the .bin files of FineWeb100B directly 2024-06-20 22:39:09 +00:00
Aleksa Gordic 5b4d11ac1d Add NCCL instruction 2024-06-21 00:16:17 +02:00
Aleksa Gordic 716808d9b9 Use C not C++ string header 2024-06-20 16:09:57 +02:00
Aleksa Gordic 3e9fb6cb36 Fix compile issues - add comments 2024-06-20 16:09:57 +02:00
Aleksa Gordic 7416ee425f Simplify logic since they all share similar args 2024-06-20 16:09:57 +02:00
Aleksa Gordic 9d141ddab9 Add ugly switching between schedulers 2024-06-20 16:09:57 +02:00
Aleksa Gordic 7308c2831c Add linear LR scheduler 2024-06-20 16:09:57 +02:00
Aleksa Gordic 6f3d099e88 Add constant lr scheduler 2024-06-20 16:09:57 +02:00
Aleksa Gordic 4a1c5a9c0f Add cyclic triangular LR scheduler 2024-06-20 16:09:57 +02:00
Aleksa Gordic 2eaea3b397 Add schedulers header 2024-06-20 16:09:57 +02:00
Aleksa Gordic 8249736233 Add learning rate schedulers 2024-06-20 16:09:57 +02:00
Aleksa Gordic bb9eab8736 Change positionX to uint 2024-06-19 21:05:28 +02:00
Andrej 2ddf128365 Merge pull request #615 from gordicaleksa/relax_grad_tests
Relax grad tensor thresholds in tests
2024-06-19 09:37:42 -07:00
Aleksa Gordic 66950f3790 Cast computation to uint 2024-06-19 17:45:34 +02:00
Aleksa Gordic 8079bd9dbb Relax grad tensor thresholds a bit 2024-06-19 10:21:59 +02:00
Andrej Karpathy fd564f7b8a add the gpt2 1558M repro script too 2024-06-19 02:26:39 +00:00
Andrej Karpathy 0bc8753bc0 update vislog to be a bit more fancy, has GPT-2/3 perf for hellaswag 2024-06-19 01:51:09 +00:00
Andrej Karpathy 3bce68b495 Merge branch 'ChrisDryden-script_to_download_tokenized_dataset' 2024-06-19 01:45:08 +00:00
Andrej Karpathy b29478e20b add starter pack .sh script for faster quickstart 2024-06-19 01:44:36 +00:00
Andrej Karpathy 48b7885f72 Merge branch 'script_to_download_tokenized_dataset' of https://github.com/ChrisDryden/llm.c into ChrisDryden-script_to_download_tokenized_dataset 2024-06-19 01:29:54 +00:00
Andrej Karpathy 483f675e92 Merge branch 'master' of github.com:karpathy/llm.c 2024-06-19 01:29:23 +00:00
Andrej Karpathy 7534c3fdf4 minor simplification inside gpt2_update of copy paste code 2024-06-19 01:29:11 +00:00
Andrej 4577d58cba Merge pull request #611 from rosslwheeler/feature/dataloader_test
Bug fixes and warning fixes for test_dataloader.c
2024-06-18 18:08:05 -07:00
Andrej 88f0543c38 Merge pull request #594 from karpathy/feature/hf_eval
add scripts to export to HF and run Eleuther evals
2024-06-18 18:04:14 -07:00
Andrej Karpathy 1a66151a80 small fixes and readme 2024-06-19 01:03:36 +00:00
Ross Wheeler 24923829c4 Adding suggested changes
Ternary change to test_dataloader and header cleanup
2024-06-18 10:38:33 -07:00
Erik Schultheis b3b4dabab0 small touchups 2024-06-18 20:03:28 +03:00
Erik Schultheis d95313d0a9 explicit sync 2024-06-18 20:02:07 +03:00
Erik Schultheis ec4424013a made buffer size more easily configurable 2024-06-18 20:00:25 +03:00
Erik Schultheis 3ebd5abe82 fixup tests 2024-06-18 19:58:03 +03:00
Andrej 6ecc52e3fc Merge pull request #614 from gordicaleksa/more_strict_tests
Stricter FP32 tests
2024-06-18 09:29:13 -07:00
Andrej 03cdca1b0d Merge pull request #601 from gordicaleksa/fix_stochastic_rounding_2
Fix stochastic rounding in encoder backward kernel
2024-06-18 09:28:39 -07:00
Andrej e12b172610 Merge pull request #613 from gordicaleksa/fix_dataloader
Fix dataloader & determinism testing
2024-06-18 08:31:01 -07:00
Aleksa Gordic e5d51c09e2 Add a comment for why pass a unique seed per param 2024-06-18 17:01:03 +02:00
Aleksa Gordic 6b2b3dc33f Pass unique seeds in stochastic rounding encoder bwd 2024-06-18 16:40:23 +02:00
Aleksa Gordic f0ffb648cc Represent loss threeshold using scientific notation 2024-06-18 16:00:00 +02:00
Aleksa Gordic 8f7f20506e Stricter FP32 tests 2024-06-18 15:41:38 +02:00
Aleksa Gordic 1120fa4465 Remove -1, point at current sample 2024-06-18 13:35:17 +02:00
Erik Schultheis 37c3815ede fix path 2024-06-18 14:35:05 +03:00
Aleksa Gordic 077f9bdd59 Fix dataloader & determinism testing 2024-06-18 13:23:13 +02:00
Erik Schultheis 33136e0b08 fix compilation with clang 2024-06-18 14:16:59 +03:00
Erik Schultheis dbeb8fc551 added missing checks 2024-06-18 13:29:12 +03:00
Erik Schultheis 188e7274e6 add unit test to CI 2024-06-18 13:23:56 +03:00
Erik Schultheis 593a71c374 use new functions for checkpointing 2024-06-18 13:23:56 +03:00
Erik Schultheis 04234d0c24 utility functions for device <-> disk IO 2024-06-18 13:23:56 +03:00
Ross Wheeler 30095bcb0b Fixes on accessing past memory buffer 2024-06-17 21:37:35 -07:00
Andrej Karpathy efbbdc8f46 Merge branch 'master' of github.com:karpathy/llm.c 2024-06-18 03:03:37 +00:00
Andrej Karpathy 4fd2df8870 Merge branch 'ngc92-check-determinism' 2024-06-18 03:03:06 +00:00
Andrej Karpathy 5b4fb0d457 delete intermediate files produced by test code 2024-06-18 03:02:53 +00:00
Andrej Karpathy 75b844898d delete warning using an explicit cast to size_t 2024-06-18 03:01:06 +00:00
rosslwheeler 4e0221de0d Merge branch 'karpathy:master' into feature/dataloader_test 2024-06-17 18:18:23 -07:00
Ross Wheeler 37023dd9a1 Cleanup for data loader 2024-06-17 18:15:29 -07:00
Andrej Karpathy 19f2c609eb Merge branch 'check-determinism' of https://github.com/ngc92/llm.c into ngc92-check-determinism 2024-06-17 22:16:21 +00:00
Andrej a33a49a7b6 Merge pull request #606 from gordicaleksa/minor_refactor_dataloader
Minor refactor of dataloader
2024-06-17 14:45:36 -07:00
Erik Schultheis 19a34d8cdd fix dataloader 2024-06-18 00:25:41 +03:00
Erik Schultheis c06ef28a0d also test dataloader as part of test script 2024-06-18 00:24:42 +03:00
Andrej f104009bc7 Merge pull request #600 from gordicaleksa/add_faster_layernorm_kernel
Use faster kernel for LayerNorm forward
2024-06-17 14:19:22 -07:00
Aleksa Gordic 88d4eff946 Swap kernel 5 for 3 for fallback 2024-06-17 23:05:01 +02:00
Erik Schultheis 674462adfe added missing check 2024-06-17 23:47:30 +03:00
Erik Schultheis 4d35bac1ae added a determinism check 2024-06-17 23:47:29 +03:00
Aleksa Gordic 7b5ee07e18 Minor refactor of dataloader 2024-06-17 21:53:18 +02:00
Andrej Karpathy d05e848010 fix pytorch code bug forward needs to know about backward_grad_sync 2024-06-17 19:49:36 +00:00
Andrej be0e10b518 Permuted DataLoader
Merge pull request #599 from karpathy/feature/permuted_dataloader

This DataLoader was shown to work slightly better than non-permuted version in a small 124M test.
Will permute all the examples on each epoch, possibly also improving multi-epoch performance.
2024-06-17 12:28:03 -07:00
Andrej Karpathy 14d9cd671e resolve merge conflict in staate file between use_master_weights and DataLoader should_shuffle 2024-06-17 19:16:42 +00:00
Andrej Karpathy 93d1450bdd malloc -> mallocCheck in dataloader.h 2024-06-17 18:56:24 +00:00
Andrej Karpathy de53e71729 small touchups 2024-06-17 18:55:16 +00:00
Andrej d57f112e16 Merge pull request #604 from gordicaleksa/minor_refactor_and_fix_guards
Fix guarding logic, refactor
2024-06-17 10:17:24 -07:00
Aleksa Gordic 9b5ec8f719 Set v memory to random floats >=0 2024-06-17 17:44:00 +02:00
Aleksa Gordic a0a8795be1 Fix guarding logic, refactor 2024-06-17 14:12:05 +02:00
Aleksa Gordic f845efd49a Add nvtx range to layernorm forward 2024-06-17 09:53:58 +02:00
Aleksa Gordic c6040eef5f Even better kernel - modification of what we use for fused residual fwd 2024-06-17 09:53:58 +02:00
Aleksa Gordic 0f36cdd15b Use faster kernel 5 for layernorm forward 2024-06-17 09:53:58 +02:00
Andrej fa7f2b6763 Merge pull request #602 from gordicaleksa/hotfix_global_norm_assert
Hotfix - proper handling of max num of block sums
2024-06-16 15:14:47 -07:00
Aleksa Gordic 0eec8cee4d Patch numpy version 2024-06-16 23:59:29 +02:00
Aleksa Gordic 8f724dbb72 Add a minor note that code is coupled 2024-06-16 23:54:09 +02:00
Aleksa Gordic 9caf2c7c2f Proper handling of max num of block sums 2024-06-16 23:49:34 +02:00
Andrej 13c893ffec Merge pull request #561 from lancerts/fix-warning
Fix the compiler warnings and errors
2024-06-16 09:25:51 -07:00
Andrej 1973631c6b Merge pull request #583 from vyom1611/FLT_MAX-fix
Replaced hard-coded max float with FLT_MAX
2024-06-16 09:25:03 -07:00
Andrej 3e7d89b81e Merge pull request #522 from gordicaleksa/add_master_weights_to_state
Add master weights to resume state
2024-06-16 09:22:10 -07:00
Andrej 8f40ef17e3 Merge pull request #597 from gordicaleksa/fix_stochastic_rounding
Fix stochastic rounding
2024-06-16 09:13:37 -07:00
Andrej 89e5d51fa4 Merge pull request #596 from gordicaleksa/fix_copy_casting
Fix copy and casting from spawning excessive threads
2024-06-16 09:02:14 -07:00
Andrej a4852cbafb Merge pull request #598 from gordicaleksa/add_new_dev_kernel
Add global norm prod kernels to dev
2024-06-16 09:00:01 -07:00
Aleksa Gordic d5601d16da Fix stuff in layernorm forward 2024-06-16 09:34:01 +02:00
Aleksa Gordic d003274b7c Remove warp size - moved to common.h 2024-06-15 20:23:59 +02:00
Aleksa Gordic 09910acc0b Remove blockreduce from classifier fused, move to common.h 2024-06-15 19:55:11 +02:00
Aleksa Gordic 0f112f1b62 Fix CI complaining about default args 2024-06-15 19:39:05 +02:00
Andrej Karpathy 1e7c31eef0 add test for shuffled, permuted, multi-process dataloader 2024-06-15 16:36:43 +00:00
Andrej Karpathy f089f90bcf add multiprocess test simple 2024-06-15 16:21:05 +00:00
Andrej Karpathy d4255f1d19 first draft of a dataloader test 2024-06-15 15:56:56 +00:00
Aleksa Gordic f3195d19af Expose deviceProp as a global var 2024-06-15 17:40:14 +02:00
Aleksa Gordic 0f6853d399 Add global norm prod kernels to dev 2024-06-15 17:29:14 +02:00
Aleksa Gordic ff14b0819f Make sure each parameter was rounded with a unique seed 2024-06-15 12:02:39 +00:00
Aleksa Gordic 20f0e7f4e8 Fix copy and casting from spawning excessive threads 2024-06-15 08:15:38 +00:00
Andrej Karpathy 60ed946641 Merge branch 'dataloader_rand' of https://github.com/gordicaleksa/llm.c into gordicaleksa-dataloader_rand 2024-06-14 22:10:55 +00:00
Aleksa Gordic c8cd023063 Full determinism 2024-06-14 21:44:20 +00:00
Andrej Karpathy d7ea4da298 add scripts to export to HF and run Eleuther evals 2024-06-14 19:49:05 +00:00
Andrej 0e69e3ae46 Merge pull request #592 from gordicaleksa/add_to_mfu_table
Add H100 PCIe to MFU table
2024-06-14 09:29:00 -07:00
Vyom Sharma ee4b056ddf Merge branch 'karpathy:master' into FLT_MAX-fix 2024-06-14 16:32:37 +05:30
Aleksa Gordic 0f144e3052 Support 100% determinism, fix race condition, update global norm 2024-06-14 11:57:07 +02:00
Aleksa Gordic 269b54ae45 Allocate memory in load state 2024-06-14 10:58:17 +02:00
Aleksa Gordic 2d23d1f08e Save master weights 2024-06-14 10:58:17 +02:00
Aleksa Gordic 328641f728 Minor refactor 2024-06-14 09:57:19 +02:00
Aleksa Gordic ef3afcb1d6 Fix num samples bug 2024-06-14 09:37:38 +02:00
Aleksa Gordic 2cd4d3985d Add H100 PCIe 2024-06-14 09:23:55 +02:00
Andrej 2bc0b474c5 Merge pull request #590 from karpathy/feature/delete_backward_act
consolidate memory
2024-06-13 18:21:25 -07:00
Andrej Karpathy cbe82e783a consolidate memory by deleting the grad activation struct, move some of the last pieces of memory to the forward pass 2024-06-14 00:54:58 +00:00
Andrej 747a5792b7 Merge pull request #582 from vyom1611/master
Benchmark modal script fixed - profiling and cuDNN (Issue #504 and PR #510 fixes)
2024-06-13 16:47:20 -07:00
Andrej 3bedc9853c Merge pull request #589 from karpathy/feature/overlap-backward-reduce
Feature/overlap backward reduce
2024-06-13 16:31:48 -07:00
Andrej Karpathy 43a903e1b4 revert old printf0, minor touchups and comments, add assert 2024-06-13 23:18:22 +00:00
Andrej Karpathy 6515f838d9 Merge branch 'comp-and-comms' of https://github.com/ngc92/llm.c into ngc92-comp-and-comms2 2024-06-13 22:34:50 +00:00
Erik Schultheis 3fc8d6192f strided all-gather 2024-06-14 01:33:43 +03:00
Aleksa Gordic 6c10cf2791 Remove local paths 2024-06-13 23:51:07 +02:00
Aleksa Gordic efd9df4c1b Remove bool, add permutation on every shard advance 2024-06-13 23:36:24 +02:00
Aleksa Gordic 7e9328f724 Remove redundant headers, refactor 2024-06-13 23:15:39 +02:00
Aleksa Gordic 1329a2c80c Change type from size_t to int to save space 2024-06-13 23:04:15 +02:00
Aleksa Gordic 814b83bbf1 Refactor 2024-06-13 22:43:57 +02:00
Aleksa Gordic 1dafa60ad9 Hardcode shuffle info in all train scripts 2024-06-13 21:27:42 +02:00
Andrej Karpathy 8367a56c03 Merge branch 'comp-and-comms' of https://github.com/ngc92/llm.c into ngc92-comp-and-comms2 2024-06-13 19:25:45 +00:00
Aleksa Gordic b7933ef62c Modify state save/load 2024-06-13 21:22:16 +02:00
Andrej afcc0a4f71 Merge pull request #586 from gordicaleksa/readme_debug_note
Add a debugging tip to README
2024-06-13 11:30:20 -07:00
Aleksa Gordic 8f61f48cba Add random inter/intra shard data shuffling 2024-06-13 20:21:49 +02:00
Aleksa Gordic 2faa5e2b6c Update other calls with null 2024-06-13 20:21:49 +02:00
Aleksa Gordic d57a52408e Make C compiler happy - no default args 2024-06-13 20:21:49 +02:00
Aleksa Gordic c3db1bb112 Set default seed value to null 2024-06-13 20:21:49 +02:00
Aleksa Gordic af0b2f51c0 Start from a random position in a random shard 2024-06-13 20:21:49 +02:00
Aleksa Gordic 2a6797b1cd Add debugging tip to README 2024-06-13 13:58:14 +02:00
Erik Schultheis f071efd94f stride-based adam and copy-and-cast 2024-06-13 02:01:26 +03:00
Erik Schultheis 99965a1e3e fix typo 2024-06-13 00:36:34 +03:00
vyom1611 e2a18d9385 Fixed missing semi-colon 2024-06-13 00:59:15 +05:30
vyom1611 dc8d8de06a Replaced hard-coded max float with FLT_MAX 2024-06-13 00:56:53 +05:30
Erik Schultheis 9d9da3b209 fewer global norm kernel calls 2024-06-12 21:55:01 +03:00
Vyom Sharma afa7978227 Merge branch 'karpathy:master' into master 2024-06-13 00:19:16 +05:30
Andrej 95cef79c17 Merge pull request #578 from ngc92/matmul
Remove cublaslt from fp32cu versions
2024-06-12 09:05:04 -07:00
Andrej Karpathy 42dff87939 clarification we trained with 150B tokens 2024-06-12 15:51:14 +00:00
Andrej Karpathy 963aa6828e add gpt-2 774M repro script 2024-06-12 15:33:19 +00:00
Erik Schultheis 8818a75157 fixup to profiling script 2024-06-12 13:16:25 +03:00
Erik Schultheis 368644959a fix non-multi-gpu case 2024-06-12 13:16:19 +03:00
Erik Schultheis 8cba9ba0c0 cleanup 2024-06-12 13:16:19 +03:00
Erik Schultheis ecca976132 re-enable reduce-scatter operation 2024-06-12 13:16:19 +03:00
Erik Schultheis ecdb0ca86e refactored code 2024-06-12 13:16:19 +03:00
Erik Schultheis fdcb30256d first working ZeRO-1 2024-06-12 13:16:19 +03:00
Erik Schultheis 6d2c72cb4b rearranged adam code 2024-06-12 13:16:19 +03:00
Erik Schultheis 50a9411114 overlap backward with nccl for non-ZeRO runs 2024-06-12 13:16:19 +03:00
Erik Schultheis d6793645f5 replace cublaslt with custom matmul 2024-06-12 09:42:19 +03:00
Erik Schultheis 7f3f3acf16 shared memory matmul 2024-06-11 21:03:43 +03:00
Erik Schultheis 071c977c79 vectorized memory access 2024-06-11 00:46:51 +03:00
Erik Schultheis bb89efdcb1 register reuse 2024-06-11 00:41:08 +03:00
Vyom Sharma 8a3a29ecb7 Merge branch 'karpathy:master' into master 2024-06-10 13:54:29 +05:30
Andrej 615ec0becd Merge pull request #523 from ngc92/testability
Testability
2024-06-09 11:02:39 -07:00
Erik Schultheis 02e30faef2 bugfix for LN recomputation 2024-06-08 21:25:27 +03:00
Erik Schultheis 8fbfba2864 allow setting master weights and recompute in tests 2024-06-08 21:03:07 +03:00
Erik Schultheis 1bf33cc66d make sure failed testing returns failure status 2024-06-08 21:03:07 +03:00
Andrej d2482ff1ad Merge pull request #570 from rosslwheeler/feature/ubicloud-gpu-ci
Adding GPU CI workflow file
2024-06-08 11:01:03 -07:00
Andrej 637c1b69c4 Merge pull request #552 from karpathy/feature/streams
Feature/streams
2024-06-08 09:25:56 -07:00
Andrej Karpathy ee6b3c9325 take out Async copies and memsets 2024-06-08 16:17:02 +00:00
Andrej 077451927f Merge pull request #546 from kmyusk/master
Fix periodic inference during training
2024-06-07 21:37:51 -07:00
Andrej d396cd18b7 Merge pull request #567 from gordicaleksa/improve_logdir_check_logic
Minor fix - better handling of out log dir
2024-06-07 09:22:32 -07:00
Andrej 58415736a4 Merge pull request #568 from gordicaleksa/improve_logger
Improve logger, add grad norm + learning rate
2024-06-07 09:20:29 -07:00
Aleksa Gordic 40e7289916 Make it cross-platform, support Windows 2024-06-07 17:49:49 +02:00
Aleksa Gordic d3790c64a1 Change default path 2024-06-07 17:38:01 +02:00
Aleksa Gordic 4090ca8fa1 Add learning rate and grad norm logging 2024-06-07 17:34:42 +02:00
Aleksa Gordic aa3d8aca86 Minor fix - better handling of out log dir 2024-06-07 17:06:53 +02:00
lancerts 089fdd7f23 Merge branch 'master' into fix-warning 2024-06-07 07:17:52 -07:00
vyom1611 ecb13ea74d missing parameter in comments added 2024-06-07 16:14:59 +05:30
vyom1611 009be4dcab Merge remote-tracking branch 'origin/master' 2024-06-07 14:21:35 +05:30
vyom1611 58c59064ec Most updated working modal script 2024-06-07 14:21:13 +05:30
Vyom Sharma c24ef4a4b4 Merge branch 'karpathy:master' into master 2024-06-07 14:20:55 +05:30
Ross Wheeler 73506dff2d Adding GPU CI workflow file 2024-06-06 23:43:01 -07:00
Andrej 8a0809b39c Merge pull request #560 from gordicaleksa/refactor_cuda_kernels
Fix mem leaks, reduce memory, refactor
2024-06-06 20:10:52 -07:00
lancerts 4b41087c64 fix the compiler warnings and errors 2024-06-06 11:46:40 -07:00
Aleksa Gordic 7b99ac96a5 Fix mem leaks, reduce memory, refactor 2024-06-06 19:58:50 +02:00
Andrej d931fc0507 -std=c++17 default, merge PR #559 from ngc92/ubuntu20.04
set compile flag and add ci check
2024-06-06 09:23:07 -07:00
Erik Schultheis 875196b672 add c++17 flag also to dev/cuda 2024-06-06 19:02:45 +03:00
Andrej 067aadc226 Merge pull request #553 from gordicaleksa/refactor_trimat
Refactor trimat
2024-06-06 08:56:12 -07:00
Erik Schultheis bb9c895400 fix 2024-06-06 12:02:33 +03:00
Erik Schultheis f4dbd88151 enable cudnn 2024-06-06 11:59:45 +03:00
Erik Schultheis 52ea988860 fixup image name 2024-06-06 11:01:10 +03:00
Erik Schultheis 43f013631b set compile flag and add ci check 2024-06-06 10:55:33 +03:00
Aleksa Gordic c116fbfa29 Fix q,k bug, move all kernels to q@k^t form 2024-06-05 22:22:55 +02:00
Aleksa Gordic 875d362381 Refactor trimat 2024-06-05 17:33:26 +02:00
Vyom Sharma 1e8efe94c0 Update Makefile - revert back 2024-06-05 12:02:09 +05:30
Andrej Karpathy 6d14c29df5 streams, everything right now is on main stream, almost 2024-06-05 01:01:24 +00:00
Andrej Karpathy d5f7458e03 Merge branch 'stream-2' of https://github.com/ngc92/llm.c into ngc92-stream-2 2024-06-05 00:23:19 +00:00
Andrej Karpathy b1171051dd Merge branch 'master' of github.com:karpathy/llm.c 2024-06-05 00:19:47 +00:00
Andrej Karpathy 9a07817380 fix minor bugs in MFU calculation 2024-06-05 00:15:08 +00:00
Erik Schultheis 9af7628063 moved remaining memcpy/memset in the training loop to main_stream 2024-06-05 02:06:32 +03:00
Andrej Karpathy 8940dcdcb1 Merge branch 'stream-2' of https://github.com/ngc92/llm.c into ngc92-stream-2 2024-06-04 22:52:26 +00:00
Erik Schultheis bc35d3438e naming the stream 2024-06-05 01:50:16 +03:00
Andrej 0ee2152106 Merge pull request #549 from rosslwheeler/feature/new-branch-ci-trigger
Trigger CI on new branch creation
2024-06-04 15:42:08 -07:00
Erik Schultheis 0de54d10da moved last remaining kernel to main stream 2024-06-05 01:32:25 +03:00
Andrej Karpathy 04b59c87be tune comments 2024-06-04 22:16:23 +00:00
Andrej 894ada5041 Merge pull request #551 from gordicaleksa/fix_loss_bug
Fix PyTorch DDP loss bug
2024-06-04 15:10:10 -07:00
Aleksa Gordic 6f20ee9f8e Remove * 2 for total batch size 2024-06-04 23:42:54 +02:00
Aleksa Gordic 34766253cf Move all reduce outside of the gradacc loop for efficiency 2024-06-04 23:39:52 +02:00
Aleksa Gordic 724085ecd6 tmp 2024-06-04 23:23:11 +02:00
Erik Schultheis 4c369eea0f overlap validity check and data movement 2024-06-04 22:47:54 +03:00
Erik Schultheis ca7caac5c5 move all kernels into a main_stream in preparation for future parallelism 2024-06-04 22:35:12 +03:00
vyom1611 ab8b040fdd Merge remote-tracking branch 'origin/master' 2024-06-04 23:39:50 +05:30
vyom1611 4564a24c69 Modal script updated for cudnn and gcc version match 2024-06-04 23:39:36 +05:30
Vyom Sharma 453fde77fd Merge branch 'karpathy:master' into master 2024-06-04 23:39:10 +05:30
Ross Wheeler d51c8db395 Trigger CI on new branch creation
With this, you can now trigger workflows in VS Code and GitHub manually.
2024-06-04 10:14:42 -07:00
Andrej 5379b1d9e9 Merge pull request #548 from gordicaleksa/fix_zerograds_bug
Fix zero grads bug
2024-06-04 09:51:44 -07:00
Aleksa Gordic 3e44d275d8 Update comment 2024-06-04 18:48:29 +02:00
Aleksa Gordic ec92ff64a2 Improve the design and comment 2024-06-04 18:47:23 +02:00
Aleksa Gordic 54ec686649 Fix zero grads bug 2024-06-04 18:40:54 +02:00
Andrej c73b4668f3 Merge pull request #547 from ngc92/includes
made implicit includes of standard headers explicit
2024-06-04 08:56:41 -07:00
Erik Schultheis b02aab1f16 made implicit includes of standard headers explicit 2024-06-04 13:22:16 +03:00
Yusik Kim 1ee0b437ca Fix periodic inference during training
Problem: Enabling the generation block crashes the program
Solution: Change the variable names for x and y so that it does not
affect downstream code
Testing: Runs as expected with tinyshakespeare

Signed-off-by: Yusik Kim <kmyusk@gmail.com>
2024-06-04 10:28:55 +02:00
Andrej cb40f458a3 Merge pull request #543 from rosslwheeler/bugfix/make-clean
Make clean fix and Windows cuDNN build fix
2024-06-03 21:06:35 -07:00
Ross Wheeler 4bfb23948d Make clean fix and Windows cuDNN build fix 2024-06-03 20:26:52 -07:00
Andrej Karpathy 45b0d1b378 copy and cast kernel move to cuda_utils, good general util 2024-06-03 23:14:59 +00:00
Andrej Karpathy d697eaecfd move the global norm kernel 2024-06-03 23:12:20 +00:00
Andrej Karpathy 8c52a652b4 move adamw kernel to llmc 2024-06-03 23:09:42 +00:00
Andrej Karpathy e484c650b5 layernorm and residual, yolo push to master 2024-06-03 23:02:46 +00:00
Andrej 7d50d3318a Merge pull request #541 from karpathy/move/matmul
move matmul
2024-06-03 15:57:09 -07:00
Andrej Karpathy 094988738b adjust the comment on extern deviceProp to be more descriptive 2024-06-03 22:57:00 +00:00
Christopher ec3921047c updated the script to run the same number of threads as files 2024-06-03 22:54:25 +00:00
Andrej Karpathy 645feac1e7 move matmul 2024-06-03 22:54:07 +00:00
Christopher 3936492620 Forgot the folder name in the readme 2024-06-03 22:51:57 +00:00
Christopher 027a43e0f5 Adds simplified way to download the already tokenized tinyshakespeare dataset and gpt2 weights 2024-06-03 22:49:50 +00:00
Andrej 299d375fbc Merge pull request #540 from karpathy/move/fused_classifier
move fused classifier
2024-06-03 15:38:24 -07:00
Andrej Karpathy 90a7bfecf1 move fused classifier 2024-06-03 22:35:20 +00:00
Andrej dd4191fbd5 Merge pull request #539 from karpathy/move/attention
move attention
2024-06-03 15:23:27 -07:00
Andrej Karpathy 2186cd6fe7 move attention 2024-06-03 22:19:47 +00:00
Andrej 3427983615 Merge pull request #537 from karpathy/move/gelu
move gelu to llmc
2024-06-03 15:08:22 -07:00
Andrej Karpathy d8d55e0a08 move gelu to llmc 2024-06-03 21:57:02 +00:00
Andrej Karpathy 49b9776e50 also move the launchers 2024-06-03 21:37:18 +00:00
Andrej Karpathy 93546aed01 move encoder, kernel 1/N 2024-06-03 21:21:28 +00:00
Andrej 3392671826 Merge pull request #535 from ngc92/speed-up-compile
moved device utilities into a separate file, and follow-up cleanup
2024-06-03 14:11:10 -07:00
Erik Schultheis 67502c7f8e moved device utilities into a separate file, and follow-up cleanup 2024-06-03 23:58:47 +03:00
Andrej 8cc39a794b Merge pull request #528 from gordicaleksa/refactor_evalloader
Refactor eval loader - remove 2 unnecessary malloc calls per iteration
2024-06-03 13:54:52 -07:00
Andrej 2a6a5b2c6c Merge pull request #529 from gordicaleksa/improve_start_training_print_info
Improve start training print info
2024-06-03 13:53:32 -07:00
Andrej 0cd9f1ac21 Merge pull request #530 from gordicaleksa/minor_refactor
Minor refactor: remove hardcoded val, delete unused vars...
2024-06-03 13:52:42 -07:00
Andrej d17893b754 Merge pull request #532 from ngc92/ln-buffers
more efficient use of memory buffers for LN recomputation
2024-06-03 13:51:22 -07:00
Andrej ac931455a4 Merge pull request #533 from karpathy/feature/refactor_part2
refactor part 2
2024-06-03 13:44:59 -07:00
Andrej Karpathy 001ad78fe1 resolve deviceProp and cudaCheck dupes 2024-06-03 19:54:32 +00:00
Andrej Karpathy 020b15ad77 small more fixes 2024-06-03 19:32:32 +00:00
Andrej Karpathy 5d9403479c some fixes, still not done 2024-06-03 18:50:41 +00:00
Andrej Karpathy 804d0693c8 refactor part 2, moving stuff into common files so that we can later nicely separate out all the kernels as well 2024-06-03 18:10:33 +00:00
Andrej e40268457f Merge pull request #526 from karpathy/try/move_cudnn_to_llmc
attempt to move cudnn attention to the llmc directory
2024-06-03 09:32:31 -07:00
Erik Schultheis 0d16b51626 more efficient use of memory buffers for LN recomputation 2024-06-03 18:32:51 +03:00
Andrej Karpathy 6434009298 attempt to remove PATH_SEP from Makefile 2024-06-03 15:30:45 +00:00
Aleksa Gordic 42ca620306 Minor refactor: remove hardcoded val, delete unused vars... 2024-06-03 16:24:22 +02:00
vyom1611 9b540fa859 Makefile updated to match script 2024-06-03 18:59:57 +05:30
vyom1611 fbc1e4983c Benchmark modal script revamped - profiling and external libraries 2024-06-03 17:00:10 +05:30
vyom1611 0867653aaf Merge remote-tracking branch 'origin/master' 2024-06-03 16:48:58 +05:30
Aleksa Gordic 7413a1a315 Remove unnecessary load filename line 2024-06-03 12:33:33 +02:00
Aleksa Gordic 728762ceed Improve start of training print info 2024-06-03 12:31:33 +02:00
Aleksa Gordic d9e867678e Remove file size - not used 2024-06-03 12:14:58 +02:00
Aleksa Gordic e86bd62374 Refactor eval loader - remove 2 unnecessary malloc calls per iteration 2024-06-03 11:43:33 +02:00
Andrej Karpathy 41dc65a0ef attempt to move cudnn attention to the llmc directory 2024-06-03 03:23:04 +00:00
Andrej Karpathy 08fc45b492 delete copy paste code for gpt2 model init 2024-06-03 02:08:29 +00:00
Andrej Karpathy c827dd2f8e Merge branch 'ngc92-mfu' 2024-06-03 01:41:31 +00:00
Andrej Karpathy 1ac30591f2 add new file llmc/mfu.h for MFU calculations for different GPUs 2024-06-03 01:41:18 +00:00
Andrej 80e0c3b8d9 Merge pull request #512 from gordicaleksa/refactor_encoder_bwd_kernel
Remove redundant CPU computation in encoder bwd
2024-06-02 16:43:10 -07:00
Andrej cb5eff0578 Merge pull request #514 from ngc92/compile-fix
fix compilation with older nvcc
2024-06-02 16:42:31 -07:00
Andrej ff98d8bda7 Merge pull request #515 from gordicaleksa/fix_adamw_kernel3_num_threads
Use local params for num blocks in adamw_kernel3
2024-06-02 16:41:59 -07:00
Andrej 1ce30fdc4a Merge pull request #516 from gordicaleksa/update_random_doc
Add clarification on Box-Muller
2024-06-02 16:37:37 -07:00
Andrej 252584954e Merge pull request #519 from gordicaleksa/mem_leak
Fix mem leak
2024-06-02 16:34:54 -07:00
Andrej Karpathy 1e893f4e11 Merge branch 'gordicaleksa-fix_zero_stage1' 2024-06-02 23:28:22 +00:00
Andrej Karpathy 4944ac5743 adjust comments 2024-06-02 23:28:13 +00:00
Aleksa Gordic a8a5a444bc Add ; - ah c my old friend :) 2024-06-02 23:04:47 +02:00
Aleksa Gordic 9a455edc25 Improve comments 2024-06-02 23:03:41 +02:00
Aleksa Gordic b4e8bd7f66 Fix zero stage 1 norm computation - WIP 2024-06-02 22:45:26 +02:00
Aleksa Gordic df0a826ef3 Close file 2024-06-02 21:37:52 +02:00
Aleksa Gordic 30d81cce15 Fix mem leak 2024-06-02 21:28:42 +02:00
“eliebak” 67379a28cc add edu fineweb support, with 10B and 1000B version 2024-06-02 14:18:44 +00:00
Aleksa Gordic 79cbc03ccf Use local params for num blocks 2024-06-02 15:06:06 +02:00
Aleksa Gordic ae9c957273 Add clarification on Box-Muller 2024-06-02 15:04:45 +02:00
Erik Schultheis a902323f36 fix compilation with older nvcc 2024-06-02 12:40:28 +03:00
Aleksa Gordic 290c00a362 Remove redundant CPU computation 2024-06-01 18:27:24 +02:00
Andrej a26041af12 Merge pull request #495 from ChrisDryden/shared_memory
Removed unnecesary shared memory due to blockreduce using static defined shared memory
2024-06-01 08:12:21 -07:00
Andrej 9cf8c2f135 Merge pull request #509 from ngc92/inplace
Inplace gelu backward
2024-06-01 08:10:56 -07:00
Andrej bd0f036ee5 Merge pull request #511 from gordicaleksa/gordicaleksa_fix_docs
Update compile cmd in the dev/cuda README
2024-06-01 08:09:31 -07:00
Aleksa Gordic 8eafd40615 Update compile cmd in the dev/cuda README 2024-06-01 16:46:32 +02:00
vyom1611 e0abfae2bc Benchmark modal script revamped - profiling and external libraries 2024-06-01 19:06:42 +05:30
Erik Schultheis 761fb2cf2b rename gelu_backward to gelu_backward_inplace, and made it obvious in the code that this is an inplace operation 2024-06-01 14:31:26 +03:00
Erik Schultheis cc4bed0806 added missing checks 2024-06-01 14:30:40 +03:00
Andrej Karpathy 9c8edaffa7 fix recompute bug 2024-05-31 22:38:03 +00:00
Andrej Karpathy 0bc24c36e6 fix bug in resumption of dataloader, have to reload shard 2024-05-31 22:19:37 +00:00
Andrej Karpathy 1de3bef0ac fix argparse comment too 2024-05-31 22:06:48 +00:00
Andrej Karpathy 60a1037db9 Merge branch 'ChrisDryden-layernorm_backwards_mem_saving' 2024-05-31 22:04:44 +00:00
Andrej Karpathy f4874a9f63 make things a bit more consistent, keep recompute 1 as default, and add a comment 2024-05-31 22:04:34 +00:00
Andrej Karpathy 629a24020d Merge branch 'layernorm_backwards_mem_saving' of https://github.com/ChrisDryden/llm.c into ChrisDryden-layernorm_backwards_mem_saving 2024-05-31 21:47:28 +00:00
Andrej Karpathy ed37d9261b fix bug in readme 2024-05-30 21:39:47 +00:00
Andrej Karpathy b48224cb0c refactor the README file because it was getting crazy 2024-05-30 21:39:06 +00:00
Andrej Karpathy 0fddf06043 more docs on scripts 2024-05-30 21:06:36 +00:00
Andrej Karpathy fe36e40f2c also move pyrun script to scripts/ 2024-05-30 20:17:16 +00:00
Andrej d98d3ed82b code refactor, introduce llmc dir and scripts dir
The main file train_gpt2.cu is getting out of control. So here we are.
2024-05-30 13:13:04 -07:00
Andrej Karpathy 22d1607f6d delete these files, they are now in scripts/ folder and better 2024-05-30 20:11:59 +00:00
Andrej Karpathy fc40ffc4aa adjust readme to new scripts/ dir 2024-05-30 20:10:04 +00:00
Andrej Karpathy a0d5dbbda6 scripts folder 2024-05-30 20:00:07 +00:00
Andrej Karpathy 35aa110244 more refactor 2024-05-30 19:39:29 +00:00
Andrej Karpathy f8f9e49ba3 Merge branch 'master' into feature/libdir 2024-05-30 19:07:51 +00:00
Andrej Karpathy 5f547601a0 further hide the c pytorch bridge logic 2024-05-30 18:20:29 +00:00
Andrej Karpathy ca27eb8fa5 bugfix don't train on val batch 2024-05-30 18:18:05 +00:00
Christopher 59432fea19 modified the recalculation parameter to be based off of 2 instead of 1 2024-05-30 17:23:47 +00:00
Christopher e5719e18b3 Rearranged layernorm forward for 1 and 2 and used layer variables instead. Validated with test file 2024-05-30 16:28:03 +00:00
Christopher 95b1fad088 First step in creating option to recalculate layernorm activations in backwards pass 2024-05-30 16:28:03 +00:00
Andrej 0e2298247a Merge pull request #496 from ChrisDryden/patch-7
Nit: duplicated line in header of project
2024-05-30 09:13:30 -07:00
Andrej Karpathy 862c957d3a fix default number of gpus from 4 to 8 for consistency with run124M.sh 2024-05-30 14:50:09 +00:00
Andrej ed805f60be Merge pull request #494 from karpathy/feature/pytorch_parity
pytorch parity v1
2024-05-30 07:47:58 -07:00
Andrej Karpathy 5959851d2d small comments for pyrun124M.sh 2024-05-30 14:47:43 +00:00
Andrej Karpathy 9b935a8f04 fix bug with advancing position in dataloader 2024-05-30 13:16:07 +00:00
Chris Dryden 3bca31fde1 Nit: duplicated line in header of project 2024-05-29 20:03:13 -07:00
Christopher 5450632237 Removed unnecesary shared memory due to blockreduce using static defined shared memory 2024-05-30 03:01:52 +00:00
Andrej Karpathy 555d07502a more feature parity work, i think it's quite close now 2024-05-30 02:29:40 +00:00
Andrej Karpathy ea88a343b5 also updated pyrun124M.sh 2024-05-30 02:05:41 +00:00
Andrej Karpathy 03c2e85601 sharded dataloader for PyTorch 2024-05-30 02:05:06 +00:00
Andrej Karpathy 3fdf88e4d4 add logfile similar to C 2024-05-29 23:59:35 +00:00
Andrej Karpathy 083c50a52e small stylistic touchups that match the c code 2024-05-29 23:50:13 +00:00
Andrej Karpathy 1f6fdfbd9a add pyrun124M.sh 2024-05-29 23:30:27 +00:00
Andrej Karpathy c413bfe0ca pytorch parity v1 2024-05-29 23:25:55 +00:00
Andrej Karpathy 4cc8b1570b hardcode grad_clip to 1.0, delete confusion around legacy code that used -c flag for it, but became inactive in the if else block 2024-05-29 19:01:12 +00:00
Andrej 55f3665b64 Merge pull request #493 from vyom1611/master-2
Modal benchmarking script updated to replace deprecated calls
2024-05-29 07:50:10 -07:00
vyom1611 062f096da1 Modal benchmarking script updated to replace deprecated calls 2024-05-29 20:01:41 +05:30
Andrej 2db7716d68 Merge pull request #476 from lancerts/layernorm-bk-fix
Fix "Mismatch of dweight at layernorm_backward.cu #428"
2024-05-29 07:30:52 -07:00
Andrej c3a65a556f Merge pull request #487 from huoyushequ/patch-1
`softmax_autoregressive_backward_kernel` does not use share memory in the kernel
2024-05-28 18:52:38 -07:00
QIAOQIAN 27cf5223e0 softmax_autoregressive_backward_kernel does not use share memory in the kernel
we do not need to launch the kernel with 256 bytes share memory, so remove it
2024-05-29 09:23:46 +08:00
Erik Schultheis daa95e04ae ampere-consumer 2024-05-29 02:27:11 +03:00
Erik Schultheis de7a9b83cd more GPUs 2024-05-29 02:20:31 +03:00
Erik Schultheis 184cc9d10a added H100 SMX version 2024-05-29 01:31:16 +03:00
Erik Schultheis e2b411362e include string view 2024-05-29 00:54:23 +03:00
Erik Schultheis 9374d751de MFU for other GPUs 2024-05-29 00:48:37 +03:00
Andrej Karpathy a664aa3094 include datasets as req 2024-05-28 21:43:11 +00:00
Andrej afe40be994 Merge pull request #479 from rosslwheeler/feature/ci-manual-trigger
Adding manual trigger for CI
2024-05-28 12:46:27 -07:00
Andrej Karpathy 88dd670c42 link to recent post from readme 2024-05-28 17:31:31 +00:00
Andrej Karpathy a6a2f042f8 keep -y 1 flag 2024-05-28 15:24:30 +00:00
Andrej Karpathy a7f01088dc repro scripts for 124M and 350M models 2024-05-28 15:08:04 +00:00
Andrej Karpathy 037ec67a5d small tweak to launch 2024-05-28 15:07:07 +00:00
Andrej Karpathy 0c83a23188 tweak the logfile parser and visualizer 2024-05-28 01:38:15 +00:00
Ross Wheeler c60d7de466 Adding manual trigger for CI 2024-05-27 16:50:06 -07:00
lancerts 47c670b885 amend 2024-05-27 15:13:11 -07:00
lancerts c4c985cb87 spell explicitly uint to unsigned int 2024-05-27 14:44:47 -07:00
lancerts f27ca4df41 fix the issue Mismatch of dweight at layernorm_backward.cu 2024-05-27 14:39:57 -07:00
Andrej Karpathy 888c0d4be3 also update the other files 2024-05-27 21:15:09 +00:00
Andrej Karpathy eed242702a experiment with adding the llmc lib directory 2024-05-27 21:10:28 +00:00
Andrej b67880a879 Merge pull request #471 from karpathy/feature/resume
resume training v1
2024-05-27 13:38:52 -07:00
Andrej Karpathy 24a8c03cb6 Merge branch 'master' into feature/resume 2024-05-27 20:30:30 +00:00
Andrej Karpathy 69d0583ac8 conditionally include dirent on not windows 2024-05-27 20:28:52 +00:00
Andrej 3b4d8bece6 Merge pull request #474 from rosslwheeler/Windows-dirent-fix
Adding dirent.h support to unistd.h
2024-05-27 13:27:58 -07:00
Ross Wheeler bb80f93618 Minor cleanup
Added static inline to the functions
2024-05-27 13:16:34 -07:00
Ross Wheeler 758acdffbf Adding dirent.h support to unistd.h 2024-05-27 13:08:07 -07:00
Andrej Karpathy 63f0e25f5a make compiler happy 2024-05-27 19:32:31 +00:00
Andrej Karpathy b75738c4eb resume optimization, seems to be working 2024-05-27 18:02:51 +00:00
Andrej Karpathy f93a30fbbf more careful with conditional MPI use 2024-05-27 15:54:55 +00:00
Andrej Karpathy d295cb8d81 part 1 of v1 of resume training functionality, writes the files but doesn't load them yet, coming up in a bit 2024-05-27 15:49:03 +00:00
Andrej 12999f7082 Merge pull request #470 from ngc92/bugfix
bugfix for layernorm correctness (dev/cuda did not trip, but end to end test `make test_gpt2cu PRECISION=FP32 && ./test_gpt2cu
` did)
2024-05-27 06:39:56 -07:00
Erik Schultheis 7dc3b7d7dc bugfix 2024-05-27 16:36:44 +03:00
Andrej 4d6f80e02b Merge pull request #469 from ahrefs/zero1-torch
Optional Zero1 changes for pytorch training
2024-05-27 06:04:00 -07:00
Andrej c3097ba308 Merge pull request #468 from ngc92/layernorm-bw-deadlock
Layernorm deadlock
2024-05-27 05:48:42 -07:00
Erik Schultheis b913905f84 fix out-of-bounds access for encoder_backward: we were allocating too little memory (/ -> DIV_CEIL) 2024-05-27 13:15:36 +03:00
Erik Schultheis 841c4dec03 update main training file 2024-05-27 13:15:36 +03:00
Erik Schultheis 5f73ecfcf1 fix C % 256 != 0 2024-05-27 13:15:36 +03:00
Erik Schultheis d35daf19c9 fail fast and hard; don't go into the deadlock 2024-05-27 13:15:36 +03:00
Erik Schultheis 3b6808269a some comments and optimized shared memory amount 2024-05-27 13:15:36 +03:00
Erik Schultheis 9eba86f0a5 fully vectorized smem access 2024-05-27 13:15:36 +03:00
Erik Schultheis cb8cc25e3a utilities for Packed and more vectorization 2024-05-27 13:15:36 +03:00
Erik Schultheis 53ee3297dd use vectorized access to shared memory 2024-05-27 13:15:36 +03:00
Chinthaka Gamanayakege f2809a6bdb zero changes for torch 2024-05-27 06:44:54 +00:00
Andrej Karpathy ee6fa5ec9f fill in the other model sizes 2024-05-26 18:04:43 +00:00
Andrej Karpathy 5ec4b774f2 simple jupyter notebook example on how to parse and vis the logfiles 2024-05-26 17:40:49 +00:00
Andrej Karpathy f145c1b101 fix the other files due to dataloader change 2024-05-26 17:01:32 +00:00
Andrej Karpathy c6680a25ae fix a bug where the max_steps was not accounting for the gradient accumulation, meaning that e.g. with grad accum 2, we'd accidentally train for 2 epochs instead of just 1 epoch. also simplify the interface a bit and the code make more readable etc 2024-05-26 16:49:30 +00:00
Andrej 5d65c5f52e Merge pull request #464 from karpathy/attempt_fix_zero1
Attempt fix zero1
2024-05-25 17:21:04 -07:00
Andrej Karpathy 05fb642ae1 attempt fix zero 1 2024-05-25 23:49:16 +00:00
Erik Schultheis 7ba3ef6bac make sure to run Adam only for the range of parameters that is available on the local device 2024-05-25 23:20:42 +03:00
Andrej Karpathy d0c014372c add ability to export fineweb100B as well 2024-05-25 17:21:51 +00:00
Andrej Karpathy f5d041ab0c careful with NULL and checkpoint correctly 2024-05-25 16:32:18 +00:00
Andrej Karpathy 273925d10d save checkpoint files to log directory, and small readability changes 2024-05-25 16:05:29 +00:00
Andrej Karpathy 90a745c40c add MFU (model flops utilization) calculation and print it. also refactor the printing a bit to make it nice. this calculation only works for A100 fp16 we should think through how to generalize it a bit 2024-05-25 14:30:59 +00:00
Andrej ab090a80c5 Merge pull request #461 from karpathy/feature/write_checkpoints
fix bug check null
2024-05-25 07:01:26 -07:00
Andrej Karpathy a13b2bd10b fix bug check null 2024-05-25 13:50:56 +00:00
Andrej fe698b32f1 Merge pull request #457 from karpathy/feature/write_checkpoints
add checkpoint function write to file
2024-05-24 17:56:53 -07:00
Andrej Karpathy 2a0f78d617 attempt to fix windows episode 2 2024-05-25 00:44:16 +00:00
Andrej Karpathy 5b96d91893 attempt to fix windows? 2024-05-25 00:39:40 +00:00
Andrej Karpathy f2ee356a06 resolve merge conflict 2024-05-25 00:33:12 +00:00
Andrej 4ff0412751 Merge pull request #450 from rosslwheeler/dataloader_win_fixes
Add glob() for windows
2024-05-24 17:30:35 -07:00
Andrej 6a7fd56d7c Merge pull request #439 from lancerts/matmul-fix
Fix the unsupported block_size in matmul_backward_bias kernel 1
2024-05-24 17:26:37 -07:00
Andrej 5ef2c04a39 Merge pull request #458 from karpathy/feature/weight_decay_2d_only
add weight decay
2024-05-24 17:22:49 -07:00
Andrej Karpathy 9f08882051 add weight decay, but only for 2D tensors, as done in GPT series and in general too. this forces us to break up our adamw kernel again into one call per tensor, so there is a small throughput hit, of about 0.5% for me. but we have to break up this kernel in near future anyway 2024-05-25 00:14:10 +00:00
Ross Wheeler 79738d2ca4 fixed mkdir change 2024-05-24 17:05:59 -07:00
Ross Wheeler e5083be900 Moved windows glob() over to dev/unistd.h
Added header guard and changed long->int64_t in dataloader.h
2024-05-24 16:56:52 -07:00
rosslwheeler 1ec081e64e Merge branch 'karpathy:master' into dataloader_win_fixes 2024-05-24 16:19:45 -07:00
Andrej Karpathy e60c4843a3 Merge branch 'ngc92-cpu-matmul' 2024-05-24 23:10:34 +00:00
Andrej Karpathy 3221e4b2d2 small cosmetic changes 2024-05-24 23:10:17 +00:00
Andrej Karpathy 6e4296fb9a Merge branch 'cpu-matmul' of https://github.com/ngc92/llm.c into ngc92-cpu-matmul 2024-05-24 22:59:24 +00:00
rosslwheeler 5e0fa452e0 Merge branch 'karpathy:master' into dataloader_win_fixes 2024-05-24 15:48:45 -07:00
Andrej Karpathy b8674cedf8 Merge branch 'ademeure-deterministic_encoder_v2' 2024-05-24 22:18:26 +00:00
Andrej Karpathy 8b951056b5 resolve merge conflict 2024-05-24 22:04:02 +00:00
Andrej Karpathy 5ddb06182c Merge branch 'ademeure-deterministic_layernorm' 2024-05-24 21:53:08 +00:00
Andrej Karpathy 25f17e6748 small formatting fix before merge 2024-05-24 21:53:00 +00:00
Andrej Karpathy dbacaf84cf Merge branch 'deterministic_layernorm' of https://github.com/ademeure/llm.c into ademeure-deterministic_layernorm 2024-05-24 21:46:25 +00:00
Andrej bf03e7fd4c Merge pull request #456 from ngc92/large-batch-fix
fix for large batch sizes
2024-05-24 12:00:41 -07:00
Andrej Karpathy 16b364d256 change -o to be not file but directory, and create a main.log inside it for the log. this is to prepare for writing of checkpoints, where we want to re-use this directory for those as well 2024-05-24 17:57:50 +00:00
Erik Schultheis 1b98637960 int -> int64_t 2024-05-24 20:11:34 +03:00
Erik Schultheis df2e0dadd2 int64_t 2024-05-24 19:59:53 +03:00
Andrej Karpathy 1d10e788af fix bug pretty sure this should just be a void* pointer 2024-05-24 16:51:26 +00:00
Andrej Karpathy 613a125216 add checkpoint function write to file 2024-05-24 16:49:20 +00:00
Erik Schultheis 2a736cb9e2 fix for large batch sizes 2024-05-24 19:45:41 +03:00
Andrej Karpathy 08fe3eef43 change default shard size for fineweb10B 2024-05-24 15:56:05 +00:00
Andrej b0f065fcf8 Merge pull request #455 from ngc92/fineweb
update fine-web preprocessing: faster, and much less RAM consumption
2024-05-24 08:54:55 -07:00
Erik Schultheis b66eb66ac7 update fine-web preprocessing: faster, and much less RAM consumption 2024-05-24 17:45:26 +03:00
rosslwheeler a241a00ad4 Merge branch 'karpathy:master' into dataloader_win_fixes 2024-05-23 23:12:15 -07:00
Andrej Karpathy dee4e42548 add option to not run hellaswag, interferes with a bunch of testing, e.g. if T is low 2024-05-24 03:47:07 +00:00
Andrej Karpathy 032e76c259 start putting llm.c and pytorch right next to each other, identical training runs with identical results and prints. almost 2024-05-24 01:02:21 +00:00
Andrej 4b88d2a976 Merge pull request #453 from otabuzzman/master
A Swift port of this project
2024-05-23 14:26:28 -07:00
Andrej Karpathy 64b6a14679 add learning rate decay schedule, now we have the full scheduler implemented 2024-05-23 21:19:00 +00:00
Andrej Karpathy 661975cc3c add learning rate warmup option 2024-05-23 20:57:04 +00:00
Andrej Karpathy 645869b6f7 add weight decay -c option and be more careful in our tests of correctness, our weight decay didn't match to pytorch. also modify the betas in AdamW to be consistent with those used in GPT-3 training 2024-05-23 20:45:59 +00:00
Andrej Karpathy 949d71a3d2 only rank 0 logs 2024-05-23 20:32:56 +00:00
Andrej Karpathy 3cb2812774 skip hellaswag eval on step 0 i think... not sure but ok for now 2024-05-23 19:27:04 +00:00
Jürgen Schuck 8af1818cb1 Merge branch 'karpathy:master' into master 2024-05-23 21:18:01 +02:00
otabuzzman 5f87b13f34 Update documentation with Swift port reference 2024-05-23 21:15:01 +02:00
Andrej 518805b7f4 Init from random weights, matching PyTorch
Merge pull request #451 from karpathy/feature/random_init
2024-05-23 12:11:05 -07:00
Andrej Karpathy 1f91bfc442 fix small bug on eval logging 2024-05-23 18:19:25 +00:00
Andrej Karpathy 86682af9a9 llm.c matches pytorch init from scratch exactly now 2024-05-23 17:50:33 +00:00
Andrej Karpathy 70a9c75348 use pytorch rand and fix dumb bug lol 2024-05-23 15:16:49 +00:00
Andrej Karpathy bc1ebc110a Merge branch 'feature/rand' of https://github.com/azret/llm.c into feature/random_init 2024-05-23 14:59:44 +00:00
Andrej Karpathy e6a7d1d3e9 allow the python script to also init from random and save those weights, so it's a good reference for our C implementation 2024-05-23 14:59:25 +00:00
Andrej Karpathy d09631807a first draft of random init, crashes with some cuBLAS error, debugging 2024-05-23 14:35:46 +00:00
Ross Wheeler d3f26951ef Add glob() for windows
Tested with tinyshakespeare and fineweb
2024-05-23 00:06:18 -07:00
Andrej Karpathy d3cf9e211c fix bug free only if initted 2024-05-23 00:49:32 +00:00
Andrej eda0c2f591 Merge pull request #447 from karpathy/feature/hellaswagc
HellaSwag eval in C
2024-05-22 15:50:12 -07:00
Andrej Karpathy da59861d38 make hellaswag optional eval yay 2024-05-22 22:42:37 +00:00
Andrej Karpathy 67239d9b8f fixed a bunch of small issues and we now utilize the full batch dimension, this is very close to ready to merge, might even be ready 2024-05-22 22:21:53 +00:00
Andrej Karpathy 051f3ca53c first draft, apparently this works. needs cleanups, and also we are not yet utilizing the full batch dimension. we actually have to load in multiple examples and fully utilize batch 2024-05-22 19:24:20 +00:00
Andrej 69f12214bf Merge pull request #444 from karpathy/feature/sharded_data
extend dataloader to be sharded
2024-05-22 06:21:48 -07:00
Andrej Karpathy 099d30f814 add a super small crappy glob for windows that only matches a single unique file. this will make CI happy but we can't train on sharded data on windows until this is improved 2024-05-22 13:08:18 +00:00
Andrej Karpathy 05be4f6d82 readme changes 2024-05-22 12:32:25 +00:00
Andrej Karpathy edb0df967a continued changes for sharded dataloader 2024-05-22 11:59:09 +00:00
ademeure 7cbeefc7f3 added new layernorm backward to /dev/cuda/ 2024-05-21 23:26:54 +01:00
ademeure 7d0891f6dd Fully deterministic layernorm (slight perf loss) 2024-05-21 22:37:09 +01:00
Andrej Karpathy 31310282e1 extend dataloader to be sharded 2024-05-21 21:33:24 +00:00
Andrej Karpathy 967420d1d1 fix print format warning for size_t vs int 2024-05-21 19:30:02 +00:00
Andrej Karpathy 587506d096 torch tensor can't handle uint16 so let's convert to int32, which is silly because we'll convert to .long right after but ok 2024-05-21 19:22:09 +00:00
Andrej 170c92b771 Merge pull request #440 from karpathy/feature/organize_data
refactor datasets
you'll have to re-generate the dataset .bin files, for example, `python dev/data/tinyshakespeare.py`.
see main README section "major changes log" for more details.
2024-05-21 11:32:38 -07:00
Andrej Karpathy 54ccbd300c docs on master-breaking changes around dataset file representation 2024-05-21 18:31:19 +00:00
Andrej Karpathy 7d58fd2abb adjust py file as well and make the errors better 2024-05-21 18:29:55 +00:00
Andrej Karpathy f7cb77f3d9 docs on master-breaking change around how we store data .bin files 2024-05-21 18:20:51 +00:00
Andrej Karpathy ccc240caab make comment more helpful 2024-05-21 18:01:35 +00:00
Andrej Karpathy d53608820a the write_datafile function accepts python list, which i think is faster but i didn't check 2024-05-21 17:55:56 +00:00
Andrej Karpathy 9bad49a99a also fix the python file. that should be it now, w.r.t. the new token format .bin files 2024-05-21 17:46:14 +00:00
Andrej Karpathy 666145e7b8 also fix the .c file and fp32 file 2024-05-21 17:30:56 +00:00
Andrej Karpathy 1defbd4a19 adjust the dataloader to load the new .bin data files, and both prod datasets to use it instead 2024-05-21 17:05:27 +00:00
ademeure a3801f01ef added algorithm header for std::sort on windows (not sure about compile time impact...) 2024-05-21 16:53:11 +01:00
Andrej Karpathy f7cbb30316 step 2 of dataloader refactor: separate out the dataloader to its own file, change its signature a little bit, and (notably) change from cudaMallocHost to a simple malloc, so that this file doesn't have to be cuda aware 2024-05-21 15:32:46 +00:00
ademeure b5e75dde8e Fully deterministic encoder backward kernels for train_gpt2.cu 2024-05-21 15:57:07 +01:00
Andrej Karpathy 4bbd01fe12 add fineweb, and add the first version of a new write_shard function that contains a header properly 2024-05-21 00:16:07 +00:00
Andrej Karpathy c51cd70cfb fix a slip in gitignore, i think i am getting tired today 2024-05-20 23:22:31 +00:00
Andrej Karpathy bced34d71e adjust gitignore 2024-05-20 23:15:44 +00:00
Andrej Karpathy 7d11b7996c i also forgot to include the readme file for the new dev/data dir 2024-05-20 23:13:42 +00:00
Andrej Karpathy 71774b3c3d oops forgot to include data_common.py 2024-05-20 23:10:49 +00:00
Andrej Karpathy f671cf9288 more changes, trying to help people out because when this merges to master it will brick everyone's code... 2024-05-20 23:02:41 +00:00
Andrej Karpathy 722e5b2fe5 refactor how we treat datasets, because we're about to have more of them and we don't want them to clutter up root dir etc. this is only step 1, i'm about to refactor a bunch of the dataloading, how the .bin files work and are loaded, how the DataLoader works, etc. This is all needed to support good evals and training at scale 2024-05-20 22:43:40 +00:00
lancer 2b0667aee1 update the utils function and assert 2024-05-20 08:00:39 -07:00
lancer 6348d4196d fix the unsupported block_size 2024-05-19 17:39:25 -07:00
Andrej 6c8bc172cc Merge pull request #436 from ChrisDryden/boundscheck
Moved bounds checks to outside of the kernel
2024-05-20 00:41:41 +01:00
Andrej 2751fa0880 Merge pull request #435 from ChrisDryden/warpsize
Added warpsize as a constant for better compile time optimization and standardization
2024-05-20 00:39:15 +01:00
Christopher 6de1137e0e Moved bounds checks outside of kernel into assertions 2024-05-19 23:17:45 +00:00
Christopher ead5d35973 Added warpsize as a constant for better compile time optimization and standardization 2024-05-19 22:54:47 +00:00
Andrej 1809740a1c Merge pull request #433 from junaire/jun/some_fix
Check if file exists using platform specific commands & Add *.o to gi…
2024-05-19 21:02:04 +01:00
Andrej 00c291c679 Merge pull request #434 from ngc92/profile-fix
initialize multi_gpu_config so profile doesn't crash
2024-05-19 20:57:11 +01:00
Erik Schultheis 299ce659ec initialize multi_gpu_config so profile doesn't crash 2024-05-19 22:25:00 +03:00
Andrej Karpathy 5a26799efb Merge branch 'ngc92-gradient-clipping' 2024-05-19 17:51:44 +00:00
Andrej Karpathy bc58cd1dc1 fix small bugs in grad clip, introduce a GPU CPU synch point to communicate the float grad_clip to print it, and small printing changes 2024-05-19 17:51:29 +00:00
Andrej Karpathy c2d12f725e small touchups to grad clip 2024-05-19 17:07:55 +00:00
Andrej Karpathy da9f4e21b0 Merge branch 'gradient-clipping' of https://github.com/ngc92/llm.c into ngc92-gradient-clipping 2024-05-19 16:03:53 +00:00
Jun Zhang 9e645314e6 Check if file exists using platform specific commands & Add *.o to gitignore
Signed-off-by: Jun Zhang <jun@junz.org>
2024-05-20 00:00:21 +08:00
Andrej 4a899f5685 Merge pull request #431 from karpathy/feature/hellaswag
add hellaswag and mmlu PyTorch reference eval
2024-05-19 16:53:38 +01:00
Andrej Karpathy 8d55c4a6ac fix a bug, we have to be careful to make sure we evaluate loss at the token just before the first completion token, because that is the prediction for the first token, and its accuracy 2024-05-19 15:45:28 +00:00
Andrej Karpathy 4e14b52285 report both acc and acc_norm 2024-05-19 15:18:04 +00:00
Andrej Karpathy 7bd2389c1f add mmlu as well and refine both a bit 2024-05-19 14:23:51 +00:00
Andrej Karpathy aec9ce5d0a move hellaswag file to dev 2024-05-19 12:18:02 +00:00
Andrej Karpathy 77b991281f add hellaswag reference eval. it looks pretty bad honestly, the examples themselves, but i'm told it is predictive... hmm 2024-05-19 11:51:01 +00:00
Erik Schultheis 66ce5766e0 fixed up dev/cuda 2024-05-18 23:06:44 +03:00
ngc92 589ead1e4f updated code to adapt to latest changes 2024-05-18 22:47:42 +03:00
Erik Schultheis c3a3b9daa5 fixed profile target 2024-05-18 22:47:40 +03:00
Erik Schultheis a9947a8315 added a flag and the missing sqrt;
testing now has clipping enabled
2024-05-18 22:47:15 +03:00
Erik Schultheis d7a81ef26f added a useful mixed precision utility for dev/cuda 2024-05-18 22:45:44 +03:00
Erik Schultheis 44d45bdd6a first draft for gradient clipping by global norm 2024-05-18 22:45:21 +03:00
Andrej Karpathy 3de05baaf4 Merge branch 'ngc92-errchk' 2024-05-18 18:35:28 +00:00
Andrej Karpathy 4374360015 adjust wte upper bound a bit, and print always because this part is really tricky and i don't trust anything other than manual inspection, even if we pass, allegedly 2024-05-18 18:34:24 +00:00
Erik Schultheis bf36a4b6b5 improved numerical error checking:
tighter tolarances
  relative tolerance based of bf16 epsilon
  less verbose output if all is OK
2024-05-18 13:27:35 +03:00
Andrej Karpathy f226a2c9d9 Merge branch 'ngc92-alloc' 2024-05-17 15:13:48 +00:00
Andrej Karpathy c8fa7a8c63 revert the adamw allocation to previous. minor 2024-05-17 15:13:36 +00:00
Erik Schultheis b24279c4db remove duplicate workspace allocation 2024-05-17 17:37:47 +03:00
Erik Schultheis 88c3bea890 print message before actual allocation for more informative OOM behaviour 2024-05-17 17:37:45 +03:00
Erik Schultheis 0f23723ae4 joined optimizer state allocation 2024-05-17 17:37:16 +03:00
Andrej b554d4d093 Merge pull request #423 from anthonix/readme_link
Add link to AMD fork
2024-05-16 21:35:18 +01:00
Anthony Blake 3113656e3f Add link to AMD fork 2024-05-16 13:33:34 -07:00
Andrej 109f516367 Merge pull request #383 from KarhouTam/feature/online-softmax-forward-without-cgs
Free cooperative groups implementation of online softmax forward
2024-05-16 21:23:56 +01:00
Andrej 9e04deabe3 Merge pull request #416 from azret/feature/cpu_dev
Setting up dev/cpu area with the first matmul_forward.c
2024-05-16 21:12:55 +01:00
Andrej Karpathy 130712215d Merge branch 'backward-bias' of https://github.com/ngc92/llm.c 2024-05-16 20:00:56 +00:00
Andrej Karpathy 68eea0bd5f Fix ZeRO-1, we were communicating way too much 2024-05-16 19:48:24 +00:00
Andrej Karpathy 6cfc7c5ce3 Merge branch 'zero-update' of https://github.com/ngc92/llm.c into ngc92-zero-update 2024-05-16 19:30:40 +00:00
Andrej bd7dc7af96 Merge pull request #422 from karpathy/feature/recompute
feature/recompute
2024-05-16 20:25:31 +01:00
Andrej Karpathy d7581fc542 make recompute be an int instead of bool, so we can strengthen it over time just like ZeRO stages, as we recompute more and more of the model in the future possibly. and make it default on because it is awesome 2024-05-16 19:07:39 +00:00
Erik Schultheis fbd8f03eea fixup profiling 2024-05-16 14:25:06 +03:00
Erik Schultheis 8b57cf6535 reduce communication overhead for ZERO stage 1 2024-05-16 14:12:29 +03:00
Erik Schultheis 57f70ea66b simplify multi-gpu logic by reducing #ifdefs 2024-05-16 14:12:14 +03:00
Andrej Karpathy 3277ccc35d Merge branch 'act-recomp' of https://github.com/ngc92/llm.c into ngc92-act-recomp 2024-05-16 10:26:23 +00:00
Andrej fbc5cba1f9 Merge pull request #415 from azret/feature/eot
train_gpt.c: Removing the hardcoded GPT2_EOT
2024-05-16 10:45:22 +01:00
Erik Schultheis d48c3a4945 (optionally) recompute gelu activations to reduce activation memory 2024-05-16 12:39:28 +03:00
Erik Schultheis 7b810c1a3b update main training script 2024-05-16 01:04:58 +03:00
Erik Schultheis 858c6e6dae deterministic kernel 2024-05-16 01:04:10 +03:00
Erik Schultheis 2ccdfb70e0 general cleanup 2024-05-16 01:04:10 +03:00
Andrej 0f48a28f12 Merge pull request #412 from karpathy/feature/grad_accum3
[wip] gradient accumulation, another attempt
2024-05-15 22:53:16 +01:00
Andrej Karpathy 8eb3a432d9 revert all streams and synchronization events, we'll bring them back but only one at a time and very very carefully, because they cause really subtle bugs and issues 2024-05-15 21:09:46 +00:00
Azret Botash 160b3bd007 Setting up dev/CPU area with the first matmul_forward.c 2024-05-14 18:22:39 -07:00
Azret Botash 16f9dad301 Update train_gpt2.c 2024-05-14 17:45:08 -07:00
Azret Botash 7de50af511 train_gpt.c: Removing the hardcoded GPT2_EOT 2024-05-14 17:38:09 -07:00
Azret Botash 3c3c965840 Adding Mersenne Twisters C 2024-05-14 17:16:24 -07:00
Andrej Karpathy a4567ae940 delete parallels, still not fixed 2024-05-14 22:13:54 +00:00
Andrej Karpathy 92fc26eba4 the nuts and bolts of gradient accumulation again, merged to master, but there is a bug and it doesn't work, debugging... 2024-05-14 21:32:47 +00:00
Erik Schultheis b2a5508b84 constness fixes 2024-05-14 22:48:59 +03:00
Erik Schultheis 3b5933ecfb considerably speed up CPU matmul while still keeping it relatively readable 2024-05-14 22:32:24 +03:00
Andrej Karpathy 2d43e5bc97 remove legacy comment 2024-05-14 19:14:07 +00:00
Andrej 222d59fa2f Merge pull request #408 from ngc92/layernorm-bw-dev-cuda
Layernorm backward updates
2024-05-14 20:13:09 +01:00
Erik Schultheis e553e2f084 update dev/cuda/layernorm_backward and improve validate_result to take into account fp epsilon when comparing results 2024-05-14 20:44:50 +03:00
Erik Schultheis dd8c9f5ec9 fix layernorm backward: accumulate weight gradient 2024-05-14 20:43:51 +03:00
Andrej Karpathy 2346cdac93 Merge branch 'ngc92-backward-bias' 2024-05-13 20:19:09 +00:00
Andrej Karpathy 37a96debbf Merge branch 'backward-bias' of https://github.com/ngc92/llm.c into ngc92-backward-bias 2024-05-13 20:18:33 +00:00
Andrej 750c5fd7e8 Merge pull request #309 from ahrefs/zero-stage1
Zero Redundancy Optimizer - Stage1
2024-05-13 21:08:53 +01:00
Erik Schultheis c66e48c06c fixup comment 2024-05-13 20:58:07 +03:00
Erik Schultheis 65727d5a4d fix CI compile by disabling kernel 5 2024-05-13 19:20:03 +03:00
Erik Schultheis 49ee3c8307 fix non-atomic version:
* accumulate instead of assign
* need dedicated argument to correctly handle the floatX == float case
2024-05-13 18:27:56 +03:00
Erik Schultheis aa41b32624 update main file 2024-05-13 17:48:25 +03:00
Erik Schultheis 081d224b21 automatically switch to buffer-less version if that can fill up the GPU 2024-05-13 17:39:32 +03:00
Erik Schultheis c0329ebdba new kernel version with fewer atomics 2024-05-13 17:18:27 +03:00
chinthysl f613ce895b Fix copy and cast params to master weights 2024-05-13 09:13:02 +00:00
chinthysl 086ce2ff4d Remove unsused template specializations and refactor 2024-05-13 06:48:22 +00:00
chinthysl b3a5d1da15 shard master_weights 2024-05-13 05:55:46 +00:00
Chinthaka Gamanayake 0d9ce74b74 Merge branch 'karpathy:master' into zero-stage1 2024-05-13 13:34:52 +08:00
Erik Schultheis 2287da0120 enable bf16 2024-05-12 19:42:32 +03:00
Andrej c1814d5d2f Merge pull request #401 from rosslwheeler/ci-add-win
Add Windows to CI
2024-05-12 18:32:13 +02:00
Andrej a2bffb5a8a Merge pull request #403 from nietras/nietras-llm-cs
Add Llm.cs by nietras to README
2024-05-12 18:29:20 +02:00
Erik Schultheis ec92368f8b added current backward bias kernel to dev cuda 2024-05-12 19:24:43 +03:00
ntr e64df91191 Add Llm.cs by nietras to README 2024-05-12 16:02:36 +02:00
Ross Wheeler b88f683569 Add Windows to CI
Change notes:
•	Add Windows build/test to CI matrix build
•	Add Cuda Windows build
•	Replace v3 checkout actions to v4 per GitHub recommendations
2024-05-11 17:49:22 -07:00
Andrej Karpathy 1cfe899121 Merge branch 'master' of github.com:karpathy/llm.c 2024-05-11 17:21:18 +00:00
Andrej Karpathy fb86c121a1 Merge branch 'Anerudhan-feature/simplify_cudnn_usage' 2024-05-11 17:21:04 +00:00
Andrej Karpathy 3dbb0bb89c bump the threshold for qkvw because flashattention expands the error here a tiny bit 2024-05-11 17:19:41 +00:00
Andrej Karpathy 58d9dda497 Merge branch 'feature/simplify_cudnn_usage' of https://github.com/Anerudhan/llm.c into Anerudhan-feature/simplify_cudnn_usage 2024-05-11 17:08:31 +00:00
Andrej 5dbda7e115 Merge pull request #397 from lancerts/comment-fix
Update the -lcublas -lcublasLt flag in the comment
2024-05-11 19:03:40 +02:00
Andrej 923b2b0a8f Merge pull request #398 from zwets/fix-cudnn-frontend-check
Fix detection of cudnn-frontend in '.' on Linux
2024-05-10 12:13:53 +02:00
Marco van Zwetselaar 6da5e63e2c Fix detection of cudnn-frontend in '.' on Linux
Plus minor legibility fixes in the cudnn-frontend explanation.
2024-05-10 10:57:39 +03:00
lancer 5c90845f7a update the -lcublas -lcublasLt flag in the comment 2024-05-09 22:04:02 -07:00
Andrej fd007d6508 Merge pull request #395 from yijunyu/patch-1
Update README.md
2024-05-10 01:18:25 +02:00
Andrej Karpathy 83ad9ac87f Fused Residual+LayerNorm layers, ~1% faster 2024-05-09 23:16:41 +00:00
Andrej Karpathy 43c48707a7 Merge branch 'fuse-ln-res-2' of https://github.com/ngc92/llm.c into ngc92-fuse-ln-res-2 2024-05-09 23:09:53 +00:00
Yijun Yu 8ccf2f97f8 Update README.md 2024-05-10 06:49:24 +08:00
ngc92 75ec629f52 remove duplicate function 2024-05-10 00:03:30 +02:00
ngc92 5920143d47 remove warning noise 2024-05-10 00:02:37 +02:00
ngc92 691c1df969 fused layernorm+residual 2024-05-09 23:48:11 +02:00
Anerudhan 01c7a331bb - Simplify graph cache and usage of cudnn.
- Fix failures in H100
2024-05-09 19:07:00 +00:00
Andrej 077b649ee1 Merge pull request #389 from vedaanta/cudnn/install
Adds cuDNN install instructions for a smaller but sufficient package
2024-05-09 20:30:48 +02:00
ngc92 99e765de8d moved bf16 boilerplate to common.h 2024-05-09 11:19:24 +02:00
Joe Halabi 69aa64cc80 Adds cuDNN install instructions for a smaller but sufficient package
- Modifies README.md to provide example apt-get cuDNN install instructions that install the cuDNN dev package. "sudo apt-get install -y cudnn" will install the default cuDNN packages, but for a minimal setup, installing the dev package will see a 50% reduction in both, download size (~850MB to 425MB now) and local storage size (~2GB to ~1GB now).

- Modifies the Makefile to point users to the README for cuDNN install instructions (through comments and the cuDNN install error message)

- Modifies attention_forward.cu comments to point users to the README for cuDNN install instructions

Signed-off-by: Vedaanta Agarwalla <vagarwalla@ipp2-1949.nvidia.com>
2024-05-08 14:53:31 -07:00
Andrej b8eaafdf9b Merge pull request #386 from Anerudhan/feature/rename_cudnn_attn
Rename cudnn_att.cu to cudnn_att.cpp to speed up compilation.
2024-05-08 23:00:57 +02:00
Andrej Karpathy 68509c89a2 Merge branch 'pjj-fix-gpt2-xl' 2024-05-08 20:20:27 +00:00
Andrej Karpathy 0c73ba6ca5 Merge branch 'fix-gpt2-xl' of https://github.com/pjj/llm.c into pjj-fix-gpt2-xl 2024-05-08 20:19:15 +00:00
Andrej Karpathy 26dbbc75c2 unbreak the tests oops 2024-05-08 20:18:33 +00:00
Andrej Karpathy 84a2aa0dd0 allow exporting all GPT-2 model sizes now, using --model flag in python script. also allow loading all the models from C, using the new flag -e, to point directly to the file to load. added some error handling for potentially common mistakes. note that loadin GPT-2 XL does not work yet and crashes, but I am imminently merging a fix as the next commit that resolves this 2024-05-08 20:13:14 +00:00
Anerudhan 0ce5fcf7e8 Rename cudnn_att.cu to cudnn_att.cpp to speed up compilation. 2024-05-08 19:53:48 +00:00
Paul Maragakis 99e5c5d283 Allow the code to train gpt2-xl
This commit is the fix by adameure described in the comment below:

https://github.com/karpathy/llm.c/pull/382#issuecomment-2100895501

To reproduce the following bug:

https://github.com/karpathy/llm.c/pull/382#issuecomment-2100648148

First switch to gpt2-xl for tinystories with the following three commands:

python prepro_tinystories.py  # download tinystories

sed -i 's/from_pretrained("gpt2")/from_pretrained("gpt2-xl")/' train_gpt2.py  # select gpt2-xl

python train_gpt2.py --input_bin data/TinyStories_train.bin --batch_size 1  # Prep the large model binaries

Now you can compile and try the code with or without the fix.  On an H100 machine I use the following command:

train_gpt2cu -i data/TinyStories  -b 10 -x 100  # gets 230k/s on one H100 node of 8 GPU
2024-05-08 15:17:25 -04:00
Andrej 1c5ef8cea3 Merge pull request #382 from ahrefs/fix-multigpu-hang
Fix incorrect GPU assignment in multi gpu setup
2024-05-08 19:50:00 +02:00
KarhouTam 6d7a99cb4e Adjust comments 2024-05-08 17:45:14 +08:00
KarhouTam 6a52d8619e Optimize codes and comments 2024-05-08 16:12:39 +08:00
chinthysl 2356be7333 set stream to main_stream in ncclAllReduce 2024-05-08 05:02:23 +00:00
chinthysl dc901d420b set correct gpu using multigpu config 2024-05-08 04:50:24 +00:00
KarhouTam b3e8a9fe67 Implementation of online softmax forward kernel without cgs. 2024-05-08 09:22:14 +08:00
Andrej Karpathy 2f6c545131 Merge branch 'ademeure-more_kernel_opt' 2024-05-07 09:36:37 +00:00
Andrej Karpathy 5b070903ce Merge branch 'more_kernel_opt' of https://github.com/ademeure/llm.c into ademeure-more_kernel_opt 2024-05-07 09:31:12 +00:00
ademeure c261240852 3 x 512 threads max for layernorm_backward to avoid cache thrashing (hacky -> better way?) 2024-05-07 03:43:27 +01:00
Andrej 0141408ca5 Merge pull request #357 from rosslwheeler/cudnn_makefile_changes
cuDNN makefile changes
2024-05-07 02:50:11 +02:00
Andrej ce72242c46 Merge pull request #376 from rosslwheeler/train_gp2cu-porting-fixes
Replace non-standard C syntax with standard C syntax - fix win compilation
2024-05-07 02:47:47 +02:00
Ross Wheeler 3a6f89cebc Merge branch 'train_gp2cu-porting-fixes' of https://github.com/rosslwheeler/llm.c into train_gp2cu-porting-fixes 2024-05-06 17:41:47 -07:00
Ross Wheeler 2550754215 Change FLT_MAX to flt_max 2024-05-06 17:41:19 -07:00
Andrej Karpathy 3cc16f1354 fix logits bug 2024-05-07 00:39:25 +00:00
rosslwheeler 60b819630d Merge branch 'karpathy:master' into train_gp2cu-porting-fixes 2024-05-06 17:25:59 -07:00
ademeure e1f89b3043 remove BF16 default from classified_fused before PR 2024-05-07 00:57:01 +01:00
ademeure 1ea7f9bf25 tiny irrelevant optimisation to final unaligned fused_classifier loop + add missing common.h changes 2024-05-07 00:54:50 +01:00
Andrej ec9fd51d5c Merge pull request #375 from Chillee/patch-1
Don't return logits during training for PyTorch baseline
2024-05-07 01:49:02 +02:00
Andrej aa54a0a808 Merge pull request #377 from ademeure/profiler_efficiency
Add DRAM BW, Tensor Core Utilisation, & Efficiency Stats
2024-05-07 01:45:33 +02:00
ademeure 9b55ea8cf5 More crazy optimisations to layernorm_backward, fused_classifier, and matmul_backward_bias. 2024-05-07 00:11:42 +01:00
Ross Wheeler f7d7760084 Non-standard C syntax replace with standard C
Also, adding unistd.h back because that has some changes required to find M_PI.
2024-05-06 13:56:17 -07:00
ademeure 5adb6ef2c1 Add tensor core and overall efficiency stats to profiler script. 2024-05-06 21:24:42 +01:00
Horace He 69f6c4f765 Don't return logits during training for PyTorch baseline
This improves perf somewhat, since currently it's always returning logits (which thus need to be materialized).
2024-05-06 13:10:17 -07:00
Ross Wheeler 134f4c7fc8 Adding two directory search for cuDNN frontend files
Search in $(HOME) and in the current directory.
2024-05-06 12:38:40 -07:00
rosslwheeler dedd86c1b3 Merge branch 'karpathy:master' into cudnn_makefile_changes 2024-05-06 11:20:22 -07:00
Andrej 4274d95b4a Merge pull request #365 from lancerts/classifier-fix
Minor update on the classifier_fused.cu
2024-05-06 12:04:16 +02:00
Andrej Karpathy 6798dd4b91 Merge branch 'rosslwheeler-Makefile-detect-gpu' 2024-05-06 10:01:57 +00:00
chinthysl 7a8f471bc7 further reorganization 2024-05-06 17:49:58 +08:00
Ross Wheeler 5344cc30b8 Merge branch 'Makefile-detect-gpu' of https://github.com/rosslwheeler/llm.c into Makefile-detect-gpu 2024-05-06 01:11:37 -07:00
Ross Wheeler 0c4908d815 Adding CI check to disable auto-detect 2024-05-06 01:11:27 -07:00
rosslwheeler 98aa83b282 Merge branch 'karpathy:master' into Makefile-detect-gpu 2024-05-06 00:58:09 -07:00
Ross Wheeler 2bfd2b6a7e Auto-detect GPU capability
3 cases tested on Windows 11 and Ubuntu 22.04
2024-05-06 00:46:28 -07:00
Chinthaka aae7145495 Merge branch 'master' into zero-stage1 2024-05-06 10:36:24 +08:00
lancer bbfe8c989c Minor update on the code 2024-05-05 16:59:58 -07:00
Andrej 9e4f9df9a2 Merge pull request #337 from lancerts/make-fix
Makefile type fix: move gelu_backward to backward block
2024-05-05 15:37:02 -07:00
Andrej a7c1e4437d Merge pull request #364 from ngc92/cleanup_may4-nvcc11
make things compile with nvcc11
2024-05-05 14:42:49 -07:00
Andrej Karpathy d6e8f5e026 Merge branch 'ademeure-cleanup_may4' 2024-05-05 21:29:41 +00:00
Erik Schultheis 804a9af1dc make things compile with nvcc11 2024-05-05 23:43:02 +03:00
ademeure 876ab93c0a Add FP16 path for atomicStochasticAdd (+remove __bfloat1622float2 to work on older CUDAs) + fixes 2024-05-05 15:05:59 +01:00
ademeure 9910a40863 Removed makefile change so we can integrate #339 instead which feels slightly cleaner 2024-05-05 14:42:32 +01:00
Ross Wheeler 42176ad4d5 Merge branch 'cudnn_makefile_changes' of https://github.com/rosslwheeler/llm.c into cudnn_makefile_changes 2024-05-04 23:27:21 -07:00
Ross Wheeler 2d4e5fd840 Minor fixes for Makefile for cudnn and windows support
Tested on Ubuntu Linux 22.04 and Windows 11
2024-05-04 23:26:57 -07:00
Ross Wheeler 9663719c62 Minor fixes for Make for cudnn and windows support
Tested on Ubuntu Linux 22.04 and Windows 11
2024-05-04 23:24:41 -07:00
ademeure c15ca1f4cf Fix makefile 2024-05-05 02:49:16 +01:00
ademeure 7789738879 Add PTX back to binary + fix whitespaces 2024-05-05 02:44:24 +01:00
ademeure 8675104b73 Compile for the user's GPU architecture using nvidia-smi query on Linux 2024-05-05 02:24:31 +01:00
ademeure 83ec4b8e70 Slightly reduce lines of code in cudnn_att 2024-05-05 02:09:52 +01:00
ademeure c261eecda8 Merge remote-tracking branch 'karpathy/master' into cleanup_may4 2024-05-05 01:49:25 +01:00
Andrej Karpathy 64b6c2a77f Merge branch 'ngc92-separate-compilation' 2024-05-05 00:41:18 +00:00
Andrej Karpathy f2224f2eb2 Merge branch 'separate-compilation' of https://github.com/ngc92/llm.c into ngc92-separate-compilation 2024-05-05 00:32:31 +00:00
Andrej Karpathy 8168b78549 Merge branch 'ademeure-less_idle_more_brrr-3' 2024-05-05 00:27:15 +00:00
Andrej Karpathy ce333de156 add ema to tok/s 2024-05-05 00:18:39 +00:00
Andrej Karpathy 2642ffd19c Merge branch 'less_idle_more_brrr' of https://github.com/ademeure/llm.c into ademeure-less_idle_more_brrr-3 2024-05-05 00:05:38 +00:00
ademeure a0b79358bd Fix loss not being calculated for forward-only paths 2024-05-05 01:01:59 +01:00
Andrej 6c179faaeb Merge pull request #349 from lancerts/fix-include
Remove the duplicated #include <assert.h>
2024-05-04 17:01:32 -07:00
Erik Schultheis b087b9c819 don't compile/link cudnn if not asked for it 2024-05-05 02:54:42 +03:00
ademeure 34304361ba fix merge conflict 2024-05-05 00:16:49 +01:00
ademeure ec0ab2d2d0 Remove arch=native as it only available on recent CUDA versions 2024-05-05 00:00:58 +01:00
ademeure 18d7ed9242 revert profile.py changes for now 2024-05-04 23:46:39 +01:00
ademeure abaaceb801 Added makefile gencode changes 2024-05-04 23:43:13 +01:00
ademeure bfb9c51446 refactoring & remove unused functions to reduce LOC (+wip profile.py improvements) 2024-05-04 23:40:15 +01:00
ademeure 66b9755ce9 Use template for blockReduce's warp reduction + add out_of_bounds parameter 2024-05-04 18:48:54 +01:00
ademeure 8d593b916e fix cuda stream bug for test_gpt2cu... by removing most of the memcpy/memset parallelism for now (keeping loss_event to avoid sync between forward and backward). 2024-05-04 18:18:51 +01:00
Erik Schultheis 19c290d7e6 improved debugging for cudnn 2024-05-04 12:31:15 +03:00
Erik Schultheis aa5bb258b6 fixed up test and profile targets 2024-05-04 11:59:59 +03:00
Chinthaka 58ac83f5ec Merge branch 'master' into zero-stage1 2024-05-04 16:24:22 +08:00
chinthysl 59b66f6c07 refactor and fix CI issue 2024-05-04 16:22:44 +08:00
lancer 03b7323d1e minor fix of the include 2024-05-03 17:38:39 -07:00
Andrej ff2fbdc2a1 Merge pull request #332 from harryjackson/llm.java
Adding a java port of the project
2024-05-03 13:58:59 -07:00
Erik Schultheis 67a82a611e first attempt at moving cudnn out of the main file for faster compiles 2024-05-03 23:08:44 +03:00
Andrej aa5dd2eeab Update README.md 2024-05-03 12:43:08 -07:00
Andrej 6260a1844b Merge pull request #345 from ngc92/profile
updated the profile script to be more robust, and adapted to recent k…
2024-05-03 12:17:22 -07:00
Erik Schultheis 977d6a6c6f fix × symbol 2024-05-03 21:57:31 +03:00
Erik Schultheis 7fabb12e2e fixed instruction counts 2024-05-03 21:54:06 +03:00
Erik Schultheis 13eeeb20c8 try to detect if we need sudo 2024-05-03 21:31:54 +03:00
Erik Schultheis 325b45647f updated the profile script to be more robust, and adapted to recent kernel changes
make it compile with cudnn
default batchsize is 24 now
2024-05-03 21:20:06 +03:00
ademeure a0b80920f1 Reducing profiler to only 3 steps instead of 10 given slight perf penalty (todo: make it configurable) 2024-05-03 19:03:06 +01:00
ademeure a19fcecb79 Set cuBLAS(Lt) to use the new main_stream 2024-05-03 18:35:46 +01:00
Andrej Karpathy 6c629ac571 Merge branch 'master' of github.com:karpathy/llm.c 2024-05-03 15:54:36 +00:00
Andrej Karpathy e780b56974 small fixes for profiling 2024-05-03 15:54:29 +00:00
ademeure b58181079a super optimised matmul_bias + cuda streams for better parallelism + fix fused_classifier cache hint + remove remaining cooperative groups 2024-05-03 16:38:46 +01:00
Chinthaka 0fd384bb21 Merge branch 'master' into zero-stage1 2024-05-03 11:08:41 +08:00
ademeure f456a75e3f add CUDA streams (no parallelism yet) + fix profiling init + new "blockReduce" function 2024-05-03 03:14:16 +01:00
lancer c81adeb6b2 move gelu_backward to backward block 2024-05-02 17:36:12 -07:00
Andrej c04c133b25 Merge pull request #335 from karpathy/feature/bias_backward2
v1 of the new matmul backward bias kernel
2024-05-02 16:52:34 -07:00
Andrej Karpathy 776541dac7 v1 of the new matmul backward bias kernel 2024-05-02 23:32:20 +00:00
Andrej Karpathy 6ebef46f83 ugh didn't notice this tiny rebasing mistake, introduced a bug. good candidate for a CI that we can overfit a single batch in the train_gpt2.cu script and get the exact same numbers as we expect in the test_gpt2.cu file 2024-05-02 23:23:28 +00:00
Andrej Karpathy 2c7960040d Merge branch 'ngc92-bias-backward' 2024-05-02 22:17:31 +00:00
Andrej Karpathy 937fbd8e36 small doc fixes 2024-05-02 22:17:22 +00:00
Erik Schultheis 6e48501d60 bias backward kernel that will use all available threads 2024-05-03 01:05:47 +03:00
Andrej Karpathy 6a3d63fb1d Merge branch 'PeterZhizhin-add_nvtx_ranges_and_timings' 2024-05-02 21:13:26 +00:00
Andrej Karpathy 79505bc6b3 resolve merge and small fixes 2024-05-02 21:13:08 +00:00
Andrej Karpathy 50714d2254 pragma unroll quick fix 2024-05-02 21:06:11 +00:00
Andrej Karpathy 9da2729c5a update encoder_forward with the latest and greatest Packed128 and incorporate into train_gpt2cu 2024-05-02 20:40:05 +00:00
Peter Zhizhin 41a0789a44 Added NVTX ranges for FlashAttention 2024-05-02 20:30:08 +00:00
Andrej Karpathy 8a9510f32b very minor stylistic preference let's keep the code style consistent 2024-05-02 19:48:56 +00:00
Andrej Karpathy 84dbc70160 Merge branch 'JaneIllario-patch-1' 2024-05-02 19:45:18 +00:00
Andrej Karpathy 3d1761b370 Merge branch 'patch-1' of https://github.com/JaneIllario/llm.c into JaneIllario-patch-1 2024-05-02 19:43:14 +00:00
Andrej Karpathy 01df41fbba careful to guard the overloaded warpReduceSum and small fixes 2024-05-02 19:24:04 +00:00
Andrej Karpathy 6879c36cf3 Merge branch 'JaneIllario-gelu_backwards-3' 2024-05-02 19:16:38 +00:00
Andrej Karpathy 99f51bae1c Merge branch 'gelu_backwards' of https://github.com/JaneIllario/llm.c into JaneIllario-gelu_backwards-3 2024-05-02 19:15:06 +00:00
Andrej Karpathy 8ccd05ec29 group warp reduce ops together 2024-05-02 19:09:53 +00:00
Andrej Karpathy d05c51c7d1 Merge branch 'ChrisDryden-removing_cg_example' 2024-05-02 19:08:26 +00:00
Andrej 2583feb7a2 Merge pull request #327 from yushengsu-thu/master
[ModuleNotFoundError]: Solution - add required package (requests) into requirements.txt
2024-05-02 12:03:46 -07:00
Andrej Karpathy d1771a7b59 quick fix master. there is some weirdness here in the adam update that is todo to understand better, for now i just want the master to be ok 2024-05-02 19:01:26 +00:00
Harry Jackson c13a730149 Merge branch 'master' into llm.java 2024-05-02 11:44:44 -07:00
Andrej Karpathy ec73086a83 Merge branch 'removing_cg_example' of https://github.com/ChrisDryden/llm.c into ChrisDryden-removing_cg_example 2024-05-02 18:14:14 +00:00
Andrej Karpathy 2202c9a51b add kernel 4 to docs. have to improve these docs more and document them better 2024-05-02 18:09:27 +00:00
Andrej Karpathy 2feb9ffcfb Fused classifier: use Packed128 and delete cg dep 2024-05-02 18:08:27 +00:00
Peter Zhizhin bcf7d4fe93 Add NSight Compute ranges, use CUDA events for timings 2024-05-02 17:23:07 +00:00
Yusheng (Ethan) Su 31e59273e1 Merge branch 'karpathy:master' into master 2024-05-02 14:04:47 +04:00
chinthysl 632caf1ce6 Refactored zero_active var 2024-05-02 07:38:49 +00:00
chinthysl d35d2e606b Generalized copy_and_cast_kernel and changes to cater model->master_weights 2024-05-02 07:31:34 +00:00
Chinthaka 07b8102e79 Merge branch 'master' into zero-stage1 2024-05-02 11:48:56 +08:00
“Christopher 398cdaf4cd Removed cooperative groups and added packed128 to fused classifier and prepare softmax blockwide 2024-05-02 02:33:51 +00:00
“Christopher 665d0a4f89 Added example of removing cooperative groups 2024-05-02 02:09:57 +00:00
Jane Illarionova d6d0d50627 fix typo 2024-05-02 00:18:55 +00:00
Jane Illarionova d0e5fea2dd update residual_forward to use util functions 2024-05-02 00:17:47 +00:00
Jane Illarionova 7203c7875f remove int cast 2024-05-02 00:17:47 +00:00
Jane Illarionova 4698164e3e update train to multiply block size 2024-05-02 00:17:47 +00:00
Jane Illarionova 52b4dd02a1 update residual forward to allow kernels to use both types 2024-05-02 00:17:47 +00:00
Jane Illarionova 32df703a05 Update train_gpt2.cu 2024-05-02 00:17:47 +00:00
Jane Illarionova 4ffcf5b90a Update residual_forward.cu 2024-05-02 00:17:47 +00:00
Jane Illarionova 3de3c53ba9 update kernel with util functions 2024-05-01 23:58:49 +00:00
Jane Illarionova d98e5ae859 remove int casting 2024-05-01 23:36:59 +00:00
Jane Illarionova e9a80b5d84 update ceildiv for gelu_backward 2024-05-01 23:36:59 +00:00
Jane Illarionova 7746217433 update gelu backward rto allow all kernels to use both types 2024-05-01 23:36:59 +00:00
Jane Illarionova ab2de05a13 Update train_gpt2.cu 2024-05-01 23:36:59 +00:00
Jane Illarionova 80072802d7 Create gelu_backward.cu 2024-05-01 23:36:59 +00:00
Andrej 4dd1ab43de Merge pull request #328 from karpathy/feature/fp32_weight_master_copy
feature/fp32 weight master copy
2024-05-01 16:07:19 -07:00
Andrej Karpathy c177c2694c make nice and print 2024-05-01 23:06:23 +00:00
Andrej Karpathy 795f8b690c fixes to keep master copy in fp32 of weights optionally 2024-05-01 22:50:22 +00:00
Yusheng (Ethan) Su 24abcbf6d2 Merge branch 'karpathy:master' into master 2024-05-02 02:19:16 +04:00
Yusheng Su ac2d635007 add required package (requests) into requirements.txt 2024-05-02 02:11:38 +04:00
Andrej Karpathy fd474feec8 Merge branch 'fp32-weights' of https://github.com/ngc92/llm.c 2024-05-01 22:02:55 +00:00
Erik Schultheis df9f0efde6 updated letter and properly free buffer 2024-05-02 01:00:36 +03:00
Andrej f4f7a987e7 Merge pull request #324 from eymay/init_weights_from_scratch
Refactoring parameter size filling
2024-05-01 14:50:39 -07:00
Erik Schultheis e9c76a1753 option to keep weights as fp32 2024-05-02 00:11:07 +03:00
Andrej ab95a11190 Merge pull request #325 from ngc92/dev-cuda-utils
mixed precision utilities for dev/cuda
2024-05-01 12:54:54 -07:00
Erik Schultheis b42db706b2 this kernel will be released later, sorry for the spoiler :) 2024-05-01 22:47:24 +03:00
Erik Schultheis 91bc72d0c0 fixup 2024-05-01 22:40:26 +03:00
Erik Schultheis 06980b6b2f mixed precision utilities for dev/cuda 2024-05-01 22:35:55 +03:00
Andrej d37639a8cb Merge pull request #316 from rosslwheeler/designated_init_fix
uint and designated initializers non-standard for C++ / Cuda
2024-05-01 12:13:12 -07:00
Andrej be5cee21e6 Merge pull request #317 from ChrisDryden/patch-5
Removed detail about layernorm inputs being f32, no longer the case
2024-05-01 12:00:01 -07:00
Andrej 38b478d8f9 Merge pull request #323 from karpathy/feature/cudnn
feature/cudnn for flash-attention
2024-05-01 11:34:16 -07:00
Andrej Karpathy 7d0042d574 make training data configurable and also dont move the engire dataset to device that is not a good idea 2024-05-01 18:15:10 +00:00
Andrej Karpathy 8aa9a9ed53 and also profile lol 2024-05-01 17:56:01 +00:00
Andrej Karpathy 1ed9381fd0 fix cudnnDestroy outside of ifdef 2024-05-01 17:55:36 +00:00
Andrej Karpathy 69d40ebfff add readme and search for cudd in home dir 2024-05-01 17:43:50 +00:00
Andrej Karpathy 35640651b3 updates to makefile: docs, override, etc 2024-05-01 17:32:52 +00:00
Andrej Karpathy 04718fbbd1 round 1 of small fixes for cudnn support 2024-05-01 17:07:10 +00:00
ademeure 1147983b8c Add missing cudnnDestroy() calls 2024-05-01 17:19:51 +01:00
ademeure d9e7a0a378 Fixed profile_gpt2.cu/test_gpt2.cu for cuDNN 2024-05-01 17:03:28 +01:00
ademeure 2b540b61da attention_forward.cu ~matches train_gpt2cu & allocate less memory for stats in "att" 2024-05-01 16:53:01 +01:00
ademeure be54842622 Rename lowp_float to floatX in /dev/cuda/attention_forward.cu 2024-05-01 15:43:28 +01:00
ademeure c778dc0098 slightly increase wte threshold to make test_gpt2.cu pass (?) 2024-05-01 15:40:30 +01:00
ademeure 57c09529d9 Merge branch 'cudnn_try2' of https://github.com/ademeure/llm.c into cudnn_try2 2024-05-01 15:32:13 +01:00
ademeure 7beff72f5f Add makefile changes for cuDNN 2024-05-01 15:31:30 +01:00
ademeure c4ecc04dc3 Merge remote-tracking branch 'karpathy/master' into cudnn_try2 2024-05-01 15:30:00 +01:00
ademeure d508463ad7 Merge remote-tracking branch 'origin/cudnn_try2' into cudnn_try2 2024-05-01 15:27:07 +01:00
ademeure d42ee0d925 Cleanup + do not allocate gradient preatt memory with cuDNN + allocate as small a cuDNN workspace as possible 2024-05-01 15:26:25 +01:00
ademeure 61ad4b15ab Fully working forward+backward cuDNN BF16 Flash Attention in train_gpt2.cu! 2024-05-01 15:02:47 +01:00
Chris Dryden 1806062a62 Removed detail about layernorm inputs being f32, no longer the case 2024-04-30 22:01:14 -07:00
Ross Wheeler 229d70b2d5 uint and designated initializers non-standard for C++
uint isn't defined by the MS compiler and designated initializers aren't part of C++ until C++ 20.
2024-04-30 18:46:54 -07:00
Andrej 40da8cf233 Merge pull request #314 from jrhemstad/patch-3
Add llm.cpp fork to README
2024-04-30 17:03:25 -07:00
Jake Hemstad 78eba5fff3 Add llm.cpp fork to README 2024-04-30 18:52:42 -05:00
Andrej Karpathy 769c911ab6 small tweaks 2024-04-30 23:48:06 +00:00
Andrej bad76f408a Merge pull request #313 from ngc92/small-improvements
fixed potential error and generalized gelu forward
2024-04-30 16:18:08 -07:00
harry 74d5bcd03a Adding a java port of the project 2024-04-30 14:51:06 -07:00
Erik Schultheis 603d862830 fix up bandwidth calculation 2024-04-30 23:52:08 +03:00
Erik Schultheis c3515cfe47 fixed potential error and generalized gelu forward 2024-04-30 23:45:34 +03:00
Andrej Karpathy 7ba1ed8a1a delete packed128 .fp32, use (float) ok 2024-04-30 19:39:20 +00:00
Andrej Karpathy 7fd0659ba8 Pack gelu_forward kernel using the new Packed128 2024-04-30 19:35:23 +00:00
Andrej Karpathy 9141e04693 fixes to the PR, careful with float/floatX etc 2024-04-30 19:35:05 +00:00
Andrej Karpathy 242981eecf Merge branch 'gelu_forwards_packing' of https://github.com/ChrisDryden/llm.c into ChrisDryden-gelu_forwards_packing 2024-04-30 19:06:39 +00:00
Andrej Karpathy 70b4de8700 add comment documenting how to reproduce python reference exactly 2024-04-30 18:51:45 +00:00
Andrej Karpathy 050cbfa42c 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 2024-04-30 18:48:12 +00:00
Andrej Karpathy 44656c3850 allow fp32 precision in the test script as well 2024-04-30 18:38:38 +00:00
Andrej Karpathy b84571f745 also add argparse to force tf32 to zero 2024-04-30 17:54:15 +00:00
Andrej Karpathy 654d6f55c5 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 2024-04-30 17:46:36 +00:00
Andrej Karpathy 52e2ca8378 Merge branch 'master' of github.com:karpathy/llm.c 2024-04-30 16:36:00 +00:00
Andrej 179b159a1d Merge pull request #311 from leloykun/fc--modal-script
Add script to run benchmarks on Modal
2024-04-30 09:35:51 -07:00
Franz Louis Cesista b239b67ae1 add modal script 2024-05-01 00:15:08 +08:00
Andrej Karpathy 9d8a6d13f0 add flag overfit_single_batch useful debugging into train_gpt2.cu 2024-04-30 16:03:33 +00:00
ademeure 435ac928b4 cuDNN Forward Flash Attention is working in both /dev/cuda/ and train_gpt2.cu! (backwards pass is broken for now as a conquence) 2024-04-30 17:00:09 +01:00
Andrej Karpathy 51face88d9 fix bug where backward/step must be outside of amp context 2024-04-30 15:48:40 +00:00
Andrej 7e8f14d2c0 Merge pull request #310 from ngc92/flush-l2-for-benchmark
clear the L2 cache between benchmarks
2024-04-30 07:23:41 -07:00
ademeure 1c516c732d re-commit of my old cudnn forward attention changes 2024-04-30 14:54:35 +01:00
Erik Schultheis d7813d281d clear the L2 cache between consecutive invokations of our microbenchmarks to get reliable results 2024-04-30 12:45:37 +03:00
chinthysl 1f442fdf7f fix gpt2_update call 2024-04-30 09:10:50 +00:00
chinthysl b3e8abdd54 optimizer update per shard and nccl all gather 2024-04-30 09:09:52 +00:00
chinthysl 2d26ec10b2 setting the zero opt configs 2024-04-30 09:08:25 +00:00
chinthysl 3254a51522 Zero Optimizations configs 2024-04-30 09:07:01 +00:00
Christopher Dryden 568615fef1 Added packing for gelu forwards kernel 2024-04-30 03:21:42 +00:00
Andrej 9464f4272e Merge pull request #283 from ChrisDryden/configurationRewrite
Changed ordering of type configuration to easily see unchanged values
2024-04-29 18:58:03 -07:00
Andrej 997864920b Merge pull request #290 from Ricardicus/ci-dev-cuda
Adding the example kernels to CI
2024-04-29 18:53:03 -07:00
Andrej 2490f7896d Merge pull request #298 from karpathy/feature/packed128
Feature/packed128
2024-04-29 16:31:02 -07:00
Andrej Karpathy 906d22f5b6 fix profile gpt2cu but note it currently needs the NO_MULTI_GPU=1 switch i think when calling it with make, todo fix later how we add the mpi flags 2024-04-29 23:27:55 +00:00
Andrej Karpathy af2bc47d22 remove forceinline. blame @arund42 if this ever does not get inlined by the compiler 2024-04-29 23:10:51 +00:00
Andrej Karpathy f02ba4b436 import updated tokenizer properly 2024-04-29 23:02:26 +00:00
Andrej Karpathy 7634f0863c add packed128 to the mainline file 2024-04-29 22:57:59 +00:00
Andrej Karpathy 5e12c66f5a move packed128 to common 2024-04-29 22:55:28 +00:00
Andrej Karpathy f43592954f reset back to precious hyperparameters 2024-04-29 22:44:25 +00:00
Andrej Karpathy ee125cffc4 make it imo clearer that the grid size is one quarter as large as it would otherwise be. because each thread now does 4X work 2024-04-29 22:40:26 +00:00
Andrej Karpathy 083b210dd4 small changes 2024-04-29 22:38:03 +00:00
Andrej Karpathy d9506a7a0a minor changes to packed128 2024-04-29 22:31:19 +00:00
Andrej Karpathy 7eefefb511 Merge branch 'yet-another-gelu' of https://github.com/ngc92/llm.c into ngc92-yet-another-gelu 2024-04-29 22:19:35 +00:00
Andrej Karpathy 39ce9be516 fix missing \n on a print 2024-04-29 21:05:44 +00:00
Andrej Karpathy f66825faba allow using flash attention via argparse and print tok/s 2024-04-29 21:05:30 +00:00
Rickard Hallerbäck 5ad0079fca Reverting changes to Makefile, adding -j4 to make 2024-04-29 22:44:21 +02:00
Andrej 0d6e096b9b Merge pull request #296 from karpathy/separate_tokenizer
separate out our Tokenizer
2024-04-29 13:44:05 -07:00
Andrej 1443edc746 Merge branch 'master' into separate_tokenizer 2024-04-29 13:38:58 -07:00
Andrej Karpathy 70f777617d Single call to NCCL AllReduce and AdamW for bf16 2024-04-29 20:36:29 +00:00
Andrej Karpathy dd60fc3991 fix prints 2024-04-29 20:36:07 +00:00
Andrej Karpathy be4aeb3d2d fix accumulated loss 2024-04-29 20:32:10 +00:00
Andrej Karpathy 938f8f78bb more defensive 2024-04-29 20:25:29 +00:00
Andrej Karpathy b4c346ac2a don't hardcode the EOT token and store it in the Tokenizer file 2024-04-29 20:23:41 +00:00
ademeure c24bb88bce Add back FP16 and disable multi_gpu all reduce when there is only 1 GPU 2024-04-29 21:15:52 +01:00
Andrej Karpathy 699c25469a separate out our Tokenizer 2024-04-29 19:54:18 +00:00
Erik Schultheis 4872c57164 comments 2024-04-29 22:46:33 +03:00
Erik Schultheis 9f07a176e0 simplify to only support cs 2024-04-29 22:43:52 +03:00
ademeure 66c454874f avoid warning for cublas_compute_type in BF16 mode 2024-04-29 20:25:07 +01:00
ademeure 1d4effd5a4 Remove FloatN and associated code for adam & allreduce 2024-04-29 20:21:34 +01:00
Erik Schultheis 896b644cb9 slight simplification 2024-04-29 22:19:12 +03:00
Andrej 1310e2f77d Merge pull request #286 from rosslwheeler/Makefile-windows-support
Makefile windows support
2024-04-29 12:11:12 -07:00
Andrej 977c688b9d Merge pull request #294 from karpathy/ademeure-layernormbf16-3
Ademeure layernormbf16 3
2024-04-29 12:10:07 -07:00
Andrej Karpathy 0e4506ac30 unify these two 2024-04-29 19:03:49 +00:00
ademeure 8be7370a55 gpt2_build_from_checkpoint 2024-04-29 19:54:32 +01:00
Erik Schultheis 45252d434c yet another gelu 2024-04-29 21:53:08 +03:00
ademeure ebedadfee8 Fix train_gpt2.py to write all weights as BF16 (+merge fixes) 2024-04-29 19:49:57 +01:00
ademeure d23bb1fd13 Merge remote-tracking branch 'karpathy/master' into layernormbf16 2024-04-29 19:41:54 +01:00
Andrej e92ec3ffd9 Merge pull request #281 from ChrisDryden/templateCleanup
Removed templates from kernels to clarify data types and simplify code
2024-04-29 10:05:29 -07:00
Andrej 237593f47b Merge pull request #287 from tairov/fix-zig-link
Fix link to zig port in README.md
2024-04-29 09:04:18 -07:00
Ross Wheeler fcef97d6a1 Adding back PFLAGS
Got missed during the merge

Missed the profiler PFLAGS

Move PFLAGS ahead to match original change
2024-04-29 08:53:50 -07:00
Rickard Hallerbäck 8915d5ab0d Adding the example kernels to CI 2024-04-29 17:02:05 +02:00
ademeure 49ffd68317 Fix sizeof(T) -> sizeof(D) (was overallocating memory) 2024-04-29 15:47:59 +01:00
Aydyn Tairov d0b943ba04 Fix link to zig port 2024-04-29 11:18:03 +01:00
Ross Wheeler 81b689ed01 Merge branch 'Makefile-windows-support' of https://github.com/rosslwheeler/llm.c into Makefile-windows-support 2024-04-29 03:00:56 -07:00
Ross Wheeler c9ab8ee123 Updating Makefile with latest changes
Add check with Windows OS
2024-04-29 03:00:46 -07:00
Ross Wheeler 4b80126913 Add check with Windows OS 2024-04-29 02:46:31 -07:00
Eymen Ünay 0e38ac7b8c Refactoring parameter size filling
Similar to  a new function is defined for parameter size table. This will be useful for reuse with e.g. initializing weights from scratch instead of a checkpoint.
2024-04-29 12:35:21 +03:00
Ross Wheeler f3b1bb70df Updating Makefile with latest changes 2024-04-29 02:28:03 -07:00
Christopher Dryden cf3e6efc79 Changed ordering of type configuration to easily see unchanged values 2024-04-29 02:28:23 +00:00
Christopher Dryden 109c62c6c6 Removed templates from kernels to clarify data types 2024-04-29 02:06:59 +00:00
ademeure 1e50f3b976 Only enable kernel4 with BF16 2024-04-29 01:55:06 +01:00
ademeure 82fef03acd Use kernel6 for train_gpt2, and add all other kernels to /dev/cuda (+bugfixes) 2024-04-29 01:51:10 +01:00
Andrej Karpathy 50acc125f3 Merge branch 'ngc92-split-file'
Separates out common error-checking wrapper utils, that are broadly useful across all file
2024-04-28 23:27:08 +00:00
Andrej Karpathy c20497cfc1 add comments pointing to the definition of the utils functions 2024-04-28 23:26:58 +00:00
Andrej Karpathy b1c80e9d7b Merge branch 'split-file' of https://github.com/ngc92/llm.c into ngc92-split-file 2024-04-28 23:23:00 +00:00
Erik Schultheis 4a3c278f3a moved checked helper functions into a separate file 2024-04-29 00:00:37 +03:00
Andrej 18b41b4df6 Merge pull request #279 from Ricardicus/prec-ci
Updating the CI to build different precisions
2024-04-28 13:47:48 -07:00
Andrej Karpathy 66a92c6791 Merge branch 'ngc92-cleanup' 2024-04-28 20:45:04 +00:00
Andrej Karpathy 518565654e Merge branch 'cleanup' of https://github.com/ngc92/llm.c into ngc92-cleanup 2024-04-28 20:43:37 +00:00
Andrej Karpathy 49228b091a add small comment on -t=0 2024-04-28 20:20:51 +00:00
Erik Schultheis 4c295c7e97 and even more cleanup 2024-04-28 23:19:36 +03:00
Erik Schultheis ca48791522 as promised, cleanup enabled by padding :) 2024-04-28 23:19:36 +03:00
Rickard Hallerbäck b522333851 Updating the CI to build different precisions 2024-04-28 22:17:47 +02:00
Andrej 10aa24e2cc Merge pull request #269 from ChrisDryden/patch-3
Enable multithreading in nvcc
2024-04-28 13:16:47 -07:00
Andrej Karpathy 4b6f68a9a9 Merge branch 'lancerts-encoder_forward-float4' 2024-04-28 20:10:48 +00:00
Andrej Karpathy 327eef3f22 incorporate faster encoder_forward kernel to fp32 CUDA version 2024-04-28 20:10:33 +00:00
Andrej Karpathy 4b6a532b9b Merge branch 'encoder_forward-float4' of https://github.com/lancerts/llm.c into lancerts-encoder_forward-float4 2024-04-28 19:49:16 +00:00
Andrej Karpathy b7972ff928 make padded vocab fixes in the .c code as well, i missed it in the previous PR, should satisfy the CI now 2024-04-28 19:33:21 +00:00
Andrej Karpathy 835060e18d padded vocab change. touched a lot of code. very stressful and error prone, but i think it is done. had to bump versions on all .bin files, invalidating the previous files. re-run the python training script to re-export the new version files. let's not do too much of things like this in the future lol. actually, fun fact i had a chance to do the padded vocab really really early in the history of llm.c development, and chose not do it, thinking i'll just do it later. i should have done it. such is life, you make mistakes, you accumulate scar tissue, and you learn, and you become better, faster, stronger. this is the mindset one must have to lead a happy and fulfilling life. it's not important that you are perfect at any point in time, it's only important that you keep improving, every day. 2024-04-28 18:47:03 +00:00
Andrej d95b8d8320 Merge pull request #265 from karpathy/feature/load_bf16
load bf16 directly, and some "quality of life" handling of fp32/fp16/bf16 precisions
2024-04-28 09:24:56 -07:00
Andrej Karpathy 9d70d9a1b2 profile and test only use bf16. but the train script can be run with fp32 or bf16 or fp16. fp16 will error, though 2024-04-28 16:17:22 +00:00
Andrej Karpathy a3f5ad9095 reshuffle the ifdefs to make bf16 the default if no PRECISION is requested via defines 2024-04-28 16:12:39 +00:00
Andrej Karpathy 4f7d8d9f44 allow user to make different precisions, add prints and error handling around precisions 2024-04-28 16:08:27 +00:00
lancer 1c7d23ab3c amend the float4 kernel 2024-04-27 20:26:19 -07:00
lancer 61c5c05ea8 amend the float4 kernel 2024-04-27 20:23:39 -07:00
lancer cfccd821e6 Include the float4 kernel 2024-04-27 19:58:18 -07:00
Andrej Karpathy 82d7907214 adjust comment 2024-04-28 00:05:03 +00:00
Andrej Karpathy 9a91b40207 bring back original ordering. i also had to bump the thresholds by 3X for some tensors and i don't exactly know why sad 2024-04-28 00:02:51 +00:00
ademeure 1d93edf010 turn on full BF16 on by default 2024-04-28 00:32:54 +01:00
ademeure 8768f6bd74 Full BF16 including layernorm at good perf by minimising BF16 atomics 2024-04-28 00:31:42 +01:00
Andrej Karpathy 00627070aa fix a really bad bug in how i was checking the gradients, where i loaded them in the old order, so yeah... 2024-04-27 23:17:22 +00:00
Chris Dryden 2954d90f30 Enable multithreading in nvcc
Tested locally and reduced compilation time by 200ms, unfortunately for me upgrading to 12.4 made my compilations times slow by 2x but at least this can make it a bit faster
2024-04-27 14:14:02 -07:00
Andrej Karpathy a58b8d5dc7 print more in the comparison 2024-04-27 16:04:55 +00:00
Andrej Karpathy 9d6fd30e3c tweak the tolerances until we pass lol 2024-04-27 15:49:54 +00:00
Andrej Karpathy e067a27395 fix dumb bug. i'll blame github copilot but i can't remember 2024-04-27 15:40:32 +00:00
Andrej Karpathy d4a642ba19 i think github copilot betrayed me on this index here, i cant remember 2024-04-27 15:38:20 +00:00
Andrej Karpathy 09d935c82b i think i am making things cleaner, but i am not fixing the problem 2024-04-27 15:27:05 +00:00
Andrej Karpathy 12da2c16cc slight docs tweaks to dev/cuda 2024-04-27 14:40:57 +00:00
Andrej a5d23e7f97 Merge pull request #266 from AnswerDotAI/master
Add a local convenience Makefile for dev/cuda/
2024-04-27 07:18:22 -07:00
austinvhuang 25d703fbec Use a make pattern rule to simplify Makefile, group kernels into forward/backward/updates/nccls, add some comments and update the README.md 2024-04-27 00:13:01 -04:00
austinvhuang d937f93aaa Add a local convenience Makefile for dev/cuda/, fix 2 minor warnings observed when building all targets. 2024-04-26 21:19:44 -04:00
Andrej Karpathy 09cd67e868 code to load bf16 weights directly, and also re-wire the position of tensors to put the layernorms at the end. the training loop seems to work ok, and the tests pass and the loss and optimization looks ok, but the gradients don't match. which can't be right. so there is a bug, but it's a bit too late in the day for me to debug right now, creating a PR and going to sleep, will fix tomorrow 2024-04-27 00:54:06 +00:00
Andrej 311d6d8441 Merge pull request #260 from ChrisDryden/patch-2
Remove unused helper functions
2024-04-26 15:32:59 -07:00
Andrej 7804a971cd Merge pull request #259 from karpathy/feature/fix_int_overflow
candidate fix for int overflow
2024-04-26 15:32:43 -07:00
Andrej Karpathy 9c2209254b also improve situation in the c code 2024-04-26 22:27:22 +00:00
Andrej Karpathy b019927100 change more ints to size_t for a bit more safety, still not perfect most likely 2024-04-26 22:20:15 +00:00
Chris Dryden 1e2df12f14 Remove unused helper functions
Removing two helper functions that are no longer used
2024-04-26 14:21:35 -07:00
Andrej Karpathy 689ccb8ec8 candidate fix for int overflow 2024-04-26 20:20:53 +00:00
Andrej ad602c59ff Merge pull request #239 from dagelf/master
Add to .gitignore: fp32 binaries, *.exe, *.obj, .venv
2024-04-26 13:07:17 -07:00
Andrej Karpathy 085270618d upgrade our .py script so it can also do torchrun with many processes, as our mixed precision mpi nccl code can do now 2024-04-26 19:50:37 +00:00
Coenraad Loubser 8e8db3d817 Merge branch 'master' into master 2024-04-26 20:44:44 +02:00
Andrej Karpathy 0c3e3e37f4 Merge branch 'PeterZhizhin-multi_gpu_support' 2024-04-26 18:23:43 +00:00
Andrej Karpathy aedf3a5cdd changes to readme for multigpu. i think this is ready now for v0 merge to master nice 2024-04-26 18:20:15 +00:00
Andrej Karpathy 29add03b64 also adjust gitignore 2024-04-26 18:14:55 +00:00
Andrej Karpathy e5fd817de8 revert fp32 version to no multigpu support 2024-04-26 18:12:44 +00:00
Andrej Karpathy a00a3a88ef fix bug with all_hostsname_hashes missing, re-arrange code a bit, make sure we only print on rank0 2024-04-26 18:11:09 +00:00
Andrej Karpathy 2e17140d33 more defensive check for openmpi install, and an install command 2024-04-26 16:27:57 +00:00
Andrej Karpathy 8389fbaa0c Merge branch 'multi_gpu_support' of https://github.com/PeterZhizhin/llm.c into PeterZhizhin-multi_gpu_support 2024-04-26 15:55:15 +00:00
Andrej Karpathy 8ac4b47696 for now save both fp32 and bf16. todo let's think this through a bit more maybe 2024-04-26 04:36:53 +00:00
Coenraad Loubser d195e1c3e0 Merge branch 'karpathy:master' into master 2024-04-25 21:07:28 +02:00
Andrej Karpathy bb56144384 we are an epsilon away from writing our model in bf16 as well, in addition to fp32, with the re-ordered layernorms 2024-04-25 18:32:08 +00:00
Andrej Karpathy 3fb7252924 round 1 of some changes. we will now always write in fp32, even if dtype is set to float16 or bfloat16. next up, we actually want to write in lower precision, when the dtype is set so 2024-04-25 17:38:01 +00:00
Andrej Karpathy 7a52a21fcb Merge branch 'master' of https://github.com/karpathy/llm.c 2024-04-25 16:31:44 +00:00
Andrej Karpathy 7fe1a4cd49 add device argparse 2024-04-25 16:31:01 +00:00
dagelf f36891675e Allow data directory to be a softlink and still ignore it 2024-04-25 11:43:03 +02:00
dagelf bcddcd6587 if / specified git doesn't ignore softlinks to subdirectories 2024-04-25 11:40:47 +02:00
Coenraad Loubser 819df1cd0f Delete .venv 2024-04-25 11:37:01 +02:00
dagelf 7476b9ec60 more specific gitignore 2024-04-25 11:33:16 +02:00
Andrej 716a2eddac Merge pull request #247 from joshcarp/patch-1
Add reference to llm.go in description
2024-04-24 20:38:05 -07:00
Andrej 1555c080aa Merge pull request #249 from ngc92/small-fixes
Small fixes
2024-04-24 20:35:35 -07:00
Erik Schultheis 533054b393 corrected calculation of num_activations 2024-04-25 01:39:24 +03:00
Erik Schultheis 45e17e3ced don't generate atomic functions for (b)float16 if those are not enabled 2024-04-25 01:39:08 +03:00
Petr Zhizhin d49e8301eb [Multi-GPU] llm.c now runs on multiple GPUs with NCCL 2024-04-24 22:11:00 +00:00
Joshua Carpeggiani 0ee2d4a965 Add llm.go in description 2024-04-24 17:48:49 -04:00
Andrej Karpathy e44ed15639 small updates on intent for fp32 and non-fp32 CUDA versions 2024-04-24 19:44:12 +00:00
Andrej Karpathy 98567a4509 turn on benchmarking in mixed precision. example command: python train_gpt2.py --write_tensors 0 --inference_only 0 --dtype bfloat16 --compile 1 --tensorcores 1 --num_iterations 50 --sequence_length 1024 2024-04-24 19:26:40 +00:00
Andrej b2c779d40e Merge pull request #231 from Saimirbaci/zig-port-ref
Added ref to the zig port
2024-04-24 10:00:08 -07:00
dagelf cd898b79d3 .add to .gitignore: fp32 binaries, *.exe, *.obj, .* 2024-04-24 17:58:49 +02:00
Andrej Karpathy 250159e140 encoder backward add, we want to get rid of atomics, hmm 2024-04-23 19:06:15 +00:00
Andrej Karpathy fd7da62564 minor comment fixes, more to come 2024-04-23 18:20:50 +00:00
Andrej Karpathy 249140287f rename positional_forward to encoder_forward as it should have been 2024-04-23 18:19:20 +00:00
Andrej Karpathy 0b59d80880 not proud of this but i have to increase the threshold for now to make test pass. later we should come back around and tighten thresholds 2024-04-23 18:13:57 +00:00
Andrej Karpathy 754627c0a2 update ci to pick up only the fp32 version for now, todo add mixed precision 2024-04-23 17:56:46 +00:00
Andrej Karpathy c5f1db6020 fix important typo on precision 2024-04-23 17:43:19 +00:00
Andrej Karpathy ed6387a695 updates to readme, and introduce the test fp32 cuda file too 2024-04-23 17:41:42 +00:00
Saimir Baci 7f279cd6cb Adding Zig port ref 2024-04-23 13:37:53 -04:00
Andrej Karpathy 095d27662b checkpoint the fp32 CUDA implementation to separate file. our mainline iteration will now continue in the new (mixed precision) file 2024-04-23 17:21:44 +00:00
Andrej 6b6ad35f29 Merge pull request #218 from ademeure/linear16
Support for FP16/BF16 in train_gpt2.cu (1.86x Perf)
2024-04-23 10:17:01 -07:00
ademeure 0333e36981 tiny fixes to make fp8 work in the future (missing conversions) 2024-04-23 14:38:41 +01:00
ademeure 91ec92f4d5 fix typo 2024-04-23 06:43:40 +01:00
ademeure a876485151 Fill in param_sizeof in a slightly less terrible way 2024-04-23 06:39:58 +01:00
ademeure 53dc40e804 Fixes for merge with latest, now ~86% faster! 2024-04-23 06:28:42 +01:00
ademeure c1992a19d5 Merge remote-tracking branch 'karpathy/master' into linear16 2024-04-23 06:18:23 +01:00
ademeure 8775856b1d Added floatN to choose FP32 vs FP16 for layernorm (parameters/gradients only) 2024-04-23 06:13:23 +01:00
ademeure 7c193bd4cc Good news: it works. Bad news: 16-bit atomics are *incredibly slow* :( (is it doing CAS?) 2024-04-23 05:00:47 +01:00
ademeure bdc661a609 BF16/FP16 Attention + bug fixes 2024-04-23 02:26:57 +01:00
ademeure 85290a5228 Updated FP16/BF16 mix-and-match ncluding stochastic rounding 2024-04-23 01:39:33 +01:00
Andrej 896f83f422 Merge pull request #225 from dagelf/master
Make msvc compile by adding the ifdefs back in gelu_backward, fixes #200
2024-04-22 14:45:29 -07:00
Coenraad 2efac6670a Make msvc compile by adding the ifdefs back in gelu_backward, fixes #200 2024-04-22 23:17:06 +02:00
Andrej Karpathy 451c1a4c99 update readme to point to SoU 2024-04-22 19:18:03 +00:00
Andrej Karpathy d3c5025a3e Merge branch 'ChrisDryden-sharedmem_layernormback' 2024-04-22 18:45:01 +00:00
Andrej Karpathy e3bcae6b2b incorporate layernorm backward kernel 2, we are now down to 73.1ms/iter on average 2024-04-22 18:44:43 +00:00
Andrej Karpathy 86b3401dd5 Merge branch 'sharedmem_layernormback' of https://github.com/ChrisDryden/llm.c into ChrisDryden-sharedmem_layernormback 2024-04-22 18:33:42 +00:00
Andrej Karpathy 5f545ca6cc Merge branch 'al0vya-matmul-backward-bias' 2024-04-22 17:51:26 +00:00
Andrej Karpathy 988489519c nice new kernel for bias reduce, we are down by ~1ms/iter, to 76.13ms/iter on average 2024-04-22 17:51:06 +00:00
ademeure d605b993da Fully working BF16 training! (very hacky Adam, and converges much slower than FP32 for now) 2024-04-22 18:38:51 +01:00
ademeure 58df2629ae Gradients now working in BF16 mode! (still need adam etc...) 2024-04-22 18:09:18 +01:00
Andrej Karpathy a42f739675 Merge branch 'matmul-backward-bias' of https://github.com/al0vya/llm.c into al0vya-matmul-backward-bias 2024-04-22 16:52:31 +00:00
Andrej f813d635b2 Merge pull request #222 from regrettable-username/patch-1
Added llm.metal to notable forks
2024-04-22 09:33:27 -07:00
Andrej Karpathy b1e55951c6 rename to _forward for the kernel, and small fixes to docs, and avoid auto 2024-04-22 16:32:03 +00:00
cip19aac 35393b4442 add more comments to explain the philosophy behind the kernel 2024-04-22 17:29:11 +01:00
cip19aac b82ec201be add comment 2024-04-22 17:05:05 +01:00
Andrej 7830cf6a12 Merge pull request #214 from ngc92/trimul
added triangular matrix multiplication kernel
2024-04-22 09:03:53 -07:00
James Thompson 38857ff1d0 Added llm.metal to notable forks
This is a fork of llm.c designed to take full advantage of Apple Silicon via the Metal framework. Just got the forward pass working and will start on the backward pass. I wrote a minimal C wrapper API over the Objective-C Metal API that was needed to set up and run compute shaders. The rest of the code is pure-C and Metal Shading Language. I'm trying to mirror the original repo as closely as possible both in code and project philosophy.
2024-04-22 09:01:21 -07:00
cip19aac 55fbd7e6ad add a faster matmul backward bias kernel that uses coalesced reads and shared memory 2024-04-22 16:35:21 +01:00
ademeure f35adbec77 WIP support for FP16/BF16 in train_gpt2.cu (compiles, not correct results) 2024-04-22 16:21:28 +01:00
Erik Schultheis 732a8b4806 added NaN-based makes for reference checks 2024-04-22 14:53:24 +03:00
Erik Schultheis caa69a5e52 added triangular matrix multiplication kernel 2024-04-22 14:45:51 +03:00
Andrej Karpathy 6984e83c6e Merge branch 'patricxu-master' 2024-04-22 00:28:00 +00:00
Andrej Karpathy 9ed0f10ee0 Merge branch 'master' of https://github.com/patricxu/llm.c into patricxu-master 2024-04-22 00:26:08 +00:00
Andrej fe85da4f3e Merge pull request #209 from Ricardicus/master
Adding a github workflow
2024-04-21 17:22:53 -07:00
Rickard Hallerbäck 3d2605bf8a Merge branch 'master' into master 2024-04-22 02:21:49 +02:00
Andrej Karpathy 4e4eba0441 simplify bool 2024-04-22 00:19:00 +00:00
Andrej a332ea5d92 Merge pull request #208 from Ricardicus/train_gpt2.py-fix
train_gpt2.py: fix for conditional use of fused and for computing gradients
2024-04-21 17:18:13 -07:00
Rickard Hallerbäck c7aad65a04 Adding a workflow that builds with and without CUDA and OMP and on CPU also tests on ubuntu and macos 2024-04-22 02:14:09 +02:00
Chris Dryden 2d97be8315 Accidental last minute change 2024-04-21 16:44:07 -07:00
Chris Dryden 2752212151 Added shared memory for the atomic additions for the layernorm_backward.cu weights and biases 2024-04-21 16:37:08 -07:00
Rickard Hallerbäck 6c163d3391 train_gpt2.py: fix for conditional use of fused and for computing gradients 2024-04-22 00:27:36 +02:00
Andrej 4f8f5fa833 Merge pull request #207 from ToJen/patch-1
add Rust port to notable forks
2024-04-21 15:06:08 -07:00
Tomisin Jenrola 43ffd2af8a add Rust port to notable forks 2024-04-21 19:23:50 -02:30
Andrej Karpathy 5ca1b939e9 small type reshuffles 2024-04-21 19:28:17 +00:00
Andrej Karpathy 48f2170e7b pad step so that we don't change horizontal size in early parts of the script 2024-04-21 18:53:08 +00:00
Andrej Karpathy 8d4f93ccbd experiment with pretty printing the run config, model config, device config, and data config 2024-04-21 18:47:36 +00:00
Andrej Karpathy e55ed0a4e6 change default lr to 3e-4, exp -> expf, and print tok/s 2024-04-21 18:22:45 +00:00
Andrej e93438efea Merge pull request #201 from ngc92/memory
Yet more memory reductions
2024-04-21 11:02:50 -07:00
Andrej 734f65f33e Merge pull request #192 from rosslwheeler/const_fixes
const changes to help with Microsoft compiler OpenMP 5.2 support
2024-04-21 10:42:15 -07:00
Andrej Karpathy 14d96e7c0d Merge branch 'Ricardicus-master' 2024-04-21 17:41:24 +00:00
Andrej Karpathy fced6d180f small adjustments to Makefile and gitignore 2024-04-21 17:40:26 +00:00
Andrej Karpathy e70bc19af5 Merge branch 'master' of https://github.com/Ricardicus/llm.c into Ricardicus-master 2024-04-21 17:34:05 +00:00
Erik Schultheis fb3bea9e2a more aggressive reuse of no longer needed memory from the forward activations 2024-04-21 19:48:21 +03:00
Erik Schultheis 8f6d1caefd merged output and scratch 2024-04-21 19:31:46 +03:00
Erik Schultheis 118f3dde12 removed yet more tensors from backward state 2024-04-21 19:23:25 +03:00
Erik Schultheis da5e8f7f56 split ActtivationTensors into separate classes for forward and backward to improve readability 2024-04-21 19:10:47 +03:00
Erik Schultheis e3cbe7d886 saved another 60 MiB 2024-04-21 18:53:25 +03:00
Erik Schultheis 2ed8eddb3e slight optimization of the python script 2024-04-21 18:43:45 +03:00
Erik Schultheis 33e81ba55f print peak memory consumption in python 2024-04-21 18:25:04 +03:00
Erik Schultheis bfdbf50189 more use of scratch buffer 2024-04-21 18:16:12 +03:00
Erik Schultheis 10fc806900 replace vaccum with generic scratch buffer 2024-04-21 17:10:29 +03:00
Erik Schultheis 0636bc4ee3 removed the union and moved the logit->probability conversion into the sampling function 2024-04-21 16:45:46 +03:00
patricxu 41c3cc1fab Merge branch 'karpathy:master' into master 2024-04-21 10:31:59 +08:00
Erik Schultheis 7f3272c38f unionize output buffer 2024-04-21 02:11:41 +03:00
Andrej Karpathy 9fb9c9185a change baseline kernel to float, for now, for fair comparison too 2024-04-20 22:57:18 +00:00
Erik Schultheis 2f81c6cbda qkv is only needed once, not per-layer 2024-04-21 01:50:10 +03:00
Andrej Karpathy 8488669d25 speed up the backward bias kernel by 45% and speed up the full running time by 1% 2024-04-20 22:46:19 +00:00
Rickard Hallerbäck 3cd6f83e32 removing some warnings 2024-04-20 22:03:35 +02:00
Andrej Karpathy 9b722ce5c0 Merge branch 'rosslwheeler-unistd_h_in_dev' 2024-04-20 20:00:39 +00:00
Andrej Karpathy 19e337b99e a comment and some breathing room 2024-04-20 20:00:24 +00:00
Rickard Hallerbäck be106f09a4 Merge branch 'master' into master 2024-04-20 22:00:02 +02:00
Andrej Karpathy 7e42c8b356 Merge branch 'unistd_h_in_dev' of https://github.com/rosslwheeler/llm.c into rosslwheeler-unistd_h_in_dev 2024-04-20 19:59:11 +00:00
Andrej aa1184d9ca Merge pull request #199 from ngc92/fused-buffer
Further memory reductions
2024-04-20 12:54:37 -07:00
Andrej Karpathy 0e79ac7d8b Merge branch 'dagelf-master' 2024-04-20 19:45:49 +00:00
Andrej Karpathy adac21e97b add a small comment as to what this is 2024-04-20 19:45:34 +00:00
Andrej Karpathy 48a9f2ec69 Merge branch 'master' of https://github.com/dagelf/llm.c into dagelf-master 2024-04-20 19:43:17 +00:00
Andrej Karpathy 7fa9f7b140 document one way to do experiments / sweeps as just an example 2024-04-20 19:39:12 +00:00
Coenraad f534e4bdfe Fixes -Ofast optimizations breaking model by skipping them for gelu_backward 2024-04-20 21:33:00 +02:00
Rickard Hallerbäck 280023bf8e Support for vanilla build without CUDA tools 2024-04-20 21:32:35 +02:00
Andrej Karpathy 4e3b7ea0f3 add a first, super simple logger 2024-04-20 19:26:25 +00:00
Andrej Karpathy acdcfeaab6 introduce a poor mans argparse because i want to start training on tinystories, and messing with some of the options of the runs, now that things are more efficient 2024-04-20 18:37:18 +00:00
Andrej 5b0d52619f Merge pull request #195 from ngc92/prof-script
profile and print out some nice stats
2024-04-20 10:48:10 -07:00
Andrej af1ee01768 Merge pull request #197 from ngc92/fused-attention-slow
slow attention kernel that demonstrates a direct translation of the CPU code
2024-04-20 10:40:04 -07:00
Erik Schultheis 24dd728953 only need one V-sized buffer that can switch from logits to dlogits in the fused backward pass 2024-04-20 20:19:16 +03:00
Erik Schultheis af9ac14b2f don't track preatt and vaccum across layers 2024-04-20 20:18:41 +03:00
Erik Schultheis 8ffc3e388b slow attention kernel that demonstrates a direct translation of the CPU code 2024-04-20 19:03:34 +03:00
Erik Schultheis d12b923c00 profile and print out some nice stats 2024-04-20 13:30:31 +03:00
Ross Wheeler d24562932f Adding unistd.h for Windows in Dev directory
Adding unistd.h for Windows in Dev directory
2024-04-20 02:37:16 -07:00
Ross Wheeler e564222113 const changes to support OpenMP 5.2 MSVC work-around
const changes to support OpenMP MSVC work-around
2024-04-20 02:07:12 -07:00
patricxu d48b0d2b2b fix indexing in matmul_forward_kernel1 2024-04-20 15:38:12 +08:00
Andrej 2d3a0fbe50 Merge pull request #181 from lancerts/const-fix-adamw
const-fix adamw.cu and fix the warning C4305: 'initializing': truncation from 'double' to 'float'
2024-04-19 22:06:11 -07:00
Andrej Karpathy d34556cf5e more profile docs 2024-04-20 01:29:14 +00:00
Andrej Karpathy 67fd37ace4 add crappy instructions for profiling basics 2024-04-20 01:06:53 +00:00
Andrej Karpathy 7b79fbb230 Merge branch 'master' of github.com:karpathy/llm.c 2024-04-20 00:59:39 +00:00
Andrej Karpathy 7df2d4d92b also requirements for quick start 2024-04-20 00:59:32 +00:00
Andrej 694cb571ec Merge pull request #189 from ngc92/profile-target
profiling script
2024-04-19 17:55:35 -07:00
Andrej Karpathy 91aab62aad fix kernel 5 of layernorm, bug from previous copy paste with loop bounds 2024-04-20 00:51:06 +00:00
Andrej Karpathy 5bdbc45806 only do 5 validation batches by default, we really don't need much precision here for this little script with current defaults 2024-04-20 00:45:13 +00:00
Andrej Karpathy 38fd42d3e3 also add a quick start CPU section 2024-04-20 00:41:13 +00:00
Andrej Karpathy da4c5c79ca very quick start section add 2024-04-20 00:35:35 +00:00
Erik Schultheis 9c1c1f0cec helper target that gets compiled with lineinfo and runs a single layer forward-backward pass 2024-04-20 03:19:23 +03:00
Andrej Karpathy 49d41ae296 add one more kernel, allocating a block per row. bad idea if C is too low, as we have it right now 2024-04-19 23:24:51 +00:00
Andrej Karpathy cb791c4ef5 new kernel that does a single pass over x on load, using a more clever variance formula. only very slightly faster on my A100 sadly 2024-04-19 22:55:30 +00:00
Andrej 816254e454 Merge pull request #187 from azret/csharp-port-ref
README: Adding C# port ref
2024-04-19 14:17:59 -07:00
Azret Botash c6bd590b47 typo 2024-04-19 14:17:13 -07:00
Azret Botash 6cd1354b3e C# port ref 2024-04-19 13:59:36 -07:00
Andrej 792812985b Merge pull request #185 from dorjeduck/update-readme-notable-forks
Added link to Mojo port in Notable Forks section
2024-04-19 13:22:32 -07:00
Andrej Karpathy b173e0237c Merge branch 'ngc92-and-one-more' 2024-04-19 20:21:50 +00:00
Martin Dudek 1f08134fb9 Added link to Mojo port in Notable Forks section 2024-04-20 01:44:50 +05:30
Erik Schultheis 44b5bb028c one more version of attention backward for another few percent perf gains 2024-04-19 22:46:40 +03:00
Erik Schultheis a360f104ae some low-level optimizations 2024-04-19 22:38:13 +03:00
Erik Schultheis 7005e09450 added missing checks 2024-04-19 22:37:35 +03:00
Andrej bba794f156 Merge pull request #180 from ngc92/yet-more-attention
separating the loops leads to yet more improvements
2024-04-19 11:05:55 -07:00
lancer 3cb993b5f6 fix the constness and the warning C4305: 'initializing': truncation from 'double' to 'float' 2024-04-19 11:04:58 -07:00
Erik Schultheis b2e6b6867f separating the loops leads to yet more improvements 2024-04-19 20:47:19 +03:00
Andrej Karpathy 810b0dd90f adjust header 2024-04-19 17:19:26 +00:00
Andrej Karpathy f7fb3e67d5 update readme 2024-04-19 17:16:48 +00:00
Andrej Karpathy 1dbf241fbe cleanups of the main script 2024-04-19 17:02:46 +00:00
Andrej Karpathy 09a117a664 add a compile option for torch compile suggested by Horace 2024-04-19 16:32:24 +00:00
Andrej b556ad971c Merge pull request #179 from ngc92/even-better-attention
towards an even better backward attention kernel
2024-04-19 09:16:22 -07:00
Andrej 1d3b6738a7 Merge pull request #169 from lancerts/fix-crossentropy-bk
Fix the bug that yields cpu, gpu results mismatch in crossentropy_softmax_backward.cu
2024-04-19 09:15:23 -07:00
Andrej c8f3981d6b Merge pull request #175 from lancerts/fix-double-f
Fix the 0.0f initialization for double
2024-04-19 09:12:58 -07:00
Erik Schultheis 50e105a250 further cleanup and ability to handle arbitrary sequence lengths 2024-04-19 19:12:37 +03:00
Erik Schultheis 4294820a8a don't do stupid redundant work; the inner loop is just a block-level sum!!!! 2024-04-19 19:12:32 +03:00
Erik Schultheis 02aa8e47dc further simplification of the loop 2024-04-19 19:12:32 +03:00
Erik Schultheis 16123c3af8 towards an even better backward attention kernel 2024-04-19 19:12:30 +03:00
lancer 0c92ea9b90 fix the 0.0f initialization for double 2024-04-18 19:35:34 -07:00
lancer e991fe1efa fix the 0.0f initialization for double 2024-04-18 19:33:30 -07:00
Andrej Karpathy 0dbfa32fd2 sorry this is a double so i shouldn't use 0.0f. i have to dedicate a whole commit to this because of my OCD 2024-04-19 02:23:21 +00:00
Andrej Karpathy 315af5f5dd micro-optimization on the order, maximizing cache re-use ty @YichengDWu for original PR, which I adapted a bit 2024-04-19 02:21:28 +00:00
Andrej Karpathy c393d828dc size_t instead of int. everywhere. this will come to bite us at one point... 2024-04-19 02:11:30 +00:00
Andrej Karpathy 213ed60afd fused classifier! we don't materialize the probs. we go from 200ms to 190ms 2024-04-19 01:33:35 +00:00
Andrej Karpathy dcff5613c6 add kernel3, which is like kernel 2 but no float4 2024-04-19 01:19:20 +00:00
Andrej Karpathy 8a2cab3431 make default dloss be 1/BT 2024-04-19 00:27:28 +00:00
Andrej Karpathy ea3e2abdcf small fixes and more comments very cool kernels 2024-04-18 23:48:39 +00:00
Andrej Karpathy 70170006ca adjust readme with latest speed 2024-04-18 22:50:52 +00:00
Andrej Karpathy 6b603fa6c2 Merge branch 'ngc92-bwd-att-coarsened' 2024-04-18 22:48:24 +00:00
Andrej Karpathy d95e62436b Merge branch 'bwd-att-coarsened' of https://github.com/ngc92/llm.c into ngc92-bwd-att-coarsened 2024-04-18 22:47:24 +00:00
Erik Schultheis 0d55421f32 updated main training script 2024-04-19 01:01:47 +03:00
Erik Schultheis 09d35a5e4d splitting the loop to get rid of the indicator 2024-04-19 01:01:47 +03:00
Erik Schultheis 6d98af1d32 another 15% by getting rid of ifs 2024-04-19 01:01:47 +03:00
Erik Schultheis 314661f516 significantly faster kernel by reusing data in registers 2024-04-19 01:01:47 +03:00
Andrej Karpathy 601e371e78 change MB to MiB when reporting capacity 2024-04-18 21:59:41 +00:00
Andrej Karpathy 5c8eaec359 report all big mallocs 2024-04-18 21:56:29 +00:00
Andrej Karpathy 6c0958eb79 Merge branch 'ngc92-memory-savings-2' 2024-04-18 21:43:02 +00:00
Andrej Karpathy 2981467415 delete a bunch of buffers that are not necessary in the backward pass, and also delete residual_backward 2024-04-18 21:42:21 +00:00
Erik Schultheis ecf072f332 remove duplication for residuals. requires memset so we remain correct 2024-04-18 23:23:57 +03:00
Erik Schultheis fab549bacb removed further storage 2024-04-18 23:23:57 +03:00
Erik Schultheis 93db09577e first step towards reduced memory usage in backwards 2024-04-18 23:23:57 +03:00
Erik Schultheis 60090dcf2e update activation backward kernels that are not accumulated into the residual stream to overwrite instead of update. 2024-04-18 23:23:55 +03:00
Andrej Karpathy de1e87158c add comments to classifier_fused 2024-04-18 16:59:45 +00:00
Andrej 2d2f1df8f1 Merge pull request #150 from ademeure/optimised_fused_classifier
Optimised version of fused classifier + bugfixes(?)
2024-04-18 08:39:05 -07:00
Andrej Karpathy eef92c3fbe make comment on comparison to PyTorch runtime with kernels as they are right now 2024-04-18 04:53:09 +00:00
Andrej Karpathy 959ea543dc Merge branch 'master' of github.com:karpathy/llm.c 2024-04-18 04:40:09 +00:00
Andrej Karpathy 33a4985d19 change readme again to keep it up to date 2024-04-18 04:39:59 +00:00
Andrej aed262d773 Merge pull request #172 from FeSens/fix/validate_result_attention
fix(attention_forward.cu): Fix the number of elements to check
2024-04-17 21:28:31 -07:00
Andrej Karpathy ac183b3ba7 small tweaks for the test script. also i noticed that i had to bump 1e-2 -> 2e-2 for tolerance, i think we changed something very minor and tripped this threshold, but the tensors look ok 2024-04-18 04:27:13 +00:00
FeSens 0c09a37568 fix(attention_forward.cu): Fix the number of elements to check when verifying attention implementations.
This was causing me problems when I was trying to debug a new implementation on fewer tokens.
2024-04-18 01:23:05 -03:00
Andrej Karpathy fdd6608281 turn on the full training by default, as things are fast enough. we're doing exactly 1 epoch of finetuning on shakespeare 2024-04-18 04:12:04 +00:00
lancer a1abcbc32d fix the bug that yields cpu, gpu results mismatch 2024-04-17 21:04:45 -07:00
Andrej 20187634ec Merge pull request #167 from ngc92/bwd-att-coalesced
coleased memory reads for faster backward pass in attention
2024-04-17 20:50:33 -07:00
ademeure 557fecd6e9 bugfix for bounds checking 2024-04-18 04:45:07 +01:00
ademeure 2fd8ada4a6 Merge remote-tracking branch 'karpathy/master' into optimised_fused_classifier 2024-04-18 03:51:01 +01:00
ademeure ea77a26473 mini fix 2024-04-18 03:42:34 +01:00
ademeure 3ba2cdc723 fixes + bounds checking 2024-04-18 03:41:27 +01:00
lancer ac055f730c fix the bug that yields cpu, gpu results mismatch 2024-04-17 17:47:11 -07:00
Erik Schultheis 58a7b57f96 enable benchmarking 2024-04-18 01:39:47 +03:00
Erik Schultheis e278112c19 coalesced reading in attention backward pass 2024-04-18 01:39:43 +03:00
Andrej 4c27dd923b Merge pull request #164 from lancerts/const-fix-ln_back
Constness fix layernorm_backward.cu
2024-04-17 14:59:04 -07:00
ademeure 4ee98f5135 Merge branch 'master' into optimised_fused_classifier 2024-04-17 22:58:51 +01:00
lancer c327923df9 Constness fix layernorm_backward.cu 2024-04-17 14:56:03 -07:00
lancer 81aa596251 Constness fix layernorm_backward.cu 2024-04-17 14:54:49 -07:00
Andrej e5d519c62b Merge pull request #162 from ngc92/layernorm-bw
warp-wide reduction for layernorm backward
2024-04-17 13:34:36 -07:00
Erik Schultheis c1bfe4a639 warp-wide reduction for layernorm backward 2024-04-17 21:52:40 +03:00
Andrej Karpathy 79586880b7 adjust readme now that the full training loop works in cud 2024-04-17 17:15:59 +00:00
Andrej Karpathy fa4ebfa663 kernel2 for softmax backward 2024-04-17 08:14:22 +00:00
Andrej 2ecb4ec3ea Merge pull request #161 from ngc92/att-bw-fixup
fix accumulation and load balancing
2024-04-17 01:12:35 -07:00
Erik Schultheis 0c16c2ce57 fix accumulation and load balancing 2024-04-17 10:59:20 +03:00
Andrej Karpathy 954077fb88 TRAINING WORKSgit add train_gpt2.cu! ITS SLOW BUT IT WORKS WOOT 2024-04-17 04:58:37 +00:00
Andrej Karpathy d74d7ae652 small adamw fixes, needs more work 2024-04-17 04:58:16 +00:00
Andrej Karpathy 66a0b5242e Merge branch 'master' of github.com:karpathy/llm.c 2024-04-17 04:51:26 +00:00
Andrej bfa69805f6 Merge pull request #159 from nopperl/adam
AdamW CUDA kernel
2024-04-16 21:48:14 -07:00
Andrej Karpathy 8a673edcc3 push my attention_backward file 2024-04-17 04:42:42 +00:00
Andrej Karpathy 3f24ac961d check the entire gradient tensor all at once nice 2024-04-17 04:41:58 +00:00
Andrej Karpathy fd6882a1e3 encoder backward really naive, just to get this finished full 2024-04-17 04:39:02 +00:00
Andrej Karpathy a3a3b82d82 also backward the attention, that was not easy lol 2024-04-17 04:23:08 +00:00
nopperl 97e61e2258 add slightly more optimized adamw kernel 2024-04-17 03:31:07 +02:00
nopperl 5684ff7e77 add adamw cuda kernel 2024-04-17 03:30:59 +02:00
Andrej Karpathy 9e71af9678 Merge branch 'master' of github.com:karpathy/llm.c 2024-04-16 20:36:42 +00:00
Andrej Karpathy 45b0df0d3a backward now up to and not including attention 2024-04-16 20:34:13 +00:00
Andrej Karpathy e96f0ee2c8 backward the layernorm 2024-04-16 20:31:21 +00:00
Andrej Karpathy bf4ac57adc backward the gelu and the matmul after 2024-04-16 20:29:08 +00:00
Andrej Karpathy d753b1872e clean up checks a bit in test script 2024-04-16 20:16:26 +00:00
Andrej Karpathy 0da60460e9 backprop through the first residual and final matmul of the final residual block 2024-04-16 20:12:17 +00:00
Andrej 136875fa5c Merge pull request #149 from ent0n29/master
-fno-finite-math-only for almost 2x speed up, fix for [#19]
2024-04-16 12:42:21 -07:00
Andrej e15691ad16 Merge pull request #151 from msharmavikram/fix-out-of-resource
simple fix for bug [#106]
2024-04-16 12:35:52 -07:00
Andrej 862852fe9d Merge pull request #152 from ngc92/fused-fixes
ensure valid test cases and incorporate fixes by @ademeure
2024-04-16 12:31:54 -07:00
Erik Schultheis 43bf9eb66a ensure valid test cases and incorporate fixes by @ademeure 2024-04-16 20:53:50 +03:00
msharmavikram c62b50015f simple fix for bug https://github.com/karpathy/llm.c/issues/106 2024-04-16 10:23:48 -05:00
ademeure b225501b7b optimised version of the cool new fused classifier + hopefully bugfixes? 2024-04-16 14:54:23 +01:00
ent0n29 66c30fb55f -fno-finite-math-only for almost 2x speed up 2024-04-16 12:49:54 +02:00
Andrej Karpathy 67e0e6d6d4 diligently cudaCheck last error after every kernel launch nice 2024-04-16 06:21:35 +00:00
Andrej Karpathy 2517fe7492 exit(1) -> exit(EXIT_FAILURE), more portable and readable 2024-04-16 06:12:23 +00:00
Andrej Karpathy de0af15576 add mallocCheck as well. is this excessive? hmm. seem ok 2024-04-16 06:04:51 +00:00
Andrej Karpathy 062ba5a30e experiment with nice macros that do error handling, fopenCheck, freadCheck, fcloseCheck, to get rid of warning errors during compile. this functionality is yet another example of something we want to make common to both the .c file and the .cu file. we have to do that refactor soon. for now just keeping it in the .cu file 2024-04-16 05:40:37 +00:00
Andrej df010fe5b5 Merge pull request #145 from lancerts/matmul_backward
Fix the backward dbias precision issue in matmul_backward
2024-04-15 20:32:16 -07:00
lancer 6ff4ad646d fix the backward dbias precision issue 2024-04-15 19:58:16 -07:00
Andrej Karpathy 6be8080501 add naive backward for layernorm too. continue backward pass for the full model. adjust gpt2_testcu, we can finally check our first parameter gradients! technically this means that we could now train some of the parameters of the model, jut the final layernorm haha. but the kernels are really inefficient so far 2024-04-16 02:10:45 +00:00
Andrej Karpathy 3de507739c minor fixes to matmul backward 2024-04-16 01:23:41 +00:00
Andrej Karpathy 9560b1a92b add matmul backward kernel, problem is the bias backward, ironically. which i thought i could cublas at first 2024-04-16 00:40:11 +00:00
Andrej Karpathy 27ccf0b595 fix bugs, must use size_t instead of int, this was bad. and also init cpu_losses to NULL if we want to free it later 2024-04-15 22:20:44 +00:00
Andrej 3c6dafeb53 Merge pull request #117 from ngc92/fused-classifier
WIP: Fully fused classification layer
2024-04-15 12:13:48 -07:00
Andrej a44b0ce0f0 Merge pull request #142 from dagelf/master
Update Makefile with -march=native for 25%+ CPU speedup
2024-04-15 12:01:25 -07:00
Andrej Karpathy 480b638606 few more words on what i want the repo to be atm 2024-04-15 18:55:50 +00:00
dagelf 9aa4efdcbe Merge remote-tracking branch 'refs/remotes/origin/master' 2024-04-15 20:33:03 +02:00
dagelf ecde56c1c8 Add -march=native to README 2024-04-15 20:32:27 +02:00
Coenraad 7064e36472 more original typos fixed 2024-04-15 20:22:33 +02:00
dagelf c75183d318 typo 2024-04-15 20:04:55 +02:00
Coenraad Loubser 752cbe068d Update Makefile with -march=native
-march=native Results in a 30% speedup on all the platforms I've tried, for train_gpt2 (Admittedly, all older Intel and AMD ones)
```
model name      : Intel(R) Core(TM) i3-9100F CPU @ 3.60GHz
step 0: train loss 5.356172 (took 15197.876634 ms)
vs 
step 0: train loss 5.356185 (took 10418.548668 ms)

model name      : AMD Ryzen 5 3600 6-Core Processor
step 0: train loss 5.356173 (took 3599.359234 ms)
vs
step 0: train loss 5.356185 (took 2708.045790 ms)
```
2024-04-15 19:52:19 +02:00
Andrej Karpathy cc5931bfdb add clarification on notable forks to readme 2024-04-15 03:43:43 +00:00
Andrej Karpathy 82761ff29b confirm that linux and macos disagree, putting in the macos version as before 2024-04-15 02:38:23 +00:00
Andrej Karpathy bcb5228ab9 also update the .c file in addition to .cu with recent changes, e.g. the tokenizer and the B*T sampling, and update the README to match 2024-04-15 02:32:26 +00:00
Andrej Karpathy 3652832073 update the .c code, not just .cu code. oh copy paste, starting to bite us haha. also fflush when printing tokens 2024-04-15 02:13:49 +00:00
Andrej be8ace4be6 Merge pull request #128 from ngc92/test-updates
Test updates
2024-04-14 19:00:26 -07:00
Andrej Karpathy c165855cd3 add optional support for the tokenizer. people have to re-run train_gpt2.py to get this feature. later, we will change the code to demand this and error if it is not found, for now i just don't want to brick people who do a pull. also i only modified the train_gpt2.cu for now, not the .c. it might be time to separate out the dataloader, the tokenizer, and any other common utilities to their own files (?). otherwise it feels a bit silly to copy paste stuff around? not sure 2024-04-15 01:57:04 +00:00
Erik Schultheis c73fa8842a fixed numerics for online kernel, and an existing race condition in kernel 2 2024-04-15 03:41:09 +03:00
Andrej Karpathy f4cfd78245 only call gpt2_forward with the exact same sizing configuration B,T right now for safety. We can later bring back the <= B, <=T forward pass, but we need to do it carefully and have tests that make sure that a strictly smaller configuration produces the exact same results for that chunk of b,t. In other words we'd want ot make sure that the entire forward pass is range invariant. Currently it is not only because of the attention kernel. I think this is fixable, not too difficult, but it requires careful thought and associated tests for ensuring the range invariance. When those tests pass we can bring back old and more efficient behavior. For now it's just causing bugs, so I am putting in this highly defensive, but a lot more likely correct code 2024-04-14 23:56:52 +00:00
Erik Schultheis 8fb966bc31 kernel 1 should also accumulate in double to ensure stability 2024-04-15 02:52:38 +03:00
Erik Schultheis 60cf281e46 use V dimension in softmax tests
adjusted tests to be more challenging
2024-04-15 02:14:48 +03:00
Erik Schultheis 99b42c1a94 nicer printing 2024-04-15 02:03:48 +03:00
Erik Schultheis fe1d8cfb47 fix timing output 2024-04-15 02:03:42 +03:00
Andrej Karpathy 9807c041ee reshuffle the validation code a bit and veirfy all inputs/targets are in [0, V) 2024-04-14 19:30:56 +00:00
Andrej 312b043697 Merge pull request #122 from ngc92/fix-illegal-access
fix #114
2024-04-14 12:19:03 -07:00
Erik Schultheis 5c9d61fd35 Prevent out-of-bounds token-ids in input padding for test generation 2024-04-14 13:05:57 +03:00
Erik Schultheis 780d8bd012 initial version 2024-04-14 03:48:56 +03:00
Andrej Karpathy 534b93a732 erase the old kernel to save LOC and do some renaming and adjust dev/cuda code to add the new largeC kernel as kernel 7 2024-04-13 22:02:31 +00:00
Andrej 0554c7e414 Merge pull request #115 from ademeure/softmax_largeC
New softmax for non-self-attention case optimised for large Cs + add more __ldcs
2024-04-13 14:49:26 -07:00
ademeure fc3dce37ce Add __ldcs to a few reads for which the memory will never be read again in the forward pass 2024-04-13 22:38:12 +01:00
ademeure 321b042657 New softmax for non-self-attention softmax case with large Cs 2024-04-13 22:31:23 +01:00
Andrej Karpathy c02bae2ebc average iterates in code 2024-04-13 21:09:37 +00:00
Andrej Karpathy c70fbd059b refine comment on time 2024-04-13 21:09:25 +00:00
Andrej Karpathy ca242a3d0b get rid of spurious handle variable creation, remnant from before, and update README 2024-04-13 21:04:49 +00:00
Andrej c65d1d305b Merge pull request #113 from ademeure/cublasTF32
Fix cuBLAS to use TF32 with pre-initialised handle
2024-04-13 13:56:13 -07:00
ademeure 5ab44973b0 Fix cuBLAS calls to use pre-initialised handle with TF32 enabled (big perf gain) 2024-04-13 20:35:53 +01:00
Andrej Karpathy dd3dc179a4 git push origin masterMerge branch 'ngc92-common' 2024-04-13 17:57:33 +00:00
Andrej Karpathy c5ee371564 make the tests own their cublas cublaslt handles and settings, instead of putting that stuff in common 2024-04-13 17:56:18 +00:00
Andrej Karpathy 82623518a4 fix comment block of gelu forward kernel 2024-04-13 17:48:53 +00:00
Andrej Karpathy 02b5b33c5f remove unused variable C 2024-04-13 17:48:11 +00:00
Andrej Karpathy 84145411d3 fix crossentropy forward, loss requires probabilities so these floats must be 0..1 instead of -1..1 or we get nans 2024-04-13 17:47:19 +00:00
Andrej Karpathy 69f0385237 add comments and handling of kernels and their correctness 2024-04-13 17:44:31 +00:00
Andrej Karpathy 035b9e76fc Merge branch 'common' of https://github.com/ngc92/llm.c into ngc92-common 2024-04-13 17:25:06 +00:00
Andrej Karpathy 552b972e52 Merge branch 'zocterminal-warnings' 2024-04-13 16:22:10 +00:00
Andrej Karpathy f347957edf change comment style 2024-04-13 16:21:57 +00:00
Erik Schultheis f459bfa833 fixed re-initialization of device memory after test 2024-04-13 18:23:12 +03:00
Erik Schultheis d035725d22 mark problematic tests 2024-04-13 18:09:28 +03:00
Erik Schultheis 59bada94d9 improved testing 2024-04-13 17:06:52 +03:00
Markus Schmidt ae63f6d008 warnings in train_gpt2.cu 2024-04-13 12:52:54 +00:00
Erik Schultheis 7e7460cfbd const'ness fixes 2024-04-13 13:38:11 +03:00
Erik Schultheis 7e8e038de8 extracted common testing functions into their own header 2024-04-13 13:00:05 +03:00
Andrej Karpathy 6b49ed1c0b incorporate the fastest kernels and adjust readme 2024-04-13 02:58:23 +00:00
Andrej Karpathy 4f75e64573 Merge branch 'ngc92-attention-softmax-2' 2024-04-13 02:39:10 +00:00
Andrej Karpathy 6f9e28a1b6 few comments 2024-04-13 02:38:57 +00:00
Andrej Karpathy 2f99b3e608 Merge branch 'attention-softmax' of https://github.com/ngc92/llm.c into ngc92-attention-softmax-2 2024-04-13 02:34:06 +00:00
Andrej Karpathy bd743f806a Merge branch 'lancerts-online_softmax' 2024-04-13 02:11:50 +00:00
Andrej Karpathy 40b78be722 small changes to comments, and very minor code modifications and movements 2024-04-13 02:11:36 +00:00
Andrej Karpathy 6250336c09 Merge branch 'online_softmax' of https://github.com/lancerts/llm.c into lancerts-online_softmax 2024-04-13 01:54:16 +00:00
Andrej Karpathy a18e517c1a alright let's start the backward pass. crossentropy_softmax_backward first, ez one 2024-04-13 01:52:55 +00:00
Andrej Karpathy 492a4b151a oops we missed a comment here, remove as we got rid of the inline mallocs now, so we're all ok 2024-04-13 01:21:31 +00:00
Andrej e9f095208a Merge pull request #100 from ademeure/cudahostmalloc
Use cudaHostMalloc for inputs/targets and cpu_losses
2024-04-12 18:19:38 -07:00
ademeure 2616a6461d Use cudaHostMalloc for inputs/targets and cpu_losses (+moved its malloc/free away from main loop). 2024-04-13 01:31:16 +01:00
Andrej Karpathy 3ab2039e32 update timings 2024-04-12 23:47:15 +00:00
Andrej Karpathy 59b9e8b96a fix handles in cublas in attention forward as well 2024-04-12 23:15:47 +00:00
Andrej 8822e7803e use cublaslt, which fuses bias, and also use tf32 when your GPU supports it (#98)
Fusing the bias into matmul and moving to tf32 speed up the code *a lot*, we're now down to 59ms/iter compared to 25ms/iter for pytorch compiled tf32. i.e. still off by 2.36X.
2024-04-12 15:32:07 -07:00
Erik Schultheis e048825f5a div_ceil 2024-04-13 01:20:09 +03:00
Erik Schultheis 8e87aac341 online softmax 2024-04-12 12:47:48 +03:00
Erik Schultheis de98a9a44a vector loads 2024-04-12 12:47:43 +03:00
Andrej Karpathy 6396e393e3 add links to discussions 2024-04-12 01:11:29 +00:00
Erik Schultheis baba3cfa01 recalculate instead of memory round-trip 2024-04-12 02:20:20 +03:00
lancer 69d0e5bac7 Include a parallel kernal of softmax_online 2024-04-11 15:30:36 -07:00
lancer ca1e5ee04d Include a parallel kernal of softmax_online 2024-04-11 15:24:41 -07:00
Andrej Karpathy dec54904b4 one more simple clarifying comment 2024-04-11 22:24:38 +00:00
Andrej Karpathy d680fbad25 one more comment on attention 2024-04-11 21:49:14 +00:00
Andrej Karpathy 9710163a60 did one pass of additional comments, potentially helpful, as i think a lot of people may be looking at this for the first time 2024-04-11 21:45:46 +00:00
Erik Schultheis 5f99fb30d2 cooperative groups and fused scale kernel 2024-04-11 23:06:46 +03:00
lancer d6bcc129b9 Improve the efficiency 2024-04-11 12:16:20 -07:00
lancer 00010c7188 Improve the efficiency 2024-04-11 12:13:34 -07:00
Andrej Karpathy 2c81198101 make the new cooperative groups layernorm kernel the default. shaves off aonly bout 1ms of the total running time though 2024-04-11 18:23:46 +00:00
Andrej Karpathy 6f5ec063f0 Merge branch 'ngc92-layer-norm' 2024-04-11 18:22:40 +00:00
Andrej Karpathy 00b9fda4a1 ceildiv in layernorm 2024-04-11 18:21:34 +00:00
Andrej Karpathy 089cd5f1dd make sure we also test correctness of mean and rstd and do it all block sizes 2024-04-11 18:12:26 +00:00
Andrej Karpathy f28613e0b9 small edits to comments and variables 2024-04-11 18:09:54 +00:00
Andrej Karpathy 2553518c97 Merge branch 'layer-norm' of https://github.com/ngc92/llm.c into ngc92-layer-norm 2024-04-11 18:07:00 +00:00
Erik Schultheis 4e91362fa2 optionally generate mean and rstd 2024-04-11 20:48:14 +03:00
lancer f524a90de3 Include the online softmax CPU code (from the paper "Online normalizer calculation for softmax") and native port to GPU kernel.
Both the CPU and GPU kernels' correctness are verified by executing the script.

Performance: the online kernel (native CPU port) is faster than the kernel1 (native CPU port) but worse than kernel 2-4.
2024-04-11 09:47:50 -07:00
Andrej 79b8af4a9b Merge pull request #56 from scotthaleen/omp-osx-intel
Detect OpenMP support - macOS Intel
2024-04-11 09:13:10 -07:00
Andrej 21a250d372 Merge pull request #64 from krrishnarraj/master
[train_gpt2.py] synchronize based on device
2024-04-11 09:12:07 -07:00
Andrej Karpathy 1751af1607 inline docs for sizes 2024-04-11 16:04:14 +00:00
Andrej Karpathy bcd0dad6c8 Merge branch 'master' of github.com:karpathy/llm.c 2024-04-11 15:50:10 +00:00
Andrej Karpathy f26cf00a61 make gelu constant be a #define, speeds up the kernel by 1% or so 2024-04-11 15:49:01 +00:00
Andrej f46d5b2e2a Merge pull request #72 from Soldy/master
-O3 cannot go with -Ofast
2024-04-11 08:38:57 -07:00
soldy 3aaed922a3 -O3 cannot go with -Ofast
Signed-off-by: soldy <4786022+soldy@users.noreply.github.com>
2024-04-11 16:27:59 +01:00
Krishnaraj Bhat 4542f8913a [train_gpt2.py] synchronize based on device 2024-04-11 17:06:51 +05:30
Scott Haleen 47b2f9312c fixed homebrew path for intel based macs 2024-04-10 23:01:07 -04:00
Erik Schultheis afddb1af5b fixed block size and improved load op hint 2024-04-11 03:47:38 +03:00
Erik Schultheis 466723bb06 fully fused layer-norm kernel 2024-04-11 02:01:41 +03:00
Andrej a08c11b60e Merge pull request #48 from lancerts/fix_matmul
fix error in small typos in matmul_forward.cu
2024-04-10 13:44:22 -07:00
lancer 68d2ce0601 fix error in small typos in matmul_forward.cu 2024-04-10 13:35:20 -07:00
Andrej Karpathy ef7a4867b1 Merge branch 'master' of github.com:karpathy/llm.c 2024-04-10 19:25:13 +00:00
Andrej Karpathy e1e0a08035 change fabs to fabsf 2024-04-10 19:23:18 +00:00
Andrej 919b2200a0 Merge pull request #20 from varunlakkur/master
Fix a typo
2024-04-10 12:21:07 -07:00
Andrej 9f5069d3b2 Merge pull request #38 from lancerts/fix_matmul
fix the consistency of the transpose notation in matmul_foward.cu
2024-04-10 12:15:53 -07:00
Andrej c0d57b1c45 Merge pull request #47 from eltociear/patch-1
update layernorm.md
2024-04-10 12:14:56 -07:00
Andrej Karpathy add9e77daa mention compilation issue that seems to be common. we may want to resolve this inside the Makefile though 2024-04-10 19:11:44 +00:00
Andrej Karpathy fd8459d858 small gitignore for now. probably iterate on 2024-04-10 18:54:23 +00:00
Andrej 15fad9bc96 Create LICENSE MIT 2024-04-10 11:47:40 -07:00
Andrej Karpathy 80f52e5290 the full forward pass of GPT-2 in one file of pure CUDA 2024-04-10 18:15:55 +00:00
Ikko Eltociear Ashimine ca137b02d5 update layernorm.md
burried -> buried
2024-04-11 00:07:22 +09:00
Andrej Karpathy d8e2a36f15 ok crossentropy is fairly ez, just have to be careful with logs of negative numbers in debug tests. so this should be all the layers now 2024-04-10 01:13:20 +00:00
Andrej Karpathy 46ee0a3944 sorry there should be no mallocs inside these functions, have to pass in buffer memory from outside. getting tired 2024-04-10 00:55:14 +00:00
Andrej Karpathy 919ae1f27b ok attention is now optimized too, by 20X, nice. actually this one was really painful, i had to make sure all the cuBLAS calls take all the right arguments and deal with the fact that cuBLAS is column major instead of row major. i go for nap now 2024-04-10 00:50:30 +00:00
lancer fbd679adec fix the consistency of the transpose notation 2024-04-09 15:19:55 -07:00
Andrej Karpathy 8386e5393c intra-warp reductions and inter-warp shared memory reductions for B,T,C parallelism of softmax kernel nice 2024-04-09 20:08:11 +00:00
Andrej Karpathy c29d70a9b3 add softmax forward, both a naive parallel over B,T version, and a more optimized parallel over B,T,C version using shared memory for reductions 2024-04-09 19:22:47 +00:00
Andrej Karpathy 327705d03f mention ongoing cuda work in main readme 2024-04-09 18:29:24 +00:00
Andrej ea1fad0337 Merge pull request #33 from lancerts/fix_M_PI_undefined
fix a potential error: identifier M_PI is undefined in the gelu kernal
2024-04-09 11:25:23 -07:00
Andrej e3e0d73fbc Merge pull request #34 from VinciGit00/patch-1
Free the memory in layernorm.c
2024-04-09 11:24:32 -07:00
Andrej Karpathy bb4f04bdf8 add flash attention, but a naive implementation there of 2024-04-09 18:22:49 +00:00
Marco Vinciguerra 8397fb1512 Free the memory in layernorm.c 2024-04-09 20:14:18 +02:00
lancer 3a4d783633 amend 2024-04-09 09:59:03 -07:00
lancer 73c92912b0 fix error: identifier M_PI is undefined 2024-04-09 09:51:28 -07:00
Andrej Karpathy 03f37cfb6b naive attention kernel only parallelizing over batch,time,heads. have to speed this up a lot 2024-04-09 14:53:18 +00:00
Andrej Karpathy 41d5f56278 might as well push the few kernels that i feel ok about so far 2024-04-09 14:34:13 +00:00
Andrej 4d98c44904 Merge pull request #23 from DominguesAddem1974/fix/torch-warning
fix: torch warning of python demo
2024-04-09 05:45:18 -07:00
DominguesAddem1974 daf1dd81b2 fix: torch warning of python demo 2024-04-09 18:17:04 +08:00
Varun 69c1f179a8 Fix a typo 2024-04-09 15:02:33 +05:30
Andrej a22c22b1ee Merge pull request #14 from Ricardicus/master
Using the compiler at hand
2024-04-08 21:29:47 -07:00
Andrej 5d9ed0cc70 Merge pull request #15 from lancerts/layer_norm_code_ptr
Include the pytorch layer_norm.cpp and layer_norm_kernel.cu code pointer in readme
2024-04-08 21:28:04 -07:00
karpathy f6b93912b7 minor clarifications 2024-04-08 21:24:34 -07:00
Rickard Hallerbäck 1df7bceff9 using the compiler at hand 2024-04-09 06:22:03 +02:00
karpathy 0d30f54d9a more notes on training/inference 2024-04-08 21:20:28 -07:00
lancer eecb234698 Include the pytorch layer_norm.cpp and layer_norm_kernel.cu code pointer for people who are interested in details 2024-04-08 21:07:28 -07:00
karpathy e4d251bad4 add comment on why we even have to implement layernorm forward pass in pytorch 2024-04-08 17:03:40 -07:00
karpathy fbc879a83c fix hallucination of github copilot lol 2024-04-08 16:48:48 -07:00
karpathy ce3fef00fb add clarification on checkpointing 2024-04-08 16:46:25 -07:00
karpathy a99b86f2ef oops fix latex expression 2024-04-08 16:30:02 -07:00
karpathy c39de59168 attach a simple tutorial 2024-04-08 16:26:38 -07:00
karpathy 3bcb9ba7d2 remove unused variables 2024-04-08 15:10:51 -07:00
Andrej b0fc807bd7 Merge pull request #2 from onuralpszr/py-tqdm
fix: add tqdm to requirements.txt
2024-04-08 14:42:56 -07:00
Andrej 90ac963826 Merge branch 'master' into py-tqdm 2024-04-08 14:42:47 -07:00
Andrej 40221ccff5 Merge pull request #4 from Ricardicus/master
Minor correction for openmp check
2024-04-08 14:41:43 -07:00
Andrej 0c2d8d771d Merge pull request #5 from alexziskind1/master
fix: add transformers to requirements.txt
2024-04-08 14:39:35 -07:00
Alex Ziskind b8f7b1c201 fix: add transformers to requirements.txt 2024-04-08 16:50:59 -04:00
Rickard Hallerbäck b2228d0e83 Minor correction for openmp check 2024-04-08 22:33:47 +02:00
Onuralp SEZER 8465847284 fix: add tqdm to requirements.txt 2024-04-08 23:26:14 +03:00
karpathy 2790d263a6 oops fix readme bug 2024-04-08 12:42:44 -07:00
karpathy e8e1628632 first commit of just the reference cpu fp32 gpt2 training 2024-04-08 12:41:21 -07:00
6 changed files with 148 additions and 397 deletions
-267
View File
@@ -1,267 +0,0 @@
# llm.c
LLMs in simple, pure C/CUDA with no need for 245MB of PyTorch or 107MB of cPython. Current focus is on pretraining, in particular reproducing the [GPT-2](https://github.com/openai/gpt-2) and [GPT-3](https://arxiv.org/abs/2005.14165) miniseries, along with a parallel PyTorch reference implementation in [train_gpt2.py](train_gpt2.py). You'll recognize this file as a slightly tweaked [nanoGPT](https://github.com/karpathy/nanoGPT), an earlier project of mine. Currently, llm.c is a bit faster than PyTorch Nightly (by about 7%). In addition to the bleeding edge mainline code in [train_gpt2.cu](train_gpt2.cu), we have a simple reference CPU fp32 implementation in ~1,000 lines of clean code in one file [train_gpt2.c](train_gpt2.c). I'd like this repo to only maintain C and CUDA code. Ports to other languages or repos are very welcome, but should be done in separate repos, and I am happy to link to them below in the "notable forks" section. Developer coordination happens in the [Discussions](https://github.com/karpathy/llm.c/discussions) and on Discord, either the `#llmc` channel on the [Zero to Hero](https://discord.gg/3zy8kqD9Cp) channel, or on `#llmdotc` on [GPU MODE](https://discord.gg/gpumode) Discord.
## quick start
The best introduction to the llm.c repo today is reproducing the GPT-2 (124M) model. [Discussion #481](https://github.com/karpathy/llm.c/discussions/481) steps through this in detail. We can reproduce other models from the GPT-2 and GPT-3 series in both llm.c and in the parallel implementation of PyTorch. Have a look at the [scripts README](scripts/README.md).
debugging tip: when you run the `make` command to build the binary, modify it by replacing `-O3` with `-g` so you can step through the code in your favorite IDE (e.g. vscode).
## quick start (1 GPU, fp32 only)
If you won't be training on multiple nodes, aren't interested in mixed precision, and are interested in learning CUDA, the fp32 (legacy) files might be of interest to you. These are files that were "checkpointed" early in the history of llm.c and frozen in time. They are simpler, more portable, and possibly easier to understand. Run the 1 GPU, fp32 code like this:
```bash
chmod u+x ./dev/download_starter_pack.sh
./dev/download_starter_pack.sh
make train_gpt2fp32cu
./train_gpt2fp32cu
```
The download_starter_pack.sh script is a quick & easy way to get started and it downloads a bunch of .bin files that help get you off the ground. These contain: 1) the GPT-2 124M model saved in fp32, in bfloat16, 2) a "debug state" used in unit testing (a small batch of data, and target activations and gradients), 3) the GPT-2 tokenizer, and 3) the tokenized [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) dataset. Alternatively, instead of running the .sh script, you can re-create these artifacts manually as follows:
```bash
pip install -r requirements.txt
python dev/data/tinyshakespeare.py
python train_gpt2.py
```
## quick start (CPU)
The "I am so GPU poor that I don't even have one GPU" section. You can still enjoy seeing llm.c train! But you won't go too far. Just like the fp32 version above, the CPU version is an even earlier checkpoint in the history of llm.c, back when it was just a simple reference implementation in C. For example, instead of training from scratch, you can finetune a GPT-2 small (124M) to output Shakespeare-like text, as an example:
```bash
chmod u+x ./dev/download_starter_pack.sh
./dev/download_starter_pack.sh
make train_gpt2
OMP_NUM_THREADS=8 ./train_gpt2
```
If you'd prefer to avoid running the starter pack script, then as mentioned in the previous section you can reproduce the exact same .bin files and artifacts by running `python dev/data/tinyshakespeare.py` and then `python train_gpt2.py`.
The above lines (1) download an already tokenized [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) dataset and download the GPT-2 (124M) weights, (3) init from them in C and train for 40 steps on tineshakespeare with AdamW (using batch size 4, context length only 64), evaluate validation loss, and sample some text. Honestly, unless you have a beefy CPU (and can crank up the number of OMP threads in the launch command), you're not going to get that far on CPU training LLMs, but it might be a good demo/reference. The output looks like this on my MacBook Pro (Apple Silicon M3 Max):
```
[GPT-2]
max_seq_len: 1024
vocab_size: 50257
num_layers: 12
num_heads: 12
channels: 768
num_parameters: 124439808
train dataset num_batches: 1192
val dataset num_batches: 128
num_activations: 73323776
val loss 5.252026
step 0: train loss 5.356189 (took 1452.121000 ms)
step 1: train loss 4.301069 (took 1288.673000 ms)
step 2: train loss 4.623322 (took 1369.394000 ms)
step 3: train loss 4.600470 (took 1290.761000 ms)
... (trunctated) ...
step 39: train loss 3.970751 (took 1323.779000 ms)
val loss 4.107781
generating:
---
Come Running Away,
Greater conquer
With the Imperial blood
the heaviest host of the gods
into this wondrous world beyond.
I will not back thee, for how sweet after birth
Netflix against repounder,
will not
flourish against the earlocks of
Allay
---
```
## datasets
The data files inside `/dev/data/(dataset).py` are responsible for downloading, tokenizing and saving the tokens to .bin files, readable easily from C. So for example when you run:
```bash
python dev/data/tinyshakespeare.py
```
We download and tokenize the [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) dataset. The output of this looks like this:
```
writing 32,768 tokens to ./dev/data/tinyshakespeare/tiny_shakespeare_val.bin
writing 305,260 tokens to ./dev/data/tinyshakespeare/tiny_shakespeare_train.bin
```
The .bin files contain a short header (1024 bytes) and then a stream of tokens in uint16, indicating the token ids with the GPT-2 tokenizer. More datasets are available in `/dev/data`.
## test
I am also attaching a simple unit test for making sure our C code agrees with the PyTorch code. On the CPU as an example, compile and run with:
```bash
make test_gpt2
./test_gpt2
```
This now loads the `gpt2_124M_debug_state.bin` file that gets written by train_gpt2.py, runs a forward pass, compares the logits and loss with the PyTorch reference implementation, then it does 10 iterations of training with Adam and makes sure the losses match PyTorch. To test the GPU version we run:
```bash
# fp32 test (cudnn not supported)
make test_gpt2cu PRECISION=FP32 && ./test_gpt2cu
# mixed precision cudnn test
make test_gpt2cu USE_CUDNN=1 && ./test_gpt2cu
```
This tests both the fp32 path and the mixed precision path. The test should pass and print `overall okay: 1`.
## tutorial
I attached a very small tutorial here, in [doc/layernorm/layernorm.md](doc/layernorm/layernorm.md). It's a simple, step-by-step guide to implementing a single layer of the GPT-2 model, the layernorm layer. This is a good starting point to understand how the layers are implemented in C.
**flash attention**. As of May 1, 2024 we use the Flash Attention from cuDNN. Because cuDNN bloats the compile time from a few seconds to ~minute and this code path is right now very new, this is disabled by default. You can enable it by compiling like this:
```bash
make train_gpt2cu USE_CUDNN=1
```
This will try to compile with cudnn and run it. You have to have cuDNN installed on your system. The [cuDNN installation instructions](https://developer.nvidia.com/cudnn) with apt-get will grab the default set of cuDNN packages. For a minimal setup, the cuDNN dev package is sufficient, e.g. on Ubuntu 22.04 for CUDA 12.x:
```bash
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt-get update
sudo apt-get -y install libcudnn9-dev-cuda-12
```
On top of this you need the [cuDNN frontend](https://github.com/NVIDIA/cudnn-frontend/tree/main), but this is just header files. Simply clone the repo to your disk. The Makefile currently looks for it in either your home directory or the current directory. If you have put it elsewhere, add `CUDNN_FRONTEND_PATH=/path/to/your/cudnn-frontend/include` to the `make` command-line.
## multi-GPU training
Make sure you install MPI and NCCL, e.g. on Linux:
```bash
sudo apt install openmpi-bin openmpi-doc libopenmpi-dev
```
For NCCL follow the instructions from the [official website](https://developer.nvidia.com/nccl/nccl-download) (e.g. network installer)
and then:
```bash
make train_gpt2cu
mpirun -np <number of GPUs> ./train_gpt2cu
```
or simply run one of our scripts under `./scripts/`.
## multi-node training
Make sure you've installed `NCCL` following instructions from [multi-GPU](#multi-gpu-training) section.
There are 3 ways we currently support that allow you to run multi-node training:
1) Use OpenMPI to exchange nccl id and initialize NCCL. See e.g. `./scripts/multi_node/run_gpt2_124M_mpi.sh` script for details.
2) Use shared file system to init NCCL. See `./scripts/multi_node/run_gpt2_124M_fs.sbatch` script for details.
3) Use TCP sockets to init NCCL. See `./scripts/multi_node/run_gpt2_124M_tcp.sbatch` script for details.
Note:
* If you're running in a slurm environment and your slurm doesn't support PMIx (which we assume will be a common situation given that `slurm-wlm` dropped PMIx support) you will have to use FS (2) or TCP (3) approach. To test whether your slurm supports PMIx run: `srun --mpi=list` and see whether you get `pmix` in the output.
* If you don't have slurm set up, you can kick off a multi-node run using `mpirun` - MPI (1).
None of these 3 methods is superior, we just offer you options so that you can run in your specific environment.
## experiments / sweeps
Just as an example process to sweep learning rates on a machine with 4 GPUs on TinyStories. Run a shell script `sweep.sh` (after you of course `chmod u+x sweep.sh`):
```bash
#!/bin/bash
learning_rates=(3e-5 1e-4 3e-4 1e-3)
for i in {0..3}; do
export CUDA_VISIBLE_DEVICES=$i
screen -dmS "tr$i" bash -c "./train_gpt2cu -i data/TinyStories -v 250 -s 250 -g 144 -l ${learning_rates[$i]} -o stories$i.log"
done
# you can bring these down with
# screen -ls | grep -E "tr[0-3]" | cut -d. -f1 | xargs -I {} screen -X -S {} quit
```
This example opens up 4 screen sessions and runs the four commands with different LRs. This writes the log files `stories$i.log` with all the losses, which you can plot as you wish in Python. A quick example of how to parse and plot these logfiles is in [dev/vislog.ipynb](dev/vislog.ipynb).
## repo
A few more words on what I want this repo to be:
First, I want `llm.c` to be a place for education. E.g. our `dev/cuda` folder is a place for a library of kernels for all the layers that are manually hand-written and very well documented, starting from very simple kernels all the way to more complex / faster kernels. If you have a new kernel with various different tradeoffs, please feel free to contribute it here.
That said, I also want `llm.c` to be very fast too, even practically useful to train networks. E.g. to start, we should be able to reproduce the big GPT-2 (1.6B) training run. This requires that we incorporate whatever fastest kernels there are, including the use of libraries such as cuBLAS, cuBLASLt, CUTLASS, cuDNN, etc. I also think doing so serves an educational purpose to establish an expert upper bound, and a unit of measurement, e.g. you could say that your manually written kernels are 80% of cuBLAS speed, etc. Then you can choose to do a super fast run, or you can choose to "drag and drop" whatever manual kernels you wish to use, and run with those.
However, as a constraint, I want to keep the mainline `llm.c` in the root folder simple and readable. If there is a PR that e.g. improves performance by 2% but it "costs" 500 lines of complex C code, and maybe an exotic 3rd party dependency, I may reject the PR because the complexity is not worth it. As a concrete example - making cuBLAS for matmuls the default in the root training loop is a no-brainer: it makes the mainline code much faster, it is a single line of interpretable code, and it is a very common dependency. On the side of this, we can have manual implementations that can compete with cuBLAS in `dev/cuda`.
Lastly, I will be a lot more sensitive to complexity in the root folder of the project, which contains the main / default files of the project. In comparison, the `dev/` folder is a bit more of a scratch space for us to develop a library of kernels or classes and share useful or related or educational code, and some of this code could be ok to be (locally) complex.
## notable forks
- AMD support
- [llm.c](https://github.com/anthonix/llm.c) by @[anthonix](https://github.com/anthonix): support for AMD devices, such as the 7900 XTX
- C#
- [llm.cs](https://github.com/azret/llm.cs) by @[azret](https://github.com/azret): a C# port of this project
- [Llm.cs](https://github.com/nietras/Llm.cs) by @[nietras](https://github.com/nietras): a C# port of this project with focus on easy to get started on any platform. Clone and run ✅
- 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 [GPU MODE Discord Server](https://discord.gg/cudamode)
- C++/CUDA
- [llm.cpp](https://github.com/zhangpiu/llm.cpp/tree/master/llmcpp) by @[zhangpiu](https://github.com/zhangpiu): a port of this project using the [Eigen](https://gitlab.com/libeigen/eigen), supporting CPU/CUDA.
- WebGPU C++
- [gpu.cpp](https://github.com/AnswerDotAI/gpu.cpp) by @[austinvhuang](https://github.com/austinvhuang): a library for portable GPU compute in C++ using native WebGPU. Aims to be a general-purpose library, but also porting llm.c kernels to WGSL.
- C++
- [llm.cpp](https://github.com/GaoYusong/llm.cpp) by @[GaoYusong](https://github.com/GaoYusong): a port of this project featuring a C++ single-header [tinytorch.hpp](https://github.com/GaoYusong/llm.cpp/blob/main/tinytorch.hpp) library
- Go
- [llm.go](https://github.com/joshcarp/llm.go) by @[joshcarp](https://github.com/joshcarp): a Go port of this project
- Java
- [llm.java](https://github.com/harryjackson/llm.java) by @[harryjackson](https://github.com/harryjackson): a Java port of this project
- Metal
- [llm.metal](https://github.com/regrettable-username/llm.metal) by @[regrettable-username](https://github.com/regrettable-username): LLM training in simple, raw C/Metal Shading Language
- Mojo
- [llm.🔥](https://github.com/dorjeduck/llm.mojo) by @[dorjeduck](https://github.com/dorjeduck): a Mojo port of this project
- OpenCL
- [llm.c](https://github.com/krrishnarraj/llm.c) by @[krrishnarraj](https://github.com/krrishnarraj): an OpenCL port of this project
- Rust
- [llm.rs](https://github.com/yijunyu/llm.rs) by @[Yijun Yu](https://github.com/yijunyu): a Rust rewrite with the aim to have same performance
- [llm.rs](https://github.com/ToJen/llm.rs) by @[ToJen](https://github.com/ToJen): a Rust port of this project
- Swift
- [llm.swift](https://github.com/otabuzzman/llm.swift) by @[otabuzzman](https://github.com/otabuzzman): a Swift port of this project
- Zig
- [llm.zig](https://github.com/Saimirbaci/llm.zig) by @[saimirbaci](https://github.com/Saimirbaci): a Zig port of this project
- Habana Gaudi2
- [llm.tpc](https://github.com/abhilash1910/llm.tpc) by @[abhilash1910](https://github.com/abhilash1910): a Habana Gaudi2 port of this project
- Nim
- [llm.nim](https://github.com/Vindaar/llm.nim) by @[Vindaar](https://github.com/Vindaar): a Nim port of this project
## discussions
Ways of organizing development:
- Experiencing a concrete issue with the repo? Use [Issues](https://github.com/karpathy/llm.c/issues).
- Have some code to contribute? Open a [PR](https://github.com/karpathy/llm.c/pulls)
- Chat about the repo, ask questions, etc.? Look at [Discussions](https://github.com/karpathy/llm.c/discussions).
- Something faster? I created a new `#llmc` channel on my [Zero to Hero Discord channel](https://discord.gg/3zy8kqD9Cp).
## license
MIT
+68 -80
View File
@@ -1,22 +1,16 @@
<!-- WEHUB_ZH_README -->
> [!NOTE]
> 本文档由 WeHub 基于上游 README 翻译整理,属于社区翻译,非官方中文文档。
> [English](./README.en.md) · [原始项目](https://github.com/karpathy/llm.c) · [上游 README](https://github.com/karpathy/llm.c/blob/HEAD/README.md)
> 原作者、版权与许可证归属以原始项目及本仓库 LICENSE 文件为准。
# llm.c
用简洁、纯粹的 C/CUDA 实现 LLM,无需 245MB PyTorch 107MB cPython。当前重点是预训练(pretraining),尤其是复现 [GPT-2](https://github.com/openai/gpt-2) [GPT-3](https://arxiv.org/abs/2005.14165) 迷你系列,同时在 [train_gpt2.py](train_gpt2.py) 中提供并行的 PyTorch 参考实现。你会认出该文件是我早期项目 [nanoGPT](https://github.com/karpathy/nanoGPT), 的稍加改动版本。目前,llm.c 比 PyTorch Nightly 略快(约 7%)。除了 [train_gpt2.cu](train_gpt2.cu) 中的前沿主线代码外,我们还在单文件 [train_gpt2.c](train_gpt2.c) 中提供了一个约 1,000 行整洁代码的简单 CPU fp32 参考实现。我希望本仓库仅维护 C 和 CUDA 代码。欢迎移植到其他语言或仓库,但应在独立仓库中进行,我很乐意在下方「notable forks」部分链接它们。开发者协作在 [Discussions](https://github.com/karpathy/llm.c/discussions) 以及 Discord 上进行,可在 [Zero to Hero](https://discord.gg/3zy8kqD9Cp) 频道的 `#llmc` 频道,或在 [GPU MODE](https://discord.gg/gpumode) Discord 的 `#llmdotc` 上交流。
LLMs in simple, pure C/CUDA with no need for 245MB of PyTorch or 107MB of cPython. Current focus is on pretraining, in particular reproducing the [GPT-2](https://github.com/openai/gpt-2) and [GPT-3](https://arxiv.org/abs/2005.14165) miniseries, along with a parallel PyTorch reference implementation in [train_gpt2.py](train_gpt2.py). You'll recognize this file as a slightly tweaked [nanoGPT](https://github.com/karpathy/nanoGPT), an earlier project of mine. Currently, llm.c is a bit faster than PyTorch Nightly (by about 7%). In addition to the bleeding edge mainline code in [train_gpt2.cu](train_gpt2.cu), we have a simple reference CPU fp32 implementation in ~1,000 lines of clean code in one file [train_gpt2.c](train_gpt2.c). I'd like this repo to only maintain C and CUDA code. Ports to other languages or repos are very welcome, but should be done in separate repos, and I am happy to link to them below in the "notable forks" section. Developer coordination happens in the [Discussions](https://github.com/karpathy/llm.c/discussions) and on Discord, either the `#llmc` channel on the [Zero to Hero](https://discord.gg/3zy8kqD9Cp) channel, or on `#llmdotc` on CUDA MODE Discord.
## quick start
当前入门 llm.c 仓库的最佳方式是复现 GPT-2 (124M) 模型。[Discussion #481](https://github.com/karpathy/llm.c/discussions/481) 详细介绍了这一过程。我们可以在 llm.c 及并行的 PyTorch 实现中复现 GPT-2 GPT-3 系列的其他模型。请查看 [scripts README](scripts/README.md)
The best introduction to the llm.c repo today is reproducing the GPT-2 (124M) model. [Discussion #481](https://github.com/karpathy/llm.c/discussions/481) steps through this in detail. We can reproduce other models from the GPT-2 and GPT-3 series in both llm.c and in the parallel implementation of PyTorch. Have a look at the [scripts README](scripts/README.md).
debugging tip: 运行 `make` 命令构建二进制文件时,将其中的 `-O3` 替换为 `-g`,即可在你喜爱的 IDE(例如 vscode)中单步调试代码。
debugging tip: when you run the `make` command to build the binary, modify it by replacing `-O3` with `-g` so you can step through the code in your favorite IDE (e.g. vscode).
## quick start (1 GPU, fp32 only)
如果你不会在多节点上训练、对混合精度不感兴趣,并且想学习 CUDA,fp32legacy)文件可能适合你。这些文件是 llm.c 早期历史中的「checkpoint」版本,被定格在那一刻。它们更简单、更可移植,也可能更容易理解。单 GPU fp32 代码可按如下方式运行:
If you won't be training on multiple nodes, aren't interested in mixed precision, and are interested in learning CUDA, the fp32 (legacy) files might be of interest to you. These are files that were "checkpointed" early in the history of llm.c and frozen in time. They are simpler, more portable, and possibly easier to understand. Run the 1 GPU, fp32 code like this:
```bash
chmod u+x ./dev/download_starter_pack.sh
@@ -25,7 +19,7 @@ make train_gpt2fp32cu
./train_gpt2fp32cu
```
`download_starter_pack.sh` 脚本是快速上手的简便方式,会下载一批有助于入门的 .bin 文件,其中包含:1) 以 fp32、bfloat16 保存的 GPT-2 124M 模型,2) 单元测试中使用的「debug state」(小批量数据及目标激活值和梯度),3) GPT-2 分词器,以及 3) 已分词化的 [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) 数据集。或者,不运行 .sh 脚本,也可按以下步骤手动重新创建这些产物:
The download_starter_pack.sh script is a quick & easy way to get started and it downloads a bunch of .bin files that help get you off the ground. These contain: 1) the GPT-2 124M model saved in fp32, in bfloat16, 2) a "debug state" used in unit testing (a small batch of data, and target activations and gradients), 3) the GPT-2 tokenizer, and 3) the tokenized [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) dataset. Alternatively, instead of running the .sh script, you can re-create these artifacts manually as follows:
```bash
pip install -r requirements.txt
@@ -35,7 +29,7 @@ python train_gpt2.py
## quick start (CPU)
「穷到连一块 GPU 都没有」专区。你仍然可以欣赏 llm.c 的训练过程!但走不了太远。与上述 fp32 版本类似,CPU 版本是 llm.c 历史上更早的一个 checkpoint,当时它还只是 C 语言的简单参考实现。例如,你可以不从头训练,而是将 GPT-2 small (124M) 微调为输出类莎士比亚文本,示例如下:
The "I am so GPU poor that I don't even have one GPU" section. You can still enjoy seeing llm.c train! But you won't go too far. Just like the fp32 version above, the CPU version is an even earlier checkpoint in the history of llm.c, back when it was just a simple reference implementation in C. For example, instead of training from scratch, you can finetune a GPT-2 small (124M) to output Shakespeare-like text, as an example:
```bash
chmod u+x ./dev/download_starter_pack.sh
@@ -44,9 +38,9 @@ make train_gpt2
OMP_NUM_THREADS=8 ./train_gpt2
```
如果你更希望不运行 starter pack 脚本,如上一节所述,可运行 `python dev/data/tinyshakespeare.py` 再运行 `python train_gpt2.py`,复现完全相同的 .bin 文件和产物。
If you'd prefer to avoid running the starter pack script, then as mentioned in the previous section you can reproduce the exact same .bin files and artifacts by running `python dev/data/tinyshakespeare.py` and then `python train_gpt2.py`.
上述命令(1)下载已分词化的 [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) 数据集并下载 GPT-2 (124M) 权重,(3)在 C 中从它们初始化,在 tineshakespeare 上用 AdamW 训练 40 步(batch size 4context length 仅 64),评估验证集损失并采样部分文本。坦白说,除非你有一块强劲的 CPU(并能在启动命令中调高 OMP 线程数),否则在 CPU 上训练 LLM 走不了太远,但这可能是个不错的演示/参考。在我的 MacBook ProApple Silicon M3 Max)上,输出如下:
The above lines (1) download an already tokenized [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) dataset and download the GPT-2 (124M) weights, (3) init from them in C and train for 40 steps on tineshakespeare with AdamW (using batch size 4, context length only 64), evaluate validation loss, and sample some text. Honestly, unless you have a beefy CPU (and can crank up the number of OMP threads in the launch command), you're not going to get that far on CPU training LLMs, but it might be a good demo/reference. The output looks like this on my MacBook Pro (Apple Silicon M3 Max):
```
[GPT-2]
@@ -84,31 +78,31 @@ Allay
## datasets
`/dev/data/(dataset).py` 内的数据文件负责下载、分词并将 token 保存为可从 C 轻松读取的 .bin 文件。例如运行:
The data files inside `/dev/data/(dataset).py` are responsible for downloading, tokenizing and saving the tokens to .bin files, readable easily from C. So for example when you run:
```bash
python dev/data/tinyshakespeare.py
```
我们会下载并对 [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) 数据集进行分词。输出如下:
We download and tokenize the [tinyshakespeare](https://raw.githubusercontent.com/karpathy/char-rnn/master/data/tinyshakespeare/input.txt) dataset. The output of this looks like this:
```
writing 32,768 tokens to ./dev/data/tinyshakespeare/tiny_shakespeare_val.bin
writing 305,260 tokens to ./dev/data/tinyshakespeare/tiny_shakespeare_train.bin
```
.bin 文件包含一个短头部(1024 字节),随后是以 uint16 存储的 token 流,表示 GPT-2 分词器的 token id。更多数据集可在 `/dev/data` 中找到。
The .bin files contain a short header (1024 bytes) and then a stream of tokens in uint16, indicating the token ids with the GPT-2 tokenizer. More datasets are available in `/dev/data`.
## test
我还附带了一个简单的单元测试,用于确保 C 代码与 PyTorch 代码一致。以 CPU 为例,编译并运行:
I am also attaching a simple unit test for making sure our C code agrees with the PyTorch code. On the CPU as an example, compile and run with:
```bash
make test_gpt2
./test_gpt2
```
该测试会加载由 train_gpt2.py 写入的 `gpt2_124M_debug_state.bin` 文件,执行前向传播,将 logits loss PyTorch 参考实现对比,然后进行 10 次 Adam 训练迭代并确保 loss 与 PyTorch 一致。要测试 GPU 版本,运行:
This now loads the `gpt2_124M_debug_state.bin` file that gets written by train_gpt2.py, runs a forward pass, compares the logits and loss with the PyTorch reference implementation, then it does 10 iterations of training with Adam and makes sure the losses match PyTorch. To test the GPU version we run:
```bash
# fp32 test (cudnn not supported)
@@ -117,19 +111,19 @@ make test_gpt2cu PRECISION=FP32 && ./test_gpt2cu
make test_gpt2cu USE_CUDNN=1 && ./test_gpt2cu
```
这会同时测试 fp32 路径和混合精度路径。测试应通过并打印 `overall okay: 1`
This tests both the fp32 path and the mixed precision path. The test should pass and print `overall okay: 1`.
## tutorial
我在此附上了一份很短的教程:[doc/layernorm/layernorm.md](doc/layernorm/layernorm.md)。这是一份简单、分步的指南,讲解如何实现 GPT-2 模型的单层——layernorm 层。这是理解 C 中如何实现各层的好起点。
I attached a very small tutorial here, in [doc/layernorm/layernorm.md](doc/layernorm/layernorm.md). It's a simple, step-by-step guide to implementing a single layer of the GPT-2 model, the layernorm layer. This is a good starting point to understand how the layers are implemented in C.
**flash attention**。截至 2024 年 5 月 1 日,我们使用 cuDNN 的 Flash Attention。由于 cuDNN 会将编译时间从几秒延长到约一分钟,且该代码路径目前非常新,因此默认禁用。可按如下方式编译以启用:
**flash attention**. As of May 1, 2024 we use the Flash Attention from cuDNN. Because cuDNN bloats the compile time from a few seconds to ~minute and this code path is right now very new, this is disabled by default. You can enable it by compiling like this:
```bash
make train_gpt2cu USE_CUDNN=1
```
这将尝试使用 cudnn 编译并运行。你的系统上必须已安装 cuDNN。[cuDNN installation instructions](https://developer.nvidia.com/cudnn) 中通过 apt-get 安装会获取默认的 cuDNN 包集。最小化安装时,cuDNN 开发包即可,例如在 Ubuntu 22.04CUDA 12.x 上:
This will try to compile with cudnn and run it. You have to have cuDNN installed on your system. The [cuDNN installation instructions](https://developer.nvidia.com/cudnn) with apt-get will grab the default set of cuDNN packages. For a minimal setup, the cuDNN dev package is sufficient, e.g. on Ubuntu 22.04 for CUDA 12.x:
```bash
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
@@ -138,45 +132,45 @@ sudo apt-get update
sudo apt-get -y install libcudnn9-dev-cuda-12
```
此外还需要 [cuDNN frontend](https://github.com/NVIDIA/cudnn-frontend/tree/main),,但这只是头文件。将仓库克隆到本地即可。Makefile 目前会在你的主目录或当前目录中查找它。若放在其他位置,请在 `make` 命令行中添加 `CUDNN_FRONTEND_PATH=/path/to/your/cudnn-frontend/include`
On top of this you need the [cuDNN frontend](https://github.com/NVIDIA/cudnn-frontend/tree/main), but this is just header files. Simply clone the repo to your disk. The Makefile currently looks for it in either your home directory or the current directory. If you have put it elsewhere, add `CUDNN_FRONTEND_PATH=/path/to/your/cudnn-frontend/include` to the `make` command-line.
## 多 GPU 训练
## multi-GPU training
请确保已安装 MPI NCCL,例如在 Linux 上:
Make sure you install MPI and NCCL, e.g. on Linux:
```bash
sudo apt install openmpi-bin openmpi-doc libopenmpi-dev
```
对于 NCCL,请按照[官方网站](https://developer.nvidia.com/nccl/nccl-download)(例如网络安装器)中的说明操作
For NCCL follow the instructions from the [official website](https://developer.nvidia.com/nccl/nccl-download) (e.g. network installer)
然后:
and then:
```bash
make train_gpt2cu
mpirun -np <number of GPUs> ./train_gpt2cu
```
或者直接运行 `./scripts/` 下的某个脚本。
or simply run one of our scripts under `./scripts/`.
## 多节点训练
## multi-node training
请确保已按照 [multi-GPU](#multi-gpu-training) 章节的说明安装 `NCCL`
Make sure you've installed `NCCL` following instructions from [multi-GPU](#multi-gpu-training) section.
目前我们支持 3 种方式来进行多节点训练:
1) 使用 OpenMPI 交换 nccl id 并初始化 NCCL。详见例如 `./scripts/multi_node/run_gpt2_124M_mpi.sh` 脚本。
2) 使用共享文件系统初始化 NCCL。详见 `./scripts/multi_node/run_gpt2_124M_fs.sbatch` 脚本。
3) 使用 TCP 套接字初始化 NCCL。详见 `./scripts/multi_node/run_gpt2_124M_tcp.sbatch` 脚本。
There are 3 ways we currently support that allow you to run multi-node training:
1) Use OpenMPI to exchange nccl id and initialize NCCL. See e.g. `./scripts/multi_node/run_gpt2_124M_mpi.sh` script for details.
2) Use shared file system to init NCCL. See `./scripts/multi_node/run_gpt2_124M_fs.sbatch` script for details.
3) Use TCP sockets to init NCCL. See `./scripts/multi_node/run_gpt2_124M_tcp.sbatch` script for details.
注意:
* 如果你在 slurm 环境中运行,且你的 slurm 不支持 PMIx(考虑到 `slurm-wlm` 已移除 PMIx 支持,我们认为这会是常见情况),则必须使用 FS(2)或 TCP(3)方式。要测试你的 slurm 是否支持 PMIx,请运行:`srun --mpi=list`,并查看输出中是否出现 `pmix`
* 如果你没有配置 slurm,可以使用 `mpirun` - MPI1)方式启动多节点运行。
Note:
* If you're running in a slurm environment and your slurm doesn't support PMIx (which we assume will be a common situation given that `slurm-wlm` dropped PMIx support) you will have to use FS (2) or TCP (3) approach. To test whether your slurm supports PMIx run: `srun --mpi=list` and see whether you get `pmix` in the output.
* If you don't have slurm set up, you can kick off a multi-node run using `mpirun` - MPI (1).
这 3 种方法并无优劣之分,我们只是提供多种选项,以便你在特定环境中运行。
None of these 3 methods is superior, we just offer you options so that you can run in your specific environment.
## 实验 / 参数扫描(sweeps
## experiments / sweeps
下面是一个示例流程:在一台配备 4 GPU 的机器上,对 TinyStories 进行学习率扫描。运行 shell 脚本 `sweep.sh`(当然,在此之前你需要先 `chmod u+x sweep.sh`):
Just as an example process to sweep learning rates on a machine with 4 GPUs on TinyStories. Run a shell script `sweep.sh` (after you of course `chmod u+x sweep.sh`):
```bash
#!/bin/bash
@@ -192,82 +186,76 @@ done
# screen -ls | grep -E "tr[0-3]" | cut -d. -f1 | xargs -I {} screen -X -S {} quit
```
此示例会开启 4 个 screen 会话,并以不同学习率运行四条命令。这会写出日志文件 `stories$i.log`,其中包含所有 loss,你可以用 Python 按需绘制。解析并绘制这些日志文件的快速示例见 [dev/vislog.ipynb](dev/vislog.ipynb)
This example opens up 4 screen sessions and runs the four commands with different LRs. This writes the log files `stories$i.log` with all the losses, which you can plot as you wish in Python. A quick example of how to parse and plot these logfiles is in [dev/vislog.ipynb](dev/vislog.ipynb).
## 仓库
## repo
关于我希望这个仓库成为什么样子,再多说几句:
A few more words on what I want this repo to be:
首先,我希望 `llm.c` 成为一个教学场所。例如,我们的 `dev/cuda` 文件夹是一个内核库,涵盖各层的手工编写、文档详尽的内核,从非常简单的内核一直到更复杂/更快的内核。如果你有具有不同权衡(tradeoffs)的新内核,欢迎在此贡献。
First, I want `llm.c` to be a place for education. E.g. our `dev/cuda` folder is a place for a library of kernels for all the layers that are manually hand-written and very well documented, starting from very simple kernels all the way to more complex / faster kernels. If you have a new kernel with various different tradeoffs, please feel free to contribute it here.
话虽如此,我也希望 `llm.c` 非常快,甚至能实际用于训练网络。例如,起步时我们应该能够复现大型 GPT-21.6B)训练运行。这要求纳入最快的内核,包括使用 cuBLAScuBLASLtCUTLASScuDNN 等库。我认为这样做也有教学意义:建立一个专家级上限和度量单位,例如你可以说你的手写内核达到 cuBLAS 速度的 80% 等。然后你可以选择超快运行,也可以选择“拖放”你想使用的任意手工内核来运行。
That said, I also want `llm.c` to be very fast too, even practically useful to train networks. E.g. to start, we should be able to reproduce the big GPT-2 (1.6B) training run. This requires that we incorporate whatever fastest kernels there are, including the use of libraries such as cuBLAS, cuBLASLt, CUTLASS, cuDNN, etc. I also think doing so serves an educational purpose to establish an expert upper bound, and a unit of measurement, e.g. you could say that your manually written kernels are 80% of cuBLAS speed, etc. Then you can choose to do a super fast run, or you can choose to "drag and drop" whatever manual kernels you wish to use, and run with those.
不过,作为约束,我希望保持根目录中的主线 `llm.c` 简单且可读。如果某个 PR 例如将性能提升 2%,但“代价”是 500 行复杂 C 代码,也许还有冷门第三方依赖,我可能会拒绝该 PR,因为复杂度不值得。举个具体例子——在根训练循环中将 cuBLAS 用于 matmul 设为默认是显而易见的:它让主线代码快得多,只是一行可理解的代码,且是非常常见的依赖。与此同时,我们可以在 `dev/cuda` 中提供能与 cuBLAS 竞争的手工实现。
However, as a constraint, I want to keep the mainline `llm.c` in the root folder simple and readable. If there is a PR that e.g. improves performance by 2% but it "costs" 500 lines of complex C code, and maybe an exotic 3rd party dependency, I may reject the PR because the complexity is not worth it. As a concrete example - making cuBLAS for matmuls the default in the root training loop is a no-brainer: it makes the mainline code much faster, it is a single line of interpretable code, and it is a very common dependency. On the side of this, we can have manual implementations that can compete with cuBLAS in `dev/cuda`.
最后,我对项目根文件夹中的复杂度会敏感得多,根目录包含项目的主/默认文件。相比之下,`dev/` 文件夹更像是我们开发内核或类库、分享有用/相关/教学代码的草稿空间,其中部分代码可以(在局部)复杂一些。
Lastly, I will be a lot more sensitive to complexity in the root folder of the project, which contains the main / default files of the project. In comparison, the `dev/` folder is a bit more of a scratch space for us to develop a library of kernels or classes and share useful or related or educational code, and some of this code could be ok to be (locally) complex.
## 值得关注的分支(fork
## notable forks
- AMD 支持
- [llm.c](https://github.com/anthonix/llm.c) by @[anthonix](https://github.com/anthonix): 支持 AMD 设备,例如 7900 XTX
- AMD support
- [llm.c](https://github.com/anthonix/llm.c) by @[anthonix](https://github.com/anthonix): support for AMD devices, such as the 7900 XTX
- C#
- [llm.cs](https://github.com/azret/llm.cs) by @[azret](https://github.com/azret): 本项目的 C# 移植版
- [Llm.cs](https://github.com/nietras/Llm.cs) by @[nietras](https://github.com/nietras): 本项目的 C# 移植版,侧重于在任何平台上易于上手。克隆即可运行
- [llm.cs](https://github.com/azret/llm.cs) by @[azret](https://github.com/azret): a C# port of this project
- [Llm.cs](https://github.com/nietras/Llm.cs) by @[nietras](https://github.com/nietras): a C# port of this project with focus on easy to get started on any platform. Clone and run
- CUDA C++
- [llm.cpp](https://github.com/gevtushenko/llm.c) by @[gevtushenko](https://github.com/gevtushenko): 使用 [CUDA C++ Core Libraries](https://github.com/NVIDIA/cccl) 对本项目的移植
- 该分支曾在 [GPU MODE Discord Server](https://discord.gg/cudamode) 的[这场讲座](https://www.youtube.com/watch?v=WiB_3Csfj_Q) 中介绍
- C++/CUDA
- [llm.cpp](https://github.com/zhangpiu/llm.cpp/tree/master/llmcpp) by @[zhangpiu](https://github.com/zhangpiu): 使用 [Eigen](https://gitlab.com/libeigen/eigen), 对本项目的移植,支持 CPU/CUDA。
- [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)
- WebGPU C++
- [gpu.cpp](https://github.com/AnswerDotAI/gpu.cpp) by @[austinvhuang](https://github.com/austinvhuang): 使用原生 WebGPU 在 C++ 中进行可移植 GPU 计算的库。目标是成为通用库,同时也将 llm.c 内核移植到 WGSL
- C++
- [llm.cpp](https://github.com/GaoYusong/llm.cpp) by @[GaoYusong](https://github.com/GaoYusong): 本项目的 C++ 移植版,提供 C++ 单头文件库 [tinytorch.hpp](https://github.com/GaoYusong/llm.cpp/blob/main/tinytorch.hpp)
- [gpu.cpp](https://github.com/AnswerDotAI/gpu.cpp) by @[austinvhuang](https://github.com/austinvhuang): a library for portable GPU compute in C++ using native WebGPU. Aims to be a general-purpose library, but also porting llm.c kernels to WGSL.
- Go
- [llm.go](https://github.com/joshcarp/llm.go) by @[joshcarp](https://github.com/joshcarp): 本项目的 Go 移植版
- [llm.go](https://github.com/joshcarp/llm.go) by @[joshcarp](https://github.com/joshcarp): a Go port of this project
- Java
- [llm.java](https://github.com/harryjackson/llm.java) by @[harryjackson](https://github.com/harryjackson): 本项目的 Java 移植版
- [llm.java](https://github.com/harryjackson/llm.java) by @[harryjackson](https://github.com/harryjackson): a Java port of this project
- Metal
- [llm.metal](https://github.com/regrettable-username/llm.metal) by @[regrettable-username](https://github.com/regrettable-username): 使用简洁、原生 C/Metal Shading Language 进行 LLM 训练
- [llm.metal](https://github.com/regrettable-username/llm.metal) by @[regrettable-username](https://github.com/regrettable-username): LLM training in simple, raw C/Metal Shading Language
- Mojo
- [llm.🔥](https://github.com/dorjeduck/llm.mojo) by @[dorjeduck](https://github.com/dorjeduck): 本项目的 Mojo 移植版
- [llm.🔥](https://github.com/dorjeduck/llm.mojo) by @[dorjeduck](https://github.com/dorjeduck): a Mojo port of this project
- OpenCL
- [llm.c](https://github.com/krrishnarraj/llm.c) by @[krrishnarraj](https://github.com/krrishnarraj): 本项目的 OpenCL 移植版
- [llm.c](https://github.com/krrishnarraj/llm.c) by @[krrishnarraj](https://github.com/krrishnarraj): an OpenCL port of this project
- Rust
- [llm.rs](https://github.com/yijunyu/llm.rs) by @[Yijun Yu](https://github.com/yijunyu): Rust 重写版,目标是达到相同性能
- [llm.rs](https://github.com/ToJen/llm.rs) by @[ToJen](https://github.com/ToJen): 本项目的 Rust 移植版
- [llm.rs](https://github.com/yijunyu/llm.rs) by @[Yijun Yu](https://github.com/yijunyu): a Rust rewrite with the aim to have same performance
- [llm.rs](https://github.com/ToJen/llm.rs) by @[ToJen](https://github.com/ToJen): a Rust port of this project
- Swift
- [llm.swift](https://github.com/otabuzzman/llm.swift) by @[otabuzzman](https://github.com/otabuzzman): 本项目的 Swift 移植版
- [llm.swift](https://github.com/otabuzzman/llm.swift) by @[otabuzzman](https://github.com/otabuzzman): a Swift port of this project
- Zig
- [llm.zig](https://github.com/Saimirbaci/llm.zig) by @[saimirbaci](https://github.com/Saimirbaci): 本项目的 Zig 移植版
- [llm.zig](https://github.com/Saimirbaci/llm.zig) by @[saimirbaci](https://github.com/Saimirbaci): a Zig port of this project
- Habana Gaudi2
- [llm.tpc](https://github.com/abhilash1910/llm.tpc) by @[abhilash1910](https://github.com/abhilash1910): 本项目的 Habana Gaudi2 移植版
- [llm.tpc](https://github.com/abhilash1910/llm.tpc) by @[abhilash1910](https://github.com/abhilash1910): a Habana Gaudi2 port of this project
- Nim
- [llm.nim](https://github.com/Vindaar/llm.nim) by @[Vindaar](https://github.com/Vindaar): 本项目的 Nim 移植版
- [llm.nim](https://github.com/Vindaar/llm.nim) by @[Vindaar](https://github.com/Vindaar): a Nim port of this project
## 讨论
## discussions
组织开发的方式:
Ways of organizing development:
- 在使用仓库时遇到具体问题?请使用 [Issues](https://github.com/karpathy/llm.c/issues).
- 有代码要贡献?请提交 [PR](https://github.com/karpathy/llm.c/pulls)
- 想聊聊仓库、提问等?请查看 [Discussions](https://github.com/karpathy/llm.c/discussions).
- 需要更快捷的沟通?我在我的 [Zero to Hero Discord 频道](https://discord.gg/3zy8kqD9Cp). 上新建了 `#llmc` 频道
- Experiencing a concrete issue with the repo? Use [Issues](https://github.com/karpathy/llm.c/issues).
- Have some code to contribute? Open a [PR](https://github.com/karpathy/llm.c/pulls)
- Chat about the repo, ask questions, etc.? Look at [Discussions](https://github.com/karpathy/llm.c/discussions).
- Something faster? I created a new `#llmc` channel on my [Zero to Hero Discord channel](https://discord.gg/3zy8kqD9Cp).
## 许可证
## license
MIT
-7
View File
@@ -1,7 +0,0 @@
# WeHub 来源说明
- 原始项目:`karpathy/llm.c`
- 原始仓库:https://github.com/karpathy/llm.c
- 导入方式:上游默认分支的最新快照
- 原作者、版权和许可证信息以原始仓库及本仓库 LICENSE 为准
- 本文件仅用于记录来源,不代表 WeHub 是原项目作者
+1 -1
View File
@@ -38,7 +38,7 @@ metrics = [
"dram__bytes_read.sum", # DRAM reads
"dram__bytes_write.sum", # DRAM writes
"lts__t_sectors_srcunit_tex_op_read.sum", # L2 reads (sectors -- 32B)
"lts__t_sectors_srcunit_tex_op_write.sum", # L2 writes (sectors -- 32B)
"lts__t_sectors_srcunit_tex_op_write.sum", # L2 reads (sectors -- 32B)
"sm__pipe_tensor_op_hmma_cycles_active.avg.pct_of_peak_sustained_active", # % of peak tensor core utilization
"smsp__inst_executed.sum", # instructions
]
+20 -15
View File
@@ -113,6 +113,8 @@ int main(int argc, char *argv[]) {
size_t V = model.config.vocab_size;
size_t Vp = model.config.padded_vocab_size;
size_t maxT = model.config.max_seq_len;
size_t L = model.config.num_layers;
size_t C = model.config.channels;
for (int i = 1; i < argc; i+=2) {
if (i + 1 >= argc) { exit(EXIT_FAILURE); } // must have arg after flag
@@ -265,26 +267,29 @@ int main(int argc, char *argv[]) {
// Also, different GPUs may use different matrix multiplication algorithms, so the
// actual errors can be hardware specific.
float grad_thresholds[NUM_PARAMETER_TENSORS] = {
5e-1f, 4e-3f, 1e-1f, 4e-2f,
5e-2f, 3.5e-2f, 2e-2f, 3e-2f,
5e-2f, 3e-2f, 3e-2f, 3e-2f,
2e-2f, 1e-2f,1e-1f,2e-2f};
float grad_thresholds[NUM_PARAMETER_TENSORS] = {5e-1f, 4e-3f, 1e-1f, 3.5e-2f, 2e-2f, 3e-2f, 5e-2f, 5e-2f, 5e-2f, 1.5e-2f, 5e-4f, 8e-3f, 1.5e-3f, 2.5e-3f, 1e-1f, 2e-2f};
#if defined(ENABLE_FP32)
for (int i = 0; i < NUM_PARAMETER_TENSORS; i++) {
grad_thresholds[i] = 1e-6f; // we can be much more precise in FP32
}
#endif
const char* names[NUM_PARAMETER_TENSORS] = {
"wte", "wpe", "ln1w", "ln1b", "qkvw", "qkvb", "attrpojw",
"attprojb", "ln2w", "ln2b", "fcw", "fcb", "fcprojw", "fcprojb",
"lnfw", "lnfb"
};
size_t* count = model.param_elements;
for(int i = 0; i < NUM_PARAMETER_TENSORS; ++i) {
allok = allok & check_tensor(tensors1[i], tensors2[i], count[i], names[i], grad_thresholds[i]);
}
allok = allok & check_tensor(tensors1[0], tensors2[0], V * C, "wte", grad_thresholds[0]);
allok = allok & check_tensor(tensors1[1], tensors2[1], maxT * C, "wpe", grad_thresholds[1]);
allok = allok & check_tensor(tensors1[2], tensors2[2], L * 3*C * C, "qkvw", grad_thresholds[2]);
allok = allok & check_tensor(tensors1[3], tensors2[3], L * 3*C, "qkvb", grad_thresholds[3]);
allok = allok & check_tensor(tensors1[4], tensors2[4], L * C * C, "attprojw", grad_thresholds[4]);
allok = allok & check_tensor(tensors1[5], tensors2[5], L * C, "attprojb", grad_thresholds[5]);
allok = allok & check_tensor(tensors1[6], tensors2[6], L * 4*C * C, "fcw", grad_thresholds[6]);
allok = allok & check_tensor(tensors1[7], tensors2[7], L * 4*C, "fcb", grad_thresholds[7]);
allok = allok & check_tensor(tensors1[8], tensors2[8], L * C * 4*C, "fcprojw", grad_thresholds[8]);
allok = allok & check_tensor(tensors1[9], tensors2[9], L * C, "fcprojb", grad_thresholds[9]);
allok = allok & check_tensor(tensors1[10], tensors2[10], L * C, "ln1w", grad_thresholds[10]);
allok = allok & check_tensor(tensors1[11], tensors2[11], L * C, "ln1b", grad_thresholds[11]);
allok = allok & check_tensor(tensors1[12], tensors2[12], L * C, "ln2w", grad_thresholds[12]);
allok = allok & check_tensor(tensors1[13], tensors2[13], L * C, "ln2b", grad_thresholds[13]);
allok = allok & check_tensor(tensors1[14], tensors2[14], C, "lnfw", grad_thresholds[14]);
allok = allok & check_tensor(tensors1[15], tensors2[15], C, "lnfb", grad_thresholds[15]);
}
float grad_norm = gpt2_calculate_grad_norm(&model, &multi_gpu_config);
+59 -27
View File
@@ -16,17 +16,17 @@ Example launches to only benchmark the speed of bfloat16 compiled GPU training:
TODO: add the actual commands
"""
import argparse
import os
import math
import glob
import inspect
from contextlib import nullcontext
from dataclasses import dataclass
import json
from pathlib import Path
import time
from typing import (
AbstractSet,
Callable,
Collection,
Dict,
Iterator,
@@ -55,6 +55,9 @@ from tiktoken.load import load_tiktoken_bpe
# -----------------------------------------------------------------------------
# PyTorch nn.Module definitions for the LLaMA 3.x model
# using a global to toggle flash-attention
FLASH = 0
# Used in Grouped Query Attention (GQA), broadcasts the key and value tensors
def repeat_kv(x: torch.Tensor, n_rep: int) -> torch.Tensor:
"""torch.repeat_interleave(x, dim=2, repeats=n_rep)"""
@@ -154,7 +157,6 @@ class CausalSelfAttention(nn.Module):
self.n_rep = self.n_head // self.n_kv_head
self.hd = config.n_embd // config.n_head
self.use_kv = config.use_kv
self.flash = config.flash
self.c_attn = nn.Linear(config.n_embd, (config.n_head + 2 * config.n_kv_head) * self.hd, bias=False) # key, query, value projections
self.c_proj = nn.Linear(config.n_embd, config.n_embd, bias=False) # output projection
@@ -184,12 +186,9 @@ class CausalSelfAttention(nn.Module):
q, k, v = map(lambda t: t.transpose(1, 2), (q, k, v)) # (B, NH, T, HD)
if self.flash:
if FLASH:
# flashattention
# if T == 1 no need to mask, otherwise the function complains
# scaled_dot_product_attention expects a mask where value of True indicates that the element should take part in attention
# our mask is the opposite, so we need to invert it
y = F.scaled_dot_product_attention(q, k, v, mask == 0 if T > 1 else None)
y = F.scaled_dot_product_attention(q, k, v, mask)
else:
# manual implementation of attention
# this materializes the large (T,T) matrix for all the queries and keys
@@ -258,7 +257,6 @@ class LlamaConfig:
use_scaled_rope: bool = True
max_gen_batch_size: int = 4
use_kv: bool = True
flash: bool = False # use flashattention?
def __init__(self, **kwargs):
for k, v in kwargs.items():
@@ -404,7 +402,7 @@ class LLaMA(nn.Module):
def from_pretrained_llama3_hf(cls, model_id):
"""Loads pretrained LLaMA model weights from HuggingFace"""
from transformers import AutoModelForCausalLM, AutoTokenizer
assert model_id == "meta-llama/Meta-Llama-3.1-8B", "Only the 8B-base model is supported for now"
assert model_id == "meta-llama/Meta-Llama-3.1-8B", "Only the 8B-bae model is supported for now"
model_args = LlamaConfig()
model = AutoModelForCausalLM.from_pretrained(model_id)
@@ -479,6 +477,7 @@ class LLaMA(nn.Module):
max_gen_len: int,
temperature: float = 0.6,
top_p: float = 0.9,
logprobs: bool = False,
echo: bool = False,
) -> Tuple[List[List[int]], Optional[List[List[float]]]]:
"""
@@ -489,28 +488,32 @@ class LLaMA(nn.Module):
max_gen_len (int): Maximum length of the generated text sequence.
temperature (float, optional): Temperature value for controlling randomness in sampling. Defaults to 0.6.
top_p (float, optional): Top-p probability threshold for nucleus sampling. Defaults to 0.9.
logprobs (bool, optional): Flag indicating whether to compute token log probabilities. Defaults to False.
echo (bool, optional): Flag indicating whether to include prompt tokens in the generated output. Defaults to False.
Returns:
Tuple[List[List[int]], Optional[List[List[float]]]]: A tuple containing generated token sequences.
Tuple[List[List[int]], Optional[List[List[float]]]]: A tuple containing generated token sequences and, if logprobs is True, corresponding token log probabilities.
Note:
This method uses the provided prompts as a basis for generating text. It employs nucleus sampling to produce text with controlled randomness.
If logprobs is True, token log probabilities are computed for each generated token.
"""
bsz = len(prompt_tokens)
assert bsz <= self.config.max_gen_batch_size, f"Batch size {bsz} exceeds the maximum generation batch size {self.config.max_gen_batch_size}"
assert bsz <= self.config.max_gen_batch_size, (bsz, self.config.max_gen_batch_size)
device = next(self.parameters()).device
min_prompt_len = min(len(t) for t in prompt_tokens)
max_prompt_len = max(len(t) for t in prompt_tokens)
assert max_prompt_len <= self.config.block_size, f"Prompt length {max_prompt_len} exceeds the maximum block size {self.config.block_size}"
assert max_prompt_len <= self.config.block_size
total_len = min(self.config.block_size, max_gen_len + max_prompt_len)
pad_id = self.tokenizer.pad_id
tokens = torch.full((bsz, total_len), pad_id, dtype=torch.long, device=device)
for idx, t in enumerate(prompt_tokens):
tokens[idx, : len(t)] = torch.tensor(t, dtype=torch.long, device=device)
for k, t in enumerate(prompt_tokens):
tokens[k, : len(t)] = torch.tensor(t, dtype=torch.long, device=device)
if logprobs:
token_logprobs = torch.zeros_like(tokens, dtype=torch.float)
prev_pos = 0
eos_reached = torch.tensor([False] * bsz, device=device)
@@ -518,6 +521,12 @@ class LLaMA(nn.Module):
if min_prompt_len == total_len:
logits, _ = self.forward(tokens, start_pos=prev_pos)
token_logprobs = -F.cross_entropy(
input=logits.transpose(1, 2),
target=tokens,
reduction="none",
ignore_index=pad_id,
)
stop_tokens = torch.tensor(list(self.tokenizer.stop_tokens)).to(device)
@@ -533,25 +542,41 @@ class LLaMA(nn.Module):
# only replace token if prompt has already been generated
next_token = torch.where(input_text_mask[:, cur_pos], tokens[:, cur_pos], next_token)
tokens[:, cur_pos] = next_token
eos_reached |= ~input_text_mask[:, cur_pos] & torch.isin(next_token, stop_tokens)
if logprobs:
token_logprobs[:, prev_pos + 1 : cur_pos + 1] = -F.cross_entropy(
input=logits.transpose(1, 2),
target=tokens[:, prev_pos + 1 : cur_pos + 1],
reduction="none",
ignore_index=pad_id,
)
eos_reached |= (~input_text_mask[:, cur_pos]) & (
torch.isin(next_token, stop_tokens)
)
prev_pos = cur_pos
if all(eos_reached):
break
out_tokens = []
if logprobs:
token_logprobs = token_logprobs.tolist()
out_tokens, out_logprobs = [], []
for i, toks in enumerate(tokens.tolist()):
# cut to max gen len
start = 0 if echo else len(prompt_tokens[i])
toks = toks[start : len(prompt_tokens[i]) + max_gen_len]
probs = None
if logprobs:
probs = token_logprobs[i][start : len(prompt_tokens[i]) + max_gen_len]
# cut to after eos tok if any
for stop_token in self.tokenizer.stop_tokens:
try:
eos_idx = toks.index(stop_token)
toks = toks[:eos_idx]
probs = probs[:eos_idx] if logprobs else None
except ValueError:
pass
out_tokens.append(toks)
return out_tokens
out_logprobs.append(probs)
return (out_tokens, out_logprobs if logprobs else None)
# -----------------------------------------------------------------------------
# sampling utils
@@ -934,16 +959,18 @@ def print0(*args, **kwargs):
print(*args, **kwargs)
if __name__ == "__main__":
import time
import argparse
print0(f"Running pytorch {torch.version.__version__}")
# default settings will overfit a tiny batch of data
# and save model weights and debug state to disk on the first iteration
parser = argparse.ArgumentParser()
parser.add_argument("--use_hf", type=int, default=1, help="use HuggingFace (default) or use Meta's model")
parser.add_argument("--ckpt_dir", type=str, default=None, help="path to llama3 model checkpoint (needed if use_hf=0)")
parser.add_argument("--tokenizer_path", type=str, default=None, help="path to llama3 tokenizer (needed if use_hf=0)")
parser.add_argument("--ckpt_dir", type=str, default=None, help="path to llama3 model checkpoint")
parser.add_argument("--tokenizer_path", type=str, default=None, help="path to llama3 tokenizer")
# file system input / output
parser.add_argument("--input_bin", type=str, default="dev/data/tinyshakespeare/tiny_shakespeare_val.bin", help="input .bin to train on")
parser.add_argument("--input_bin", type=str, default="dev/data/tinystories/TinyStories_val.bin", help="input .bin to train on")
parser.add_argument("--input_val_bin", type=str, default="", help="input .bin to eval validation loss on")
parser.add_argument("--output_dir", type=str, default="", help="output directory to which to write logs and checkpoints")
parser.add_argument("--model", type=str, default="meta-llama/Meta-Llama-3.1-8B", help="chose the llama model")
@@ -955,7 +982,7 @@ if __name__ == "__main__":
parser.add_argument("--num_iterations", type=int, default=10, help="number of iterations to run")
parser.add_argument("--inference_only", type=int, default=0, help="only run inference")
# optimization
parser.add_argument("--learning_rate", type=float, default=1e-5, help="learning rate warmup iterations")
parser.add_argument("--learning_rate", type=float, default=1e-4, help="learning rate warmup iterations")
parser.add_argument("--warmup_iters", type=int, default=0, help="learning rate warmup iterations")
parser.add_argument("--learning_rate_decay_frac", type=float, default=1.0, help="learning rate warmup iterations")
parser.add_argument("--weight_decay", type=float, default=0.0, help="weight decay")
@@ -971,6 +998,7 @@ if __name__ == "__main__":
# memory management
parser.add_argument("--device", type=str, default="", help="by default we autodetect, or set it here")
parser.add_argument("--compile", type=int, default=0, help="torch.compile the model")
parser.add_argument("--flash", type=int, default=0, help="use flash attention")
parser.add_argument("--dtype", type=str, default="bfloat16", help="float32|float16|bfloat16")
parser.add_argument("--zero_stage", type=int, default=0, help="zero redundancy optimizer stage (0/1/2/3)")
# python -> C bridge
@@ -1024,9 +1052,9 @@ if __name__ == "__main__":
device = "cuda"
elif hasattr(torch.backends, "mps") and torch.backends.mps.is_available():
device = "mps"
device_type = 'cuda' if 'cuda' in device else 'cpu'
assert device_type in {'cuda'}, "GPU required to run LLaMA 3" # we need to load LLaMA as bf16 on CUDA
print(f"using device: {device}")
device_type = 'cuda' if 'cuda' in device else 'cpu'
assert device_type in {'cuda'} # we need to load LLaMA as bf16 on CUDA
# calculate gradient accumulation from the desired total batch size and the current run configuration
tokens_per_fwdbwd = B * T * ddp_world_size
@@ -1049,12 +1077,16 @@ if __name__ == "__main__":
if args.tensorcores:
torch.set_float32_matmul_precision('high')
# turn on/off flash attention
assert args.flash in {0, 1}
FLASH = args.flash
# init the model
assert args.ckpt_dir is not None and os.path.exists(args.ckpt_dir), f"llama3 ckpt dir {args.ckpt_dir} does not exist"
assert args.tokenizer_path is not None and os.path.exists(args.tokenizer_path), f"llama3 tokenizer path {args.tokenizer_path} does not exist"
if args.use_hf:
model = LLaMA.from_pretrained_llama3_hf(args.model)
else: # use Meta's checkpoint
assert args.ckpt_dir is not None and os.path.exists(args.ckpt_dir), f"llama3 ckpt dir {args.ckpt_dir} does not exist"
assert args.tokenizer_path is not None and os.path.exists(args.tokenizer_path), f"llama3 tokenizer path {args.tokenizer_path} does not exist"
model = LLaMA.from_pretrained_llama3_meta(args.ckpt_dir, args.tokenizer_path)
model.train()
@@ -1169,7 +1201,7 @@ if __name__ == "__main__":
else: # Meta
prompt_tokens = [model.tokenizer.encode(x, bos=True, eos=False) for x in prompts]
generation_tokens = model.generate(prompt_tokens, max_gen_len=64, temperature=0.6, top_p=0.9, echo=False)
generation_tokens, _ = model.generate(prompt_tokens, max_gen_len=64, temperature=0.6, top_p=0.9, logprobs=False, echo=False)
results = [{"generation": model.tokenizer.decode(t)} for t in generation_tokens]
for prompt, result in zip(prompts, results):
print(prompt, end="")