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
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
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
Erik Schultheis
2287da0120
enable bf16
2024-05-12 19:42:32 +03:00
Erik Schultheis
ec92368f8b
added current backward bias kernel to dev cuda
2024-05-12 19:24:43 +03:00
lancer
5c90845f7a
update the -lcublas -lcublasLt flag in the comment
2024-05-09 22:04:02 -07: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
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
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
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
ademeure
9b55ea8cf5
More crazy optimisations to layernorm_backward, fused_classifier, and matmul_backward_bias.
2024-05-07 00:11:42 +01:00
lancer
bbfe8c989c
Minor update on the code
2024-05-05 16:59:58 -07:00
lancer
c81adeb6b2
move gelu_backward to backward block
2024-05-02 17:36:12 -07: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
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
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
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
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
“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
52b4dd02a1
update residual forward to allow kernels to use both types
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
7746217433
update gelu backward rto allow all kernels to use both types
2024-05-01 23:36:59 +00:00
Jane Illarionova
80072802d7
Create gelu_backward.cu
2024-05-01 23:36:59 +00: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
ademeure
1147983b8c
Add missing cudnnDestroy() calls
2024-05-01 17:19:51 +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
c4ecc04dc3
Merge remote-tracking branch 'karpathy/master' into cudnn_try2
2024-05-01 15:30:00 +01:00
Andrej Karpathy
769c911ab6
small tweaks
2024-04-30 23:48:06 +00: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
