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
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
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
1c516c732d
re-commit of my old cudnn forward attention changes
2024-04-30 14:54:35 +01:00
Erik Schultheis
8ffc3e388b
slow attention kernel that demonstrates a direct translation of the CPU code
2024-04-20 19:03:34 +03: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
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
69f0385237
add comments and handling of kernels and their correctness
2024-04-13 17:44:31 +00: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
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
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
59b9e8b96a
fix handles in cublas in attention forward as well
2024-04-12 23:15:47 +00: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
Erik Schultheis
baba3cfa01
recalculate instead of memory round-trip
2024-04-12 02:20:20 +03:00
Erik Schultheis
5f99fb30d2
cooperative groups and fused scale kernel
2024-04-11 23:06:46 +03: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
Andrej Karpathy
bb4f04bdf8
add flash attention, but a naive implementation there of
2024-04-09 18:22:49 +00: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
