cuda: Add Q5_1, Q5_0, Q4_1 and Q4_0 to F32 conversion support. (#10976) #12000
Conversation
This file contains hidden or 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
github-actions
bot
added
Nvidia GPU
ggml/src/ggml-cuda/cpy.cu
Outdated
Comment on lines
435
to
488
| static void ggml_cpy_q5_1_f32_cuda( | ||
| const char * cx, char * cdst, const int ne, | ||
| const int ne00, const int ne01, const int ne02, | ||
| const int nb00, const int nb01, const int nb02, | ||
| const int nb03, const int ne10, const int ne11, const int ne12, | ||
| const int nb10, const int nb11, const int nb12, const int nb13, | ||
| cudaStream_t stream) { | ||
| const int num_blocks = ne; | ||
| cpy_q_f32<cpy_blck_q_f32<dequantize_q5_1, QK5_1>, QK5_1><<<num_blocks, 1, 0, stream>>>( | ||
| cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, | ||
| ne10, ne11, ne12, nb10, nb11, nb12, nb13); | ||
| } | ||
|
|
||
| static void ggml_cpy_q5_0_f32_cuda( | ||
| const char * cx, char * cdst, const int ne, | ||
| const int ne00, const int ne01, const int ne02, | ||
| const int nb00, const int nb01, const int nb02, | ||
| const int nb03, const int ne10, const int ne11, const int ne12, | ||
| const int nb10, const int nb11, const int nb12, const int nb13, | ||
| cudaStream_t stream) { | ||
| const int num_blocks = ne; | ||
| cpy_q_f32<cpy_blck_q_f32<dequantize_q5_0, QK5_0>, QK5_0><<<num_blocks, 1, 0, stream>>>( | ||
| cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, | ||
| ne10, ne11, ne12, nb10, nb11, nb12, nb13); | ||
| } | ||
|
|
||
| static void ggml_cpy_q4_1_f32_cuda( | ||
| const char * cx, char * cdst, const int ne, | ||
| const int ne00, const int ne01, const int ne02, | ||
| const int nb00, const int nb01, const int nb02, | ||
| const int nb03, const int ne10, const int ne11, const int ne12, | ||
| const int nb10, const int nb11, const int nb12, const int nb13, | ||
| cudaStream_t stream) { | ||
| const int num_blocks = ne; | ||
| cpy_q_f32<cpy_blck_q_f32<dequantize_q4_1, QK4_1>, QK4_1><<<num_blocks, 1, 0, stream>>>( | ||
| cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, | ||
| ne10, ne11, ne12, nb10, nb11, nb12, nb13); | ||
| } | ||
|
|
||
| static void ggml_cpy_q4_0_f32_cuda( | ||
| const char * cx, char * cdst, const int ne, | ||
| const int ne00, const int ne01, const int ne02, | ||
| const int nb00, const int nb01, const int nb02, | ||
| const int nb03, const int ne10, const int ne11, const int ne12, | ||
| const int nb10, const int nb11, const int nb12, const int nb13, | ||
| cudaStream_t stream) { | ||
| const int num_blocks = ne; | ||
| cpy_q_f32<cpy_blck_q_f32<dequantize_q4_0, QK4_0>, QK4_0><<<num_blocks, 1, 0, stream>>>( | ||
| cx, cdst, ne, ne00, ne01, ne02, nb00, nb01, nb02, nb03, | ||
| ne10, ne11, ne12, nb10, nb11, nb12, nb13); | ||
| } | ||
|
|
There was a problem hiding this comment.
Please keep the order of quants consistent. The order that I usually use for CUDA code is q4_0, q4_1, q5_0, q5_1, q8_0.
There was a problem hiding this comment.
Some of the existing code doesn't respect this order. I think it's better to not clean that (the existing code) up in the same patch though, as it would just add noise for reviewing. Can be done in a follow-up if you want.
|
@JohannesGaessler Incorporated your comments, plus reorder the newly added functions as requested. The order in the file is not totally consistent now, but as said above I think it's better to address that in a follow-up. |
JohannesGaessler
approved these changes
Feb 22, 2025
orca-zhang
pushed a commit
to orca-zhang/llama.cpp
that referenced
this pull request
Feb 26, 2025
arthw
pushed a commit
to arthw/llama.cpp
that referenced
this pull request
Feb 26, 2025
mglambda
pushed a commit
to mglambda/llama.cpp
that referenced
this pull request
Mar 8, 2025
mostlyuseful
pushed a commit
to mostlyuseful/llama.cpp
that referenced
this pull request
May 12, 2025
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.
Using templates and reusing the
dequant_qX_Yfunctions.