GPU-accelerated token generation (new quantization format) #1412

JohannesGaessler · 2023-05-12T09:49:18Z

Build instructions (Linux):

git clone https://github.com/JohannesGaessler/llama.cpp llama.cpp-johannesgaessler
cd llama.cpp-johannesgaessler                               
git fetch
git switch dequantize-matmul-4
make LLAMA_CUBLAS=1

For building on Windows, read the llama.cpp README.

This PR is a replacement for #1375 that works with the new quantization format. Read that PR for more information. People with no git experience are already using that version so I'm making a new branch and PR to avoid unnecessary confusion, especially with the breaking quantization changes.

The goals of this PR:

Provide an implementation of GPU-accelerated matrix vector multiplication for all quantization types.
Provide an implementation that works on as broad of a range of hardware as possible.
Load weights directly into VRAM to reduce RAM usage.

Not the goals of this PR:

Squeeze out every last % of performance. It seems that GPU performance optimizations like varying block sizes strongly depend on the specific model of GPU so optimizing that seems like it will take a long time and require feedback from many people. I would like to do that at a later time.
Accelerating operations other than matrix vector multiplications and avoiding the copying of vectors between CPU and GPU.
Multi-GPU support. As of right now I don't have a machine that I could test this on.
iGPU support. Again, I don't have a machine to test this on.

In other news, the quantization changes make a big difference for my kernel implementation. I can now get 14.53 t/s with a GTX 1070 for 7b which is 16% faster than with the old quantization method. I think the reason is memory coalescing when reading the vector values.

JohannesGaessler · 2023-05-12T11:27:48Z

I implemented q4_1 by making a dequantize_mul_mat_vec kernel template that takes accepts another kernel to determine how the dequantization should be done. I'm aware that the README tells you not to use templates but I think that this is the most elegant solution. It allows you to write the matrix vector multiplication code only once; you only need to provide different kernels for dequantization. The dequantization kernels can then also be reused to dequantize matrices for cuBLAS prompt processing. @ggerganov please provide feedback regarding the software design.

Folko-Ven · 2023-05-12T12:57:43Z

Hello, thank you very much for your contribution to speed up llama.cpp As I understand the acceleration now only works with nvidia gpu? I would like to ask if there is any way to keep the option of not copying weights to vram, for those of us who use igpu? New igpus are quite powerful, for example the amd 780M is more powerful than the 1650. There are also rumors that Intel Meteor Lake will be even more powerful.

JohannesGaessler · 2023-05-12T13:08:38Z

I unfortunately don't have a machine on which I could test an iGPU implementation. I was thinking I would implement the kernels for discrete GPUs first and then someone else could make a follow-up PR that supports iGPUs.

JohannesGaessler · 2023-05-12T13:09:49Z

@SlyEcho somehow managed to run the previous version that I implemented on AMD, but I don't know how he did it and I don't have an AMD GPU to test my implementation on.

SlyEcho · 2023-05-12T13:34:55Z

You can just start from the master branch here:

git pull origin master

# apply the ROCm PR:
curl -L https://github.com/ggerganov/llama.cpp/pull/1087.diff | git apply -

# apply this PR:
curl -L https://github.com/ggerganov/llama.cpp/pull/1412.diff | git apply -

Then you can follow the instructions in #1087.

If there is a mess and you want to get back to the start then just git reset --hard HEAD should put you back.

EDIT: I put some Docker instructions here: rocm.Dockerfile, with containers there is no need to install any drivers or SDKs or special compilers.

ggml-cuda.cu

JohannesGaessler · 2023-05-12T19:28:31Z

I've tried looking into how memory management is done in llama.cpp but it's more tricky than I thought. At least I think that when using --mmap the current version is not too bad: after the initial loading the unused parts of the file should be automatically cleared from memory although the GPU-accelerated parts of the model will need to be loaded into memory twice if the amount of RAM is smaller than the model size.

However, when using --mlock the entire input file is locked into memory, including the parts that are offloaded to the GPU anyways. This particular problem can be fixed by not memlocking the first part of the file until the CPU weights start being loaded; the weights for e.g. normalization in the GPU layers are relatively small and will probably not be too relevant for performance. Overall I feel like the memory management in llama.cpp is relatively complicated; maybe this should be postponed to a later date? Can the people that wrote the memory management code perhaps give some insight?

github-actions

clang-tidy made some suggestions

github-actions · 2023-05-12T19:41:47Z

llama.cpp

@@ -879,6 +883,7 @@ static void llama_model_load_internal(
        ggml_type memory_type,
        bool use_mmap,
        bool use_mlock,
+        int gpu_layers,


warning: unused parameter 'gpu_layers' [clang-diagnostic-unused-parameter]

int gpu_layers, ^

SlyEcho · 2023-05-12T19:48:15Z

The layers that are loaded to the GPU could be unmapped from the memory map, although, I don't think they are lying on page boundaries, but since the loading is sequential it may be possible to at least partially unmap them?

JohannesGaessler · 2023-05-12T20:00:50Z

My current plan: I'll prioritize cleaning up the code and finishing the kernels for the various quantization types. After that as far as I'm concerned I think this can be merged (unless someone wants to suggest changes to the software design?). I'll probably be done tomorrow morning. I'll make follow-up PRs to fix things like the inefficient use of memory.

SlyEcho · 2023-05-12T20:39:16Z

ggml-cuda.cu

@@ -173,6 +206,41 @@ static __global__ void dequantize_block_q8_0(const void * vx, float * y) {
    }
 }

+template <int block_size, int qk, dequantize_kernel_t dequantize_kernel> static __global__ void dequantize_mul_mat_vec(const void * vx, const float * y, float * dst, const int ncols) {


Could the template<...> be on a separate line? It's getting very long otherwise.

ggerganov · 2023-05-12T20:45:54Z

@JohannesGaessler

Thank you for these contributions

In past discussions, I've been strongly opposed to changing ggml.h interface with GPU-specific elements, like the backend member in this PR. But given the huge inference speed-up, I am starting to change my opinion 😄
I'm OK to merge this.

I wish to hear @slaren's opinion as he had worked on most of the CUDA backend and he had ideas for offloading the tensors to the GPU. I'm open to alternative implementations that keep the ggml.h interface changes minimal and potentially can offer additional perplexity computation speed-up as well.

slaren · 2023-05-12T21:30:19Z

ggml.c

+    if (ggml_cuda_can_mul_mat(src0, src1, dst) ||
+        ggml_cuda_can_dequantize_mul_mat_vec(src0, src1, dst) ||
+        src0->backend == GGML_BACKEND_CUDA) {


Is there any reason to put this here instead of leaving all the checks in ggml_cuda_can_mul_mat?

I was initially thinking it would be better this way: to have a separate method for dequantize mul mat since you can then also use that method for the logic inside ggml_cuda_mul_mat_q_f32. I was thinking I would just push it and see what other people have to say since it turned out kind of convoluted.

I think that it is better to keep the complexity in ggml_cuda_can_mul_mat than to spread it all around, but I am not entirely sure if you have any reason to split the check in two parts.

slaren · 2023-05-12T21:31:07Z

I think this is good if it can bring a performance improvement right now, the design will need changes if we want to go further and keep as much of the computation in the GPU as possible, but we can iterate over it.

Green-Sky · 2023-05-12T22:19:45Z

Load weights directly into VRAM to reduce RAM usage.

if this is done correctly, I can now run 65B size q4_0 models o.O

32gig ram + 8gig vram

SlyEcho · 2023-05-12T22:34:34Z

This PR is pretty neat and makes my generation speed ~~double~~ three times as fast (for 7B at least).

It really makes my GPU scream. Literally! I should make a video...

JohannesGaessler · 2023-05-13T05:00:00Z

The GPU "screaming" is a sign of inefficiency though. Currently the GPU has to briefly stop between layers because parts of the model like the norms are still CPU only, so the GPU is rapidly turning on and off. The change in current induces Lorentz forces that cause vibrations, i.e. sounds. Theoretically the vibrations could damage the GPU via resonance catastrophe if you hit the eigenfrequencies of the components but I don't think that this is a realistic problem.

Green-Sky · 2023-05-13T11:46:23Z

I also get coil whine with my rtx 2070 mobile 😄
edit: so the smaller the quantization, the faster, the more coil whine. also with q4_0 i observe almost full powerdraw with only ~60% util

github-actions

clang-tidy made some suggestions

llama.cpp

slaren · 2023-05-13T14:12:53Z

ggml-cuda.cu

+        case GGML_TYPE_Q8_0:
+            return dequantize_mul_mat_vec_q8_0_cuda;
+        case GGML_TYPE_F16:
+            return dequantize_mul_mat_vec_q8_0_cuda;


Shouldn't this be convert_mul_mat_vec_f16_cuda instead?

I think so - just pushed a fix, although this branch is never reached

JohannesGaessler · 2023-05-13T15:10:09Z

My 3090 has arrived in the mail. I'll still need to run more detailed performance tests but it's looking good so far: 43.7 t/s for q4_0 7b and 26.6 t/s for q4_0 13b. Thanks everyone for the feedback and help.

kevkid · 2023-05-13T16:14:53Z

Has anyone got this to build in windows? cant seem to get it to build. link to discussion: #1431

Green-Sky · 2023-05-13T16:35:38Z

@kevkid Check the CI. or grab it from the latest release. https://github.com/ggerganov/llama.cpp/releases/tag/master-bda4d7c

kevkid · 2023-05-13T16:54:43Z

@Green-Sky I got it to build thank you for replying. it compiled, but is there a good way to ensure its using gpu? gpu usage never goes above 10%.
Edit:

but I am getting:

llama_model_load_internal: [cublas] offloading 0 layers to GPU
llama_model_load_internal: [cublas] total VRAM used: 0 MB

moejay · 2023-05-13T17:23:44Z

@Green-Sky I got it to build thank you for replying. it compiled, but is there a good way to ensure its using gpu? gpu usage never goes above 10%. Edit: but I am getting:
llama_model_load_internal: [cublas] offloading 0 layers to GPU
llama_model_load_internal: [cublas] total VRAM used: 0 MB

I believe you need to add -ngl N, --n-gpu-layers argument

Green-Sky · 2023-05-13T17:52:20Z

@moejay is right, loading parts(or all) of the model to the gpu is opt-in.

kevkid · 2023-05-13T19:17:31Z

@Green-Sky I got it to build thank you for replying. it compiled, but is there a good way to ensure its using gpu? gpu usage never goes above 10%. Edit: but I am getting:
llama_model_load_internal: [cublas] offloading 0 layers to GPU
llama_model_load_internal: [cublas] total VRAM used: 0 MB
I believe you need to add -ngl N, --n-gpu-layers argument

Thank you.

Green-Sky · 2023-05-13T21:44:57Z

@JohannesGaessler or @ggerganov please add a hint to the

llama_model_load_internal: [cublas] offloading 0 layers to GPU
llama_model_load_internal: [cublas] total VRAM used: 0 MB

which CLI flag to set or lookup in the help. People have been complaining. :)

JohannesGaessler · 2023-05-13T21:49:10Z

On the other hand, you can always just use --help and you will be enlightened.

Green-Sky · 2023-05-13T21:54:11Z

sure, but we already had 2 ask for help here in the last 8h 😆

CRD716 · 2023-05-14T03:05:10Z

there should also be some documentation on how many layers is best. in my initial testing there seems to be a number of layers offloaded that's worse than none (in my case with 65B it's anything under ~15 layers, with my 13B tests it seems to be higher) before it starts getting faster.

regstuff · 2023-05-14T05:05:04Z

I've got an AMD GPU, and compiling with CLBlast flags enables me to use the GPU for prompt ingestion. Would CLBlast also enable me to use the GPU for generation?

JohannesGaessler · 2023-05-14T07:11:05Z

there should also be some documentation on how many layers is best.

Almost always: as many as you can fit into VRAM. I don't think that there is a feasibly way to determine specifics for all possible hardware configurations.

Would CLBlast also enable me to use the GPU for generation?

In theory yes, but no one has implemented OpenCL token generation.

zakkor · 2023-05-14T13:15:50Z

What GPUs & OSes is this currently compatible with?

Green-Sky · 2023-05-14T14:52:37Z

What GPUs & OSes is this currently compatible with?

Nvidia windows and linux.

JohannesGaessler mentioned this pull request May 12, 2023

CLI args use - instead of _, backwards compatible #1416

Merged

JohannesGaessler mentioned this pull request May 12, 2023

Proof of concept: GPU-accelerated token generation #1375

Closed

ggerganov reviewed May 12, 2023

View reviewed changes

ggml-cuda.cu Show resolved Hide resolved

ggerganov reviewed May 12, 2023

View reviewed changes

ggml-cuda.cu Show resolved Hide resolved

github-actions bot reviewed May 12, 2023

View reviewed changes

JohannesGaessler added 4 commits May 12, 2023 21:42

CUDA kernel for q4_0 dequant. + mat. vec. mult.

637be12

Added q4_1 via template

12fc292

Added missing __syncthreads();

7dc2f57

--gpu_layers -> --gpu-layers

f0af475

JohannesGaessler force-pushed the dequantize-matmul-4 branch from 1e735d2 to f0af475 Compare May 12, 2023 19:44

SlyEcho reviewed May 12, 2023

View reviewed changes

slaren reviewed May 12, 2023

View reviewed changes

JohannesGaessler added 4 commits May 13, 2023 07:03

Shorter dequantize_mul_mat_vec line

0986c2f

q5_0 dequantize_mul_mat kernel

9da44fd

More readable dequantize_mul_mat_vec logic

5a0ecf7

dequantize_mul_mat_vec kernels for q5_1, q8_0, f16

bb0993e

JohannesGaessler force-pushed the dequantize-matmul-4 branch from 991ef9e to bb0993e Compare May 13, 2023 07:01

llama : offload "output" tensor to GPU too + coding style fixes

ad8a9e6

ggerganov approved these changes May 13, 2023

View reviewed changes

ggerganov merged commit 905d87b into ggerganov:master May 13, 2023

github-actions bot reviewed May 13, 2023

View reviewed changes

llama.cpp Show resolved Hide resolved

slaren reviewed May 13, 2023

View reviewed changes

ggerganov added a commit that referenced this pull request May 13, 2023

cuda : fix convert function (#1412)

08737ef

JohannesGaessler mentioned this pull request May 13, 2023

More GPU threads for dequantization #1341

Closed

AshleySchaeffer mentioned this pull request May 14, 2023

Update to latest upstream LLaMA implementation rustformers/llm#210

Closed

JohannesGaessler mentioned this pull request May 14, 2023

Deduplicated dequantization code #1453

Merged

garrettsutula mentioned this pull request May 14, 2023

Add llamacpp n_gpu_layers langchain-ai/langchain#4677

Closed

AlphaAtlas mentioned this pull request May 14, 2023

Add llama.cpp GPU offload option oobabooga/text-generation-webui#2060

Merged

garrettsutula mentioned this pull request May 14, 2023

Add n_gpu_layers to llama.cpp langchain-ai/langchain#4679

Closed

mudler mentioned this pull request May 18, 2023

feat: add support for cublas/openblas in the llama.cpp backend mudler/LocalAI#258

Merged

edp1096 mentioned this pull request May 22, 2023

CUDA/OpenCL error, out of memory when reload. #1456

Closed

GPU-accelerated token generation (new quantization format) #1412

GPU-accelerated token generation (new quantization format) #1412

Conversation

JohannesGaessler commented May 12, 2023 • edited Loading

JohannesGaessler commented May 12, 2023 • edited Loading

Folko-Ven commented May 12, 2023

JohannesGaessler commented May 12, 2023

JohannesGaessler commented May 12, 2023

SlyEcho commented May 12, 2023 • edited Loading

JohannesGaessler commented May 12, 2023 • edited Loading

github-actions bot left a comment

Choose a reason for hiding this comment

github-actions bot May 12, 2023

Choose a reason for hiding this comment

SlyEcho commented May 12, 2023

JohannesGaessler commented May 12, 2023

SlyEcho May 12, 2023

Choose a reason for hiding this comment

ggerganov commented May 12, 2023

slaren May 12, 2023

Choose a reason for hiding this comment

JohannesGaessler May 12, 2023

Choose a reason for hiding this comment

slaren May 12, 2023

Choose a reason for hiding this comment

slaren commented May 12, 2023

Green-Sky commented May 12, 2023 • edited Loading

SlyEcho commented May 12, 2023 • edited Loading

JohannesGaessler commented May 13, 2023

Green-Sky commented May 13, 2023 • edited Loading

github-actions bot left a comment

Choose a reason for hiding this comment

slaren May 13, 2023

Choose a reason for hiding this comment

ggerganov May 13, 2023

Choose a reason for hiding this comment

JohannesGaessler commented May 13, 2023

kevkid commented May 13, 2023 • edited Loading

Green-Sky commented May 13, 2023

kevkid commented May 13, 2023 • edited Loading

moejay commented May 13, 2023

Green-Sky commented May 13, 2023

kevkid commented May 13, 2023

Green-Sky commented May 13, 2023

JohannesGaessler commented May 13, 2023

Green-Sky commented May 13, 2023

CRD716 commented May 14, 2023 • edited Loading

regstuff commented May 14, 2023

JohannesGaessler commented May 14, 2023

zakkor commented May 14, 2023

Green-Sky commented May 14, 2023

JohannesGaessler commented May 12, 2023 •

edited

Loading

JohannesGaessler commented May 12, 2023 •

edited

Loading

SlyEcho commented May 12, 2023 •

edited

Loading

JohannesGaessler commented May 12, 2023 •

edited

Loading

Green-Sky commented May 12, 2023 •

edited

Loading

SlyEcho commented May 12, 2023 •

edited

Loading

Green-Sky commented May 13, 2023 •

edited

Loading

kevkid commented May 13, 2023 •

edited

Loading

kevkid commented May 13, 2023 •

edited

Loading

CRD716 commented May 14, 2023 •

edited

Loading