Refactoring & Improvements to reduce LOC #2
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…ernel changes make it compile with cudnn default batchsize is 24 now
updated the profile script to be more robust, and adapted to recent k…
Adding a java port of the project
ngc92
reviewed
May 4, 2024
| const int warp_size = 32; | ||
| const int block_size = 512; | ||
| const int OC_per_warp = warp_size * x128::size; // 256 at BF16 | ||
| const int block_size_x = 32; | ||
| const int block_size_y = block_size / block_size_x; // 16 | ||
| const int grid_size_x = OC / OC_per_warp; // e.g. 3 horizontal blocks for 768 OCs at BF16 | ||
| const int grid_size_y = max(1, cuda_threads_per_SM * cuda_num_SMs / (block_size * grid_size_x)); // full GPU! | ||
| const int grid_size_y = max(1, deviceProp.maxThreadsPerMultiProcessor * deviceProp.multiProcessorCount |
There was a problem hiding this comment.
+1 for just storing the entire deviceProp
ngc92
reviewed
May 4, 2024
| @@ -1636,29 +1504,28 @@ void matmul_backward(floatX* dinp, floatX* dweight, floatX* dbias, | |||
| dim3(block_size_x, block_size_y), | |||
| OC_per_warp * sizeof(float), main_stream>>>(dbias_buffer, dout, B, T, OC); | |||
| cast_and_add_kernel<<<CEIL_DIV(OC, 256), 256, 0, main_stream>>>(dbias, dbias_buffer, OC); | |||
| cudaCheck(cudaGetLastError()); | |||
There was a problem hiding this comment.
I figured it'd be more consistent to always have a single "cudaCheck(cudaGetLastError())" at the end of every function.
That should be more than enough for the initial debug of where the problem comes from without cluttering e.g. attention_forward/backward with 3 error checks each.
Remove the duplicated #include <assert.h>
… into ademeure-less_idle_more_brrr-3
…into ngc92-separate-compilation
…h feels slightly cleaner
…work on older CUDAs) + fixes
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Refactoring and removing unused functions to reduce the number of lines of code and make everything slightly more consistent (while still having space for the code to breathe).
Also update encoder_backward with my version from more_stochastic branch in order to delete atomicAddX() from the codebase, while improving accuracy very slightly by using stochastic rounding (it's literally the only place in the entire code where we are not accumulating in FP32!)
This is based off my other PR: karpathy#343 - assuming everyone likes these changes, I will try to merge this back with latest version of the main branch at some point next week once that PR has been integrated, then create a new PR.