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 |
|
Andrej Karpathy
|
7ba1ed8a1a
|
delete packed128 .fp32, use (float) ok
|
2024-04-30 19:39:20 +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 |
|
Franz Louis Cesista
|
b239b67ae1
|
add modal script
|
2024-05-01 00:15:08 +08: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 |
|
Christopher Dryden
|
568615fef1
|
Added packing for gelu forwards kernel
|
2024-04-30 03:21:42 +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
|
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 |
|
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 |
|
Erik Schultheis
|
896b644cb9
|
slight simplification
|
2024-04-29 22:19:12 +03:00 |
|
Erik Schultheis
|
45252d434c
|
yet another gelu
|
2024-04-29 21:53:08 +03:00 |
|
ademeure
|
d23bb1fd13
|
Merge remote-tracking branch 'karpathy/master' into layernormbf16
|
2024-04-29 19:41:54 +01:00 |
|
ademeure
|
49ffd68317
|
Fix sizeof(T) -> sizeof(D) (was overallocating memory)
|
2024-04-29 15:47:59 +01: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
|
327eef3f22
|
incorporate faster encoder_forward kernel to fp32 CUDA version
|
2024-04-28 20:10:33 +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
|
12da2c16cc
|
slight docs tweaks to dev/cuda
|
2024-04-27 14:40:57 +00: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
|
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
|
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 |
|
Petr Zhizhin
|
d49e8301eb
|
[Multi-GPU] llm.c now runs on multiple GPUs with NCCL
|
2024-04-24 22:11:00 +00: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
|
249140287f
|
rename positional_forward to encoder_forward as it should have been
|
2024-04-23 18:19:20 +00:00 |
|