 Erik Schultheis
|
2b76155976
|
small fixes based on clang-tidy
|
2024-07-09 18:37:09 +02:00 |
|
|
Erik Schultheis
|
486f98e3cc
|
zero-grad is async and part of backward call
|
2024-07-02 00:47:41 +02: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 |
|
|
Aleksa Gordic
|
74e5ba7260
|
Add micro step guard
|
2024-07-01 12:11:02 +02:00 |
|
|
ademeure
|
fa5a196be0
|
Refactor matmul to only use cuBLASLt and optionally use fused GELU
|
2024-06-30 20:45:37 +01:00 |
|
|
Andrej Karpathy
|
e12d090483
|
resolve merge conflicts
|
2024-06-26 04:05:46 +00:00 |
|
|
Erik Schultheis
|
7b13f20343
|
moved targets to backward/validate
|
2024-06-26 01:56:23 +03:00 |
|
|
Erik Schultheis
|
8d994782cf
|
moved loss calculation to backward
|
2024-06-25 21:28:52 +03:00 |
|
|
Andrej Karpathy
|
329aee333e
|
fix training loop in other parts
|
2024-06-25 04:30:34 +00:00 |
|
|
Andrej
|
76b133d3f4
|
Merge pull request #631 from gordicaleksa/minor_refactor2
Minor refactor
|
2024-06-24 14:38:27 -07:00 |
|
|
Aleksa Gordic
|
c420a522b2
|
Fix arg permutation
|
2024-06-24 14:38:49 +00:00 |
|
|
Aleksa Gordic
|
a259eeabe5
|
Fix CI errors
|
2024-06-24 16:26:12 +02: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 |
|
|
Aleksa Gordic
|
8079bd9dbb
|
Relax grad tensor thresholds a bit
|
2024-06-19 10:21:59 +02:00 |
|
|
Andrej
|
6ecc52e3fc
|
Merge pull request #614 from gordicaleksa/more_strict_tests
Stricter FP32 tests
|
2024-06-18 09:29:13 -07: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
|
077f9bdd59
|
Fix dataloader & determinism testing
|
2024-06-18 13:23:13 +02:00 |
|
|
Erik Schultheis
|
dbeb8fc551
|
added missing checks
|
2024-06-18 13:29:12 +03:00 |
|
|
Andrej Karpathy
|
5b4fb0d457
|
delete intermediate files produced by test code
|
2024-06-18 03:02:53 +00:00 |
|
|
Erik Schultheis
|
c06ef28a0d
|
also test dataloader as part of test script
|
2024-06-18 00:24:42 +03: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 |
|
|
Erik Schultheis
|
6d2c72cb4b
|
rearranged adam code
|
2024-06-12 13:16:19 +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 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 Karpathy
|
8b951056b5
|
resolve merge conflict
|
2024-05-24 22:04:02 +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 |
|
|
ademeure
|
b5e75dde8e
|
Fully deterministic encoder backward kernels for train_gpt2.cu
|
2024-05-21 15:57:07 +01:00 |
|
|
Erik Schultheis
|
a9947a8315
|
added a flag and the missing sqrt;
testing now has clipping enabled
|
2024-05-18 22:47:15 +03: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 |
|
|
Erik Schultheis
|
57f70ea66b
|
simplify multi-gpu logic by reducing #ifdefs
|
2024-05-16 14:12:14 +03: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
|
26dbbc75c2
|
unbreak the tests oops
|
2024-05-08 20:18:33 +00: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
|
bfb9c51446
|
refactoring & remove unused functions to reduce LOC (+wip profile.py improvements)
|
2024-05-04 23:40:15 +01:00 |
|
|
ademeure
|
a19fcecb79
|
Set cuBLAS(Lt) to use the new main_stream
|
2024-05-03 18:35:46 +01: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 |
|
|
Andrej Karpathy
|
1ed9381fd0
|
fix cudnnDestroy outside of ifdef
|
2024-05-01 17:55:36 +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
|
c778dc0098
|
slightly increase wte threshold to make test_gpt2.cu pass (?)
|
2024-05-01 15:40:30 +01:00 |
|
|
ademeure
|
c4ecc04dc3
|
Merge remote-tracking branch 'karpathy/master' into cudnn_try2
|
2024-05-01 15:30:00 +01:00 |
|
|
Andrej Karpathy
|
44656c3850
|
allow fp32 precision in the test script as well
|
2024-04-30 18:38:38 +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 |
|
|
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 |
|