Feature Request: Way to use on Tesla P40

[narikm]

Prerequisites

• I am running the latest code. Mention the version if possible as well.
• I carefully followed the README.md.
• I searched using keywords relevant to my issue to make sure that I am creating a new issue that is not already open (or closed).
• I reviewed the Discussions, and have a new and useful enhancement to share.

Feature Description

Is there a way to use this code on legacy Tesla P40?

Motivation

Is there a way to use this repo on a old Tesla P40? I tried to deactivate flash attention and use:
cmake -B build ^
-DGGML_CUDA=ON ^
-DGGML_BLAS=OFF ^
-DGGML_CUDA_ARCH=61 ^
-DGGML_CUDA_GRAPH=OFF ^
-DGGML_CUDA_FORCE_MMQ=OFF ^
-DGGML_CUDA_DMMV_X=32 ^
-DGGML_CUDA_MMQ_ENABLE=OFF ^

according to Chat GPT, is there a way to compile it for old devices?

I only get CUDA errors like:
CUDA error: an illegal memory access was encountered
current device: 0, in function launch_mul_mat_q at D:\ik_llama.cpp\ggml\src\ggml-cuda\template-instances../mmq.cuh:4008
cudaFuncSetAttribute(mul_mat_q<type, mmq_x, 8, false>, cudaFuncAttributeMaxDynamicSharedMemorySize, shmem)
D:\ik_llama.cpp\ggml\src\ggml-cuda.cu:110: CUDA error

Possible Implementation

No response

[firecoperana]

P40 should work. Try this command:
-DGGML_CUDA=ON ^
-DGGML_BLAS=OFF ^
-DCMAKE_CUDA_ARCHITECTURES="61" ^
-DGGML_CUDA_USE_GRAPHS=OFF ^
-DGGML_CUDA_FORCE_MMQ=OFF ^
-DGGML_CUDA_DMMV_X=32 ^
-DGGML_CUDA_MMQ_ENABLE=OFF

Flash attention should work too with -DGGML_CUDA_FA_ALL_QUANTS=ON

[narikm]

Same error.
INFO [ update_slots] kv cache rm [p0, end) | tid="5232" timestamp=1753453942 id_slot=0 id_task=0 p0=0
CUDA error: an illegal memory access was encountered
current device: 0, in function launch_mul_mat_q at G:\ik_llama.cpp\ggml\src\ggml-cuda\template-instances../mmq.cuh:4008
cudaFuncSetAttribute(mul_mat_q<type, mmq_x, 8, false>, cudaFuncAttributeMaxDynamicSharedMemorySize, shmem)
G:\ik_llama.cpp\ggml\src\ggml-cuda.cu:110: CUDA error

The compiler said -DGGML_CUDA_MMQ_ENABLE=OFF was not used. Is this related?

[firecoperana]

Can you post your full command line including the model name?

[narikm]

The launch command line:
cd /d G:\ik_llama.cpp\build\bin\Release

llama-server ^
--alias DeepSeek-R1-0528-IQ2_K_R4 ^
--model "G:\DeepSeek-R1-0528-IQ2_K_R4\DeepSeek-R1-0528-IQ2_K_R4-00001-of-00005.gguf" ^
-rtr ^
--ctx-size 12288 ^
-ctk q8_0 ^
-mla 2 ^
-amb 512 ^
-fmoe ^
--n-gpu-layers 63 ^
--override-tensor exps=CPU ^
--parallel 1 ^
--threads 32 ^
--host 0.0.0.0 ^
--port 8008

The build line is what you gave me:
cmake -B build ^
-DGGML_CUDA=ON ^
-DGGML_BLAS=OFF ^
-DCMAKE_CUDA_ARCHITECTURES="61" ^
-DGGML_CUDA_USE_GRAPHS=OFF ^
-DGGML_CUDA_FORCE_MMQ=OFF ^
-DGGML_CUDA_DMMV_X=32 ^
-DGGML_CUDA_MMQ_ENABLE=OFF

[firecoperana]

You are using the ik quants file. Add -DGGML_IQK_FA_ALL_QUANTS=1 to see if it works. You can also use regular quants here.

[narikm]

Will do. Do i need -fa in the launch ? The compiler said:

Manually-specified variables were not used by the project:

GGML_CUDA_MMQ_ENABLE.

Is this OK?

[firecoperana]

-fa is supported, but since you are not using ctv, it's not required I think. I would remove GGML_CUDA_MMQ_ENABLE since it's off by default. If you still have issue, you can remove -ctk q8_0.

[firecoperana]

You can add -DGGML_CUDA_FA_ALL_QUANTS=ON to make sure all quants are supported.

[ikawrakow]

DeepSeek flash attention does not work on a P40.

[narikm]

ChatGPT said the same, so i removed the -fa in the launch args before posting. It changed the error to the one i posted. I still get the same error each time. It load the model but when i ask something, error, the cpu run for ten seconds, then the program crash.

[firecoperana]

It most likely not gonna work, but if you like to try, can you remove line 4008-4009 or 4008-4010 in mmq.cuh, and run it without fa?

[Ph0rk0z]

DeepSeek flash attention does not work on a P40.

Is this because of having to use BF16? Or something else like not enough smem?

-DGGML_CUDA_FA_ALL_QUANTS=ON

I think this is just to use different K and V quantization. Like Q4/Q8.

This thread saved me from re-installing the P40s I have at least. Any other caveats, say for qwen coder or others? Assume P100 is the same story too. What about turning cards?

[firecoperana]

ChatGPT said the same, so i removed the -fa in the launch args before posting. It changed the error to the one i posted. I still get the same error each time. It load the model but when i ask something, error, the cpu run for ten seconds, then the program crash.

Does your card work in llama.cpp if you use regular quants of Deepseek R1?