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