Skip to content

ggml-cuda: HIP: use __builtin_amdgcn_perm for q1_0 unpacking - #29927

Merged
JohannesGaessler merged 1 commit into
ggml-org:masterfrom
thelittlefireman:optimize-hip-q1-amdgcn-perm
Oct 17, 2026
Merged

JohannesGaessler merged 1 commit into
ggml-org:masterfrom
thelittlefireman:optimize-hip-q1-amdgcn-perm

Conversation

@thelittlefireman

Copy link
Copy Markdown
Contributor

Overview

Clean version of #28398. This PR handle both function vec_dot_q1_0_q8_1 and ggml_cuda_mmq_load_tiles_q1_0 by replacing __byte_perm by __builtin_amdgcn_perm for performance.

Additional information

Test master __builtin_amdgcn_perm Gain Speedup
pp2048 2512.65 t/s 3570.25 t/s +42.09% 1.421×
tg128 66.99 t/s 128.96 t/s +92.51% 1.925×

AMD Instinct MI50 (gfx906), m=4096, k=14336.

n master PR Speedup
1 0.263 TFLOPS 2.26 TFLOPS 8.59x
2 0.448 TFLOPS 4.06 TFLOPS 9.07x
3 0.402 TFLOPS 2.51 TFLOPS 6.24x
4 0.475 TFLOPS 4.71 TFLOPS 9.92x
5 0.565 TFLOPS 5.70 TFLOPS 10.08x
8 0.560 TFLOPS 5.96 TFLOPS 10.65x
512 6.94 TFLOPS 10.81 TFLOPS 1.56x

Requirements

@github-actions github-actions Bot added ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Oct 3, 2026
@thelittlefireman

thelittlefireman commented Oct 3, 2026 •

Copy link
Copy Markdown
Contributor Author

Marked as DRAFT as i already have too many open PR, but it's ready for a review. @IMbackK

@JohannesGaessler

Copy link
Copy Markdown
Contributor

@ggerganov this developer is making good contributions. I would like to retain them and have been thinking about inviting them to take a more active role in the project anyways. Can you make them a member? For context, their earliest PR was Aug 28.

@ggerganov

Copy link
Copy Markdown
Member

@ggerganov this developer is making good contributions. I would like to retain them and have been thinking about inviting them to take a more active role in the project anyways. Can you make them a member? For context, their earliest PR was Aug 28.

I just sent them an invite to the @ggml-org/ggml-cuda team.

@thelittlefireman
thelittlefireman force-pushed the optimize-hip-q1-amdgcn-perm branch from b6c4ff7 to efb7318 Compare October 15, 2026 20:02
@thelittlefireman
thelittlefireman force-pushed the optimize-hip-q1-amdgcn-perm branch from efb7318 to d6f0bd9 Compare October 17, 2026 13:29
@thelittlefireman

Copy link
Copy Markdown
Contributor Author

@JohannesGaessler @IMbackK Can we merge this one ?

@JohannesGaessler
JohannesGaessler merged commit 9d1878a into ggml-org:master Oct 17, 2026
14 of 15 checks passed
edwardyoon pushed a commit to edwardyoon/focus-llama that referenced this pull request Oct 24, 2026
auRose94 pushed a commit to auRose94/llama.cpp that referenced this pull request Oct 28, 2026
@BrewTestBot BrewTestBot mentioned this pull request Oct 28, 2026
1 task done
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CUDA Related to the CUDA backend ggml changes relating to the ggml tensor library for machine learning

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants