CUDA: faster softmax via shared memory + fp16 math

[JohannesGaessler]

This PR adds performance optimizations for CUDA softmax calculation. According to profiling data the softmax kernel has become 2.4 times faster. However, because the softmax kernel only takes up a small percentage of the runtime the overall impact on performance is small:

Model	GPU	test	t/s master	t/s PR	Speedup
llama 7B Q4_0	NVIDIA GeForce RTX 3090	pp512	3462.88	3527.40	1.02
llama 7B Q4_0	NVIDIA GeForce RTX 3090	tg128	136.90	136.73	1.00
llama 7B Q4_0	Tesla P40	pp512	894.99	908.75	1.02
llama 7B Q4_0	Tesla P40	tg128	56.67	56.84	1.00
llama 7B Q4_0	AMD Radeon RX 6800	pp512	1241.66	1248.01	1.01
llama 7B Q4_0	AMD Radeon RX 6800	tg128	74.41	74.31	1.00

The performance optimizations include:

• Using shared memory to avoid loading the data multiple times from global memory.
• Template parameter to optionally inform the compiler of the number of columns and the block size. If both are specified the compiler can unroll loops, and avoid checks for correctness. Aside from a generic kernel that can handle any number of columns fast variants for powers of 2 are compiled. These are the sizes used for prompt processing where softmax (on master) is taking up a relatively high percentage of the runtime compared to token generation.
• On NVIDIA cards that are Volta or newer FP16 arithmetic instead of FP32 arithmetic is used. This is both faster and reduces the amount of shared memory required. Perplexity for LLaMA 2 q4_0 changes from $5.9634 \pm 0.03349$ to $5.9629 \pm 0.03349$. For AMD FP16 arithmetic is disabled because the HIP documentation does not seem to contain an equivalent for __hmax2.

[slaren]

test-backend-ops fails with this change:

$ tests/test-backend-ops -b CUDA0 -o SOFT_MAX                                                                                                            Testing 2 backends

Backend 1/2 (CPU)
  Skipping
Backend 2/2 (CUDA0)
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3090 Ti, compute capability 8.6, VMM: yes
  Backend name: CUDA
  SOFT_MAX(type=f32,ne=[10,10,10,10]): [SOFT_MAX] NMSE = 4.772921 sentinel mismatch: sent_2 FAIL
  876/877 tests passed
  Backend CUDA: FAIL

1/2 backends passed
**FAIL**

The sentinel mismatch indicates a buffer overflow.

We need to be careful about downcasting to F16 without the user explicitly requesting this, as some models are susceptible to precision issues. We have had some issues already with F16 matrix multiplication, with Phi2 it causes issue in the KV cache, and in whisper.cpp it causes degraded quality. I think that stable diffusion is also susceptible to precision issues. Currently we have a function to set the precision of matrix multiplications, but I think that we need a more general solution.

[JohannesGaessler]

The sentinel mismatch indicates a buffer overflow.

Thank you. I didn't notice with compute-sanitizer because this only seems to happen for batch sizes with % 32 != 1.

We need to be careful about downcasting to F16 without the user explicitly requesting this, as some models are susceptible to precision issues.

Good point. I changed the check to choose simply based on GGML_CUDA_F16.

More generally, what is the long-term plan for CUDA FP16 support in ggml? I think to remember that one of your unmerged PRs had more general FP16 support so I was under the impression that it is of at least moderate priority.

[slaren]

And at the moment my priority is to move the CUDA backend implementation fully to the ggml-backend interface, and then do some refactoring of the CUDA backend. Essentially what this means is that partial offloading will be handled entirely by ggml-backend with ggml_backend_sched, and most of the code that deals with copying data between the CPU and the GPU in the CUDA backend will be removed. Then we can finally remove most of the globals in ggml-cuda, make it thread safe, add better support for multiple devices and so on.

So adding support for FP16 is not a priority for me at the moment. I had a partial implementation in the first iteration of the ggml-backend interface, and some of that may still be useful today. The way it worked was that ggml_context had a parameter that allowed setting the type of the operations, in that way operations could return FP16 tensors instead of always returning FP32 as it does now. In the CUDA backend side, all kernels were made templates that could use any floating point type. I think there were some mistakes there, for instance there is no reason to use FP16 arithmetic unless using SIMD with half2, but overall I think the design could still work.

[JohannesGaessler]

for instance there is no reason to use FP16 arithmetic unless using SIMD with half2

I would say this depends on the computing task. Currently the focus is on inference where the static weights take up the most memory but e.g. for training neural networks you need a lot of memory to store gradients. In such scenarios FP16 would be useful simply to save space compared to FP32. But for the current state of the software I definitely agree.

[slaren]

Sorry, what I meant was just the arithmetic, my kernels in general used the type of the destination tensor as the type of the arithmetic. Storing the results as FP16 would be useful just to avoid the conversions to FP16 in the matrix multiplications, even if the arithmetic in the kernels is still FP32.

[kalomaze]

Are there any plans to move sampling logic / memory to the GPU as well?

As it currently stands, multiple sampler passes (for Top K, Temperature, Min P, repetition penalty, etc) are running on CPU without any threading. The softmax is just one operation that is currently happening on CPU; certain more complex samplers are also not trivially inexpensive on CPU due to large vocabulary sizes, especially the custom Grammar sampler (which I helped mitigate with the last PR that I submitted by avoiding grammar checks when the selected token adheres to the Grammar).

This is particularly troublesome for batched throughput use cases, less so for single user inference, but it could have more of an impact than previously believed maybe?

[JohannesGaessler]

I'm not overly familiar with the sampling code but if I remember correctly it has a lot of branching. GPUs are only faster than CPUs for code where you can apply the same instructions to multiple data streams at the same time. So I'm not sure that sampling on the GPU would even be any faster.

[slaren]

Btw, test-backend-ops also has a mode for testing performance, just add perf to the command line. I get these results:

New test cases:

test_cases.emplace_back(new test_soft_max(GGML_TYPE_F32, {4096, 1, 1, 1}));
test_cases.emplace_back(new test_soft_max(GGML_TYPE_F32, {1024, 512, 1, 1}));
test_cases.emplace_back(new test_soft_max(GGML_TYPE_F32, {4096, 512, 1, 1}));

master:

SOFT_MAX(type=f32,ne=[10,10,10,10]):                           8192 runs -     5.69 us/run -       78 kB/run -   13.09 GB/s
SOFT_MAX(type=f32,ne=[4096,1,1,1]):                            8192 runs -     7.15 us/run -       32 kB/run -    4.27 GB/s
SOFT_MAX(type=f32,ne=[1024,512,1,1]):                          8192 runs -    16.97 us/run -     4096 kB/run -  230.13 GB/s
SOFT_MAX(type=f32,ne=[4096,512,1,1]):                          2049 runs -    24.97 us/run -    16384 kB/run -  625.66 GB/s

PR:

SOFT_MAX(type=f32,ne=[10,10,10,10]):                           8192 runs -     6.04 us/run -       78 kB/run -   12.34 GB/s
SOFT_MAX(type=f32,ne=[4096,1,1,1]):                            8192 runs -     6.37 us/run -       32 kB/run -    4.79 GB/s
SOFT_MAX(type=f32,ne=[1024,512,1,1]):                          8192 runs -    12.82 us/run -     4096 kB/run -  304.65 GB/s
SOFT_MAX(type=f32,ne=[4096,512,1,1]):                          2049 runs -    23.56 us/run -    16384 kB/run -  663.12 GB/s

PR + LLAMA_CUDA_F16:

SOFT_MAX(type=f32,ne=[10,10,10,10]):                           8192 runs -     5.51 us/run -       78 kB/run -   13.52 GB/s
SOFT_MAX(type=f32,ne=[4096,1,1,1]):                            8192 runs -     6.18 us/run -       32 kB/run -    4.94 GB/s
SOFT_MAX(type=f32,ne=[1024,512,1,1]):                          8192 runs -     8.99 us/run -     4096 kB/run -  434.42 GB/s
SOFT_MAX(type=f32,ne=[4096,512,1,1]):                          2049 runs -    22.80 us/run -    16384 kB/run -  685.35 GB/s

Building with LLAMA_CUDA_F16 also makes the tests fail, which I guess it is expected, however the error in the first case seems too high:

  SOFT_MAX(type=f32,ne=[10,10,10,10]): [SOFT_MAX] NMSE = 4.806375 FAIL
  SOFT_MAX(type=f32,ne=[4096,1,1,1]): [SOFT_MAX] NMSE = 0.000000 FAIL
  SOFT_MAX(type=f32,ne=[1024,512,1,1]): [SOFT_MAX] NMSE = 0.000000 FAIL
  SOFT_MAX(type=f32,ne=[4096,512,1,1]): [SOFT_MAX] NMSE = 0.000000 FAIL

[JohannesGaessler]

however the error in the first case seems too high

Yes, it seems the checks for arbitrary tensor sizes were still not quite correct.

I've also committed the changes I made to test-backend-ops.cpp for debugging. The output now looks like this:

> ./tests/test-backend-ops -o SOFT_MAX                                
Testing 2 backends

Backend 1/2 (CPU)
  Backend name: CPU
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 CUDA devices:
  Device 0: NVIDIA GeForce RTX 3090, compute capability 8.6, VMM: yes
  SOFT_MAX(type=f32,ne=[10,10,10,10]): OK
  877/877 tests passed
  Backend CPU: OK

Backend 2/2 (CUDA0)
  Backend name: CUDA
  SOFT_MAX(type=f32,ne=[10,10,10,10]): [SOFT_MAX] NMSE = 2.362E-07 > 1.000E-07 FAIL
  876/877 tests passed
  Backend CUDA: FAIL

1/2 backends passed
FAIL

I think this is better because it tells you how large some metric is compared to some threshold. "NMSE" to me doesn't mean anything without knowing what the code does so I wouldn't know how good or bad a value of 4.8 is. But with a comparison to a threshold of 1e-7 it is clear that there is still a bug somewhere.

[slaren]

Two more issues:

• F16 fails with an odd number of columns
• Too many columns fails due to not enough shared memory. This is not probably not an issue in llama.cpp, but soft_max is used for sampling, and in that case the number of columns can be very large (the vocab size).

  SOFT_MAX(type=f32,ne=[10,10,10,10]): [SOFT_MAX] NMSE = 2.371E-07 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[11,10,10,10]): [SOFT_MAX] NMSE = 2.313E-07 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[4097,1,1,1]): [SOFT_MAX] NMSE = 9.993E-01 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[1025,512,1,1]): [SOFT_MAX] NMSE = 9.962E-01 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[4097,512,1,1]): [SOFT_MAX] NMSE = 9.991E-01 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[50000,1,1,1]): CUDA error: invalid argument

[JohannesGaessler]

By the way, the kernel sizes that you are using for performance testing are way to small. On Ampere each SM can fit 6 of the new kernels. An RTX 3090 can therefore process 492 rows per wave while an RTX 3090 ti can process 504 rows per wave. With 1 row 99.8% of the GPU is idling while with 512 rows the GPU only has an utilization of ~50% due to tail effects. During perplexity calculation for a 7b model with a batch size of 512 the tensor shape was [4096, 16384, 1, 1]. With this shape the evaluation on my RTX 3090 took on average 85 µs.

[slaren]

Larger batch sizes generally reduce the difference with master.

SOFT_MAX(type=f32,ne=[4096,16384,1,1]):                          65 runs -   630.09 us/run -   524288 kB/run -  793.53 GB/s

SOFT_MAX(type=f32,ne=[4096,16384,1,1]):                          65 runs -   596.89 us/run -   524288 kB/run -  837.67 GB/s

With this shape the evaluation on my RTX 3090 took on average 85 µs.

I don't think this can be right, it would imply a throughput in the thousands of GB/s. Am I wrong? (4096×16384×4×2 bytes) / 85us −> GiB/s = 5882 GiB/s

[FSSRepo]

@slaren Sometimes the current softmax kernel has race condition issues, causing the same prompt and seed to generate different results at times. This occurs in llama.cpp as well as in stable-diffusion.cpp.

compute-sanitizer --tool racecheck cuda program and args

Example:

compute-sanitizer --tool racecheck main -m models/llama-13b-v2/ggml-model-q4_0.gguf -p "Building a website can be done in 10 simple steps:\nStep 1:" -n 400 -e

For greater ease, it's better try with test-backend-ops since it supports all softmax cases.

It is solved by commenting out this line:

static void soft_max_f32_cuda(const float * x, const float * y, float * dst, const int ncols_x, const int nrows_x, const int nrows_y, const float scale, cudaStream_t stream) {
    int nth = WARP_SIZE;
    //while (nth < ncols_x && nth < CUDA_SOFT_MAX_BLOCK_SIZE) nth *= 2; <---
    const dim3 block_dims(nth,     1, 1);
    const dim3 block_nums(nrows_x, 1, 1);
    soft_max_f32<<<block_nums, block_dims, 0, stream>>>(x, y, dst, ncols_x, nrows_y, scale);
}

[JohannesGaessler]

I pushed a fix for very large tensors. In that case instead of shared memory global memory is used to hold the values (not ideal for performance).

I don't think this can be right, it would imply a throughput in the thousands of GB/s. Am I wrong? (4096×16384×4×2 bytes) / 85us −> GiB/s = 5882 GiB/s

The 85 µs are for the implementation using FP16. When using FP32 NVIDIA Nsight Systems reports 150 µs. But I agree that the implied memory throughput in either case seems too high.

[JohannesGaessler]

I pushed a fix for ncols = 33, 65, 129, ... where the issue seemed to be that the shared memory buffer was not being padded correctly.

[JohannesGaessler]

@FSSRepo how long is the compute sanitizer supposed to run for your example using LLaMA 2 13b? I killed it after 15 minutes because I need my GPU for something else. Also you seem to have forgotten -ngl 99 for you example. I cannot reproduce any race conditions using tests/test-backend-ops. Can you instead check whether the race conditions still occur with this PR?

[slaren]

Recap of the performance that I see with test-backend-ops, 3090 Ti, WSL:

test	master	PR	PR F16	speedup	speedup F16
SOFT_MAX(type=f32,ne=[10,10,10,10])	20.25 GB/s	20.28 GB/s	22.07 GB/s	1.001	1.089
SOFT_MAX(type=f32,ne=[4096,1,1,1])	5.37 GB/s	7.34 GB/s	7.51 GB/s	1.366	1.398
SOFT_MAX(type=f32,ne=[1024,512,1,1])	221.22 GB/s	313.56 GB/s	544.22 GB/s	1.417	2.460
SOFT_MAX(type=f32,ne=[4096,512,1,1])	637.94 GB/s	683.83 GB/s	685.38 GB/s	1.071	1.074
SOFT_MAX(type=f32,ne=[4096,16384,1,1])	791.93 GB/s	837.56 GB/s	837.13 GB/s	1.057	1.057

There may still be some issues with F16, I tried some random dimensions and some tests have a very large error:

  SOFT_MAX(type=f32,ne=[27068,3,1,1]): [SOFT_MAX] NMSE = 1.000E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[60364,480,1,1]): [SOFT_MAX] NMSE = 2.123E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[40397,67,1,1]): [SOFT_MAX] NMSE = 1.753E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[60604,151,1,1]): [SOFT_MAX] NMSE = 2.240E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[95781,111,1,1]): [SOFT_MAX] NMSE = 6.218E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[91341,24,1,1]): [SOFT_MAX] NMSE = 8.201E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[92999,316,1,1]): [SOFT_MAX] NMSE = 5.111E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[59815,366,1,1]): [SOFT_MAX] NMSE = 2.363E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[85203,225,1,1]): [SOFT_MAX] NMSE = 4.698E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[34963,461,1,1]): [SOFT_MAX] NMSE = 1.000E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[90831,115,1,1]): [SOFT_MAX] NMSE = 5.474E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[69272,34,1,1]): [SOFT_MAX] NMSE = 5.267E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[83132,229,1,1]): [SOFT_MAX] NMSE = 4.389E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[59246,17,1,1]): [SOFT_MAX] NMSE = 3.770E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[81631,48,1,1]): [SOFT_MAX] NMSE = 4.178E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[59536,390,1,1]): [SOFT_MAX] NMSE = 1.860E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[86513,329,1,1]): [SOFT_MAX] NMSE = 4.791E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[64085,127,1,1]): [SOFT_MAX] NMSE = 2.969E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[99567,221,1,1]): [SOFT_MAX] NMSE = 5.428E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[68392,54,1,1]): [SOFT_MAX] NMSE = 2.938E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[73394,387,1,1]): [SOFT_MAX] NMSE = 3.807E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[45188,221,1,1]): [SOFT_MAX] NMSE = 1.728E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[44265,162,1,1]): [SOFT_MAX] NMSE = 1.675E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[86377,410,1,1]): [SOFT_MAX] NMSE = 4.713E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[92523,330,1,1]): [SOFT_MAX] NMSE = 4.958E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[67839,104,1,1]): [SOFT_MAX] NMSE = 3.744E+00 > 1.000E-07 FAIL
  SOFT_MAX(type=f32,ne=[64483,77,1,1]): [SOFT_MAX] NMSE = 2.855E+00 > 1.000E-07 FAIL

To generate the tests I used this code:

    std::random_device rd;
    std::default_random_engine rng(rd());
    std::uniform_int_distribution<> dist1(1, 100000);
    std::uniform_int_distribution<> dist2(1, 512);
    for (int n = 0; n < 100; n++) {
        int64_t ne0 = dist1(rng);
        int64_t ne1 = dist2(rng);
        test_cases.emplace_back(new test_soft_max(GGML_TYPE_F32, {ne0, ne1, 1, 1}));
    }

I am not convinced about the scientific notation in the NMSE error, it makes harder to see at a glance what errors are huge. Maybe we could increase the number of decimals printed?

Edit: the tests seem to fail with ne0 > 24512.

It would be good to add the cases that are known to be edge cases to test-backend-ops to avoid regressions in the future, with low batch sizes to avoid adding too much time to the test.

[bobqianic]

It would be good to add the cases that are known to be edge cases to test-backend-ops to avoid regressions in the future, with low batch sizes to avoid adding too much time to the test.

I'm thinking it might be a good idea to implement a comprehensive stress test on GGML, ensuring we cover all operators with 100% accuracy. We could start by establishing a standard range for inputs, and then conduct our tests within this specified range for every possible dimension. This approach would allow us to confidently assure that, as long as inputs fall within this range, the outcomes will be accurate. If users encounter issues with inputs outside our defined range, they can report it to us. We can then consider expanding our testing range to address these problems.

For our routine testing, we can casually choose random dimensions as our focus, but it's important to periodically conduct a comprehensive stress test to identify any potential issues.

[JohannesGaessler]

This approach would allow us to confidently assure that, as long as inputs fall within this range, the outcomes will be accurate.

That is literally impossible. For floating point addition $c = a + b$ will always give you garbage results for $a \approx -b$ and $|a - b| << |a|$ due to cancellation.

[ggerganov]

Maybe we could increase the number of decimals printed?

I would prefer that as well

I'm thinking it might be a good idea to implement a comprehensive stress test on GGML

We can discuss this further in a separate discussion to avoid derailing the conversation in this PR

[bobqianic]

That is literally impossible. For floating point addition c=a+b will always give you garbage results for a≈−b and |a−b|<<|a| due to cancellation.

No, what I'm referring to is accuracy in terms of alignment with the ground truth values. Or, to put it another way, achieving numerical consistency across all backends when the same data type is used.

[FSSRepo]

@JohannesGaessler I still can't get the idea out of my head that flash attention could help optimize the performance of the attention layer. We briefly discussed this before, but it was quickly dismissed #778 because there was no perceived performance gain, leaving only the ggml_flash_attn operation. However, this was only because it was implemented on the CPU, which is incorrect, as this technique is intended to optimize performance and reduce GPU memory usage by leveraging shared memory.

I've been reviewing and trying to interpret the code from the original repository, but I'm lost. I could simply try to compile the code and implement it externally to ggml, but it requires cutlass, and that library is too heavy, and very complex. It's a shame that all these techniques are inherently implemented for use with PyTorch, making their reuse challenging.

[JohannesGaessler]

One thing at a time. I did this PR even though the performance impact is small because it was fairly low-risk and straightforward. My priorities after that will be to try and implement faster general matrix matrix multiplication using quantized data. And after that I will work on better support for quantizing the KV cache and training. At some point I will maybe look into this but the limiting factor for me is how much time I have and I want to first get the fundamentals straight before I do anything else.

[FSSRepo]

Don't worry, everything will happen in due time. I wish you luck; your work is very refined and excellent.

[slaren]

We could start by establishing a standard range for inputs, and then conduct our tests within this specified range for every possible dimension.

We cannot test for every possible input for obvious reasons, but I would like to implement fuzz testing at some point.

[JohannesGaessler]

I pushed a fix for FP16 soft max. The problem was that due to the conversion from FP32 to FP16 you no longer had a one-to-one conversion between dst and tmp values. So if you then use the dst buffer to store temporary values you get race conditions between warps. I fixed it by padding the tmp addresses for the dst buffer.

I also expanded the soft max test code to now test pseudo-random shapes with an increasing scale for ne0. This allows you to estimate how NMSE evolves: for large ne0 values it increases because there are a lot of very low probability outputs that have large relative errors. I seeded the RNG to make the test results deterministic.

I changed the NMSE display to show a flat 9 decimals after the point.

[LostRuins]

Hi, after this commit I am unable to build for CUDA anymore, with a few errors (all variants of the same error)

ggml-cuda.cu(5474): error : identifier "__hmax" is undefined
ggml-cuda.cu(5462): error : identifier "__hmax2" is undefined

Before this commit, it was fine.

Am attempting to use CUDA 11.4.4 on Windows.
Any ideas?

[JohannesGaessler]

What specific commands are you running and which GPU are you using? It could be an issue with compute capability because IIRC certain half instructions are not defined for compute capabilities < 6.0 (Pascal).

[henk717]

Does that mean this breaks compatibility with my maxwell GPU?

[JohannesGaessler]

It should not. There are kernels for both FP16 and FP32, the latter of which should work without issue on Maxwell.

[LostRuins]

@JohannesGaessler the device I am compiling on does not have a GPU, thought it's using CUDA Toolkit Version: 11.4.152.

I am building for multiple architectures, I have set(CMAKE_CUDA_ARCHITECTURES "37;52;61;70") - so far this has worked prior to this commit.

[cebtenzzre]

Builds fine for Maxwell, albeit with some warnings:

/home/jared/src/forks/llama.cpp/ggml-cuda.cu(598): warning #177-D: function "warp_reduce_sum(half2)" was declared but never referenced
                                   half2 warp_reduce_sum(half2 a) {
                                         ^

Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

/home/jared/src/forks/llama.cpp/ggml-cuda.cu(619): warning #177-D: function "warp_reduce_max(half2)" was declared but never referenced
                                   half2 warp_reduce_max(half2 x) {
                                         ^

I didn't know that anyone was using llama.cpp on Kepler. Should we be testing that in CI?

[JohannesGaessler]

The lines associated with the build errors are wrapped in #if !(defined(GGML_USE_HIPBLAS) && defined(__HIP_PLATFORM_AMD__)) && __CUDA_ARCH__ >= CC_PASCAL. And For CUDA 11.4.4 __hmax and __hmax2 are both defined. So I don't understand why the build is failing.

[JohannesGaessler]

I remembered that at my physics institute there is a machine with CUDA 11.0 and 4x some old Maxwell GPU for testing purposes. On that machine I can build llama.cpp both for compute_52 and compute_37 without error.

[JohannesGaessler]

Okay, I can reproduce the issue. The problem does not seem to stem from compute capabilities 3.7 or 5.2 but rather it is the builds for 6.1 and 7.0 which are failing. So for some reason __hmax and __hmax2 really are not defined.

[LostRuins]

Thanks! Was worried it was something I did wrong

[LostRuins]

Fixed by #4862