mirrors / karpathy--llm.c

提交

提交图

  • 59b66f6c07 refactor and fix CI issue chinthysl 2024-05-04 16:22:44 +08:00
  • 03b7323d1e minor fix of the include lancer 2024-05-03 17:38:39 -07:00
  • ff2fbdc2a1 Merge pull request #332 from harryjackson/llm.java Andrej 2024-05-03 13:58:59 -07:00
  • 67a82a611e first attempt at moving cudnn out of the main file for faster compiles Erik Schultheis 2024-05-03 23:08:44 +03:00
  • aa5dd2eeab Update README.md Andrej 2024-05-03 12:43:08 -07:00
  • 6260a1844b Merge pull request #345 from ngc92/profile Andrej 2024-05-03 12:17:22 -07:00
  • 977d6a6c6f fix × symbol Erik Schultheis 2024-05-03 21:57:31 +03:00
  • 7fabb12e2e fixed instruction counts Erik Schultheis 2024-05-03 21:54:06 +03:00
  • 13eeeb20c8 try to detect if we need sudo Erik Schultheis 2024-05-03 21:31:54 +03:00
  • 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 Erik Schultheis 2024-05-03 21:14:39 +03:00
  • a0b80920f1 Reducing profiler to only 3 steps instead of 10 given slight perf penalty (todo: make it configurable) ademeure 2024-05-03 19:03:06 +01:00
  • a19fcecb79 Set cuBLAS(Lt) to use the new main_stream ademeure 2024-05-03 18:35:46 +01:00
  • 6c629ac571 Merge branch 'master' of github.com:karpathy/llm.c Andrej Karpathy 2024-05-03 15:54:36 +00:00
  • e780b56974 small fixes for profiling Andrej Karpathy 2024-05-03 15:54:29 +00:00
  • b58181079a super optimised matmul_bias + cuda streams for better parallelism + fix fused_classifier cache hint + remove remaining cooperative groups ademeure 2024-05-03 16:38:46 +01:00
  • f133ee2294 Merge pull request #342 from ngc92/feature/grad_accum feature/grad_accum Andrej 2024-05-03 08:37:42 -07:00
  • d2bd2d6f21 fixed activation gradient resetting for backward pass Erik Schultheis 2024-05-03 17:57:01 +03:00
  • 0fd384bb21 Merge branch 'master' into zero-stage1 Chinthaka 2024-05-03 11:08:41 +08:00
  • f456a75e3f add CUDA streams (no parallelism yet) + fix profiling init + new "blockReduce" function ademeure 2024-05-03 03:14:16 +01:00
  • c81adeb6b2 move gelu_backward to backward block lancer 2024-05-02 17:36:12 -07:00
  • c04c133b25 Merge pull request #335 from karpathy/feature/bias_backward2 Andrej 2024-05-02 16:52:34 -07:00
  • 776541dac7 v1 of the new matmul backward bias kernel feature/bias_backward2 Andrej Karpathy 2024-05-02 23:32:20 +00:00
  • 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 Andrej Karpathy 2024-05-02 23:23:28 +00:00
  • 2c7960040d Merge branch 'ngc92-bias-backward' Andrej Karpathy 2024-05-02 22:17:31 +00:00
  • 937fbd8e36 small doc fixes Andrej Karpathy 2024-05-02 22:17:22 +00:00
  • 6e48501d60 bias backward kernel that will use all available threads Erik Schultheis 2024-05-03 01:05:47 +03:00
  • 6a3d63fb1d Merge branch 'PeterZhizhin-add_nvtx_ranges_and_timings' Andrej Karpathy 2024-05-02 21:13:26 +00:00
  • 79505bc6b3 resolve merge and small fixes Andrej Karpathy 2024-05-02 21:13:08 +00:00
  • 50714d2254 pragma unroll quick fix Andrej Karpathy 2024-05-02 21:06:11 +00:00
  • 9da2729c5a update encoder_forward with the latest and greatest Packed128 and incorporate into train_gpt2cu Andrej Karpathy 2024-05-02 20:40:05 +00:00
  • 41a0789a44 Added NVTX ranges for FlashAttention Peter Zhizhin 2024-05-02 20:30:08 +00:00
  • 8a9510f32b very minor stylistic preference let's keep the code style consistent Andrej Karpathy 2024-05-02 19:48:56 +00:00
  • 84dbc70160 Merge branch 'JaneIllario-patch-1' Andrej Karpathy 2024-05-02 19:45:18 +00:00
  • 3d1761b370 Merge branch 'patch-1' of https://github.com/JaneIllario/llm.c into JaneIllario-patch-1 Andrej Karpathy 2024-05-02 19:43:14 +00:00
  • 01df41fbba careful to guard the overloaded warpReduceSum and small fixes Andrej Karpathy 2024-05-02 19:24:04 +00:00
  • 6879c36cf3 Merge branch 'JaneIllario-gelu_backwards-3' Andrej Karpathy 2024-05-02 19:16:38 +00:00
  • 99f51bae1c Merge branch 'gelu_backwards' of https://github.com/JaneIllario/llm.c into JaneIllario-gelu_backwards-3 Andrej Karpathy 2024-05-02 19:15:06 +00:00
  • 8ccd05ec29 group warp reduce ops together Andrej Karpathy 2024-05-02 19:09:53 +00:00
  • d05c51c7d1 Merge branch 'ChrisDryden-removing_cg_example' Andrej Karpathy 2024-05-02 19:08:26 +00:00
  • 2583feb7a2 Merge pull request #327 from yushengsu-thu/master Andrej 2024-05-02 12:03:46 -07:00
  • 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 Andrej Karpathy 2024-05-02 19:01:26 +00:00
  • c13a730149 Merge branch 'master' into llm.java Harry Jackson 2024-05-02 11:44:44 -07:00
  • ec73086a83 Merge branch 'removing_cg_example' of https://github.com/ChrisDryden/llm.c into ChrisDryden-removing_cg_example Andrej Karpathy 2024-05-02 18:14:14 +00:00
  • 2202c9a51b add kernel 4 to docs. have to improve these docs more and document them better Andrej Karpathy 2024-05-02 18:09:27 +00:00
  • 2feb9ffcfb Fused classifier: use Packed128 and delete cg dep Andrej Karpathy 2024-05-02 18:08:27 +00:00
  • bcf7d4fe93 Add NSight Compute ranges, use CUDA events for timings Peter Zhizhin 2024-04-28 00:12:31 +00:00
  • 31e59273e1 Merge branch 'karpathy:master' into master Yusheng (Ethan) Su 2024-05-02 14:04:47 +04:00
  • 632caf1ce6 Refactored zero_active var chinthysl 2024-05-02 07:38:49 +00:00
  • d35d2e606b Generalized copy_and_cast_kernel and changes to cater model->master_weights chinthysl 2024-05-02 07:31:34 +00:00
  • 07b8102e79 Merge branch 'master' into zero-stage1 Chinthaka 2024-05-02 11:48:56 +08:00
  • 398cdaf4cd Removed cooperative groups and added packed128 to fused classifier and prepare softmax blockwide “Christopher 2024-05-02 02:33:51 +00:00
  • 665d0a4f89 Added example of removing cooperative groups “Christopher 2024-05-02 02:09:57 +00:00
  • d6d0d50627 fix typo Jane Illarionova 2024-05-02 00:18:55 +00:00
  • d0e5fea2dd update residual_forward to use util functions Jane Illarionova 2024-05-02 00:17:41 +00:00
  • 7203c7875f remove int cast Jane Illarionova 2024-05-02 00:14:41 +00:00
  • 4698164e3e update train to multiply block size Jane Illarionova 2024-05-01 04:49:23 +00:00
  • 52b4dd02a1 update residual forward to allow kernels to use both types Jane Illarionova 2024-05-01 04:47:57 +00:00
  • 32df703a05 Update train_gpt2.cu Jane Illarionova 2024-04-29 20:02:25 -07:00
  • 4ffcf5b90a Update residual_forward.cu Jane Illarionova 2024-04-29 19:58:16 -07:00
  • 3de3c53ba9 update kernel with util functions Jane Illarionova 2024-05-01 23:58:49 +00:00
  • d98e5ae859 remove int casting Jane Illarionova 2024-05-01 23:36:55 +00:00
  • e9a80b5d84 update ceildiv for gelu_backward Jane Illarionova 2024-05-01 04:55:18 +00:00
  • 7746217433 update gelu backward rto allow all kernels to use both types Jane Illarionova 2024-05-01 04:26:19 +00:00
  • ab2de05a13 Update train_gpt2.cu Jane Illarionova 2024-04-29 21:49:32 -07:00
  • 80072802d7 Create gelu_backward.cu Jane Illarionova 2024-04-29 21:46:35 -07:00
  • 4dd1ab43de Merge pull request #328 from karpathy/feature/fp32_weight_master_copy Andrej 2024-05-01 16:07:19 -07:00
  • c177c2694c make nice and print feature/fp32_weight_master_copy Andrej Karpathy 2024-05-01 23:06:23 +00:00
  • 795f8b690c fixes to keep master copy in fp32 of weights optionally Andrej Karpathy 2024-05-01 22:50:22 +00:00
  • 24abcbf6d2 Merge branch 'karpathy:master' into master Yusheng (Ethan) Su 2024-05-02 02:19:16 +04:00
  • ac2d635007 add required package (requests) into requirements.txt Yusheng Su 2024-05-02 02:11:38 +04:00
  • fd474feec8 Merge branch 'fp32-weights' of https://github.com/ngc92/llm.c Andrej Karpathy 2024-05-01 22:02:55 +00:00
  • df9f0efde6 updated letter and properly free buffer Erik Schultheis 2024-05-02 01:00:36 +03:00
  • f4f7a987e7 Merge pull request #324 from eymay/init_weights_from_scratch Andrej 2024-05-01 14:50:39 -07:00
  • e9c76a1753 option to keep weights as fp32 Erik Schultheis 2024-05-02 00:11:07 +03:00
  • ab95a11190 Merge pull request #325 from ngc92/dev-cuda-utils Andrej 2024-05-01 12:54:54 -07:00
  • b42db706b2 this kernel will be released later, sorry for the spoiler :) Erik Schultheis 2024-05-01 22:47:24 +03:00
  • 91bc72d0c0 fixup Erik Schultheis 2024-05-01 22:40:26 +03:00
  • 06980b6b2f mixed precision utilities for dev/cuda Erik Schultheis 2024-05-01 22:34:49 +03:00
  • d37639a8cb Merge pull request #316 from rosslwheeler/designated_init_fix Andrej 2024-05-01 12:13:12 -07:00
  • be5cee21e6 Merge pull request #317 from ChrisDryden/patch-5 Andrej 2024-05-01 12:00:01 -07:00
  • 38b478d8f9 Merge pull request #323 from karpathy/feature/cudnn Andrej 2024-05-01 11:34:16 -07:00
  • 7d0042d574 make training data configurable and also dont move the engire dataset to device that is not a good idea feature/cudnn Andrej Karpathy 2024-05-01 18:15:10 +00:00
  • 8aa9a9ed53 and also profile lol Andrej Karpathy 2024-05-01 17:56:01 +00:00
  • 1ed9381fd0 fix cudnnDestroy outside of ifdef Andrej Karpathy 2024-05-01 17:55:36 +00:00
  • 69d40ebfff add readme and search for cudd in home dir Andrej Karpathy 2024-05-01 17:43:50 +00:00
  • 35640651b3 updates to makefile: docs, override, etc Andrej Karpathy 2024-05-01 17:32:52 +00:00
  • 04718fbbd1 round 1 of small fixes for cudnn support Andrej Karpathy 2024-05-01 17:07:10 +00:00
  • 1147983b8c Add missing cudnnDestroy() calls ademeure 2024-05-01 17:19:51 +01:00
  • d9e7a0a378 Fixed profile_gpt2.cu/test_gpt2.cu for cuDNN ademeure 2024-05-01 17:03:28 +01:00
  • 2b540b61da attention_forward.cu ~matches train_gpt2cu & allocate less memory for stats in "att" ademeure 2024-05-01 16:53:01 +01:00
  • be54842622 Rename lowp_float to floatX in /dev/cuda/attention_forward.cu ademeure 2024-05-01 15:43:28 +01:00
  • c778dc0098 slightly increase wte threshold to make test_gpt2.cu pass (?) ademeure 2024-05-01 15:40:30 +01:00
  • 57c09529d9 Merge branch 'cudnn_try2' of https://github.com/ademeure/llm.c into cudnn_try2 ademeure 2024-05-01 15:32:13 +01:00
  • 7beff72f5f Add makefile changes for cuDNN ademeure 2024-05-01 15:31:30 +01:00
  • c4ecc04dc3 Merge remote-tracking branch 'karpathy/master' into cudnn_try2 ademeure 2024-05-01 15:30:00 +01:00
  • d508463ad7 Merge remote-tracking branch 'origin/cudnn_try2' into cudnn_try2 ademeure 2024-05-01 15:27:07 +01:00
  • d42ee0d925 Cleanup + do not allocate gradient preatt memory with cuDNN + allocate as small a cuDNN workspace as possible ademeure 2024-05-01 15:26:25 +01:00
  • 61ad4b15ab Fully working forward+backward cuDNN BF16 Flash Attention in train_gpt2.cu! ademeure 2024-05-01 15:02:47 +01:00
  • 662f7bf545 gradient accumulation preview, doesn't work yet but i think it is 95% of the way there or so Andrej Karpathy 2024-05-01 05:41:39 +00:00
  • 1806062a62 Removed detail about layernorm inputs being f32, no longer the case Chris Dryden 2024-04-30 22:01:14 -07:00