Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by deredede on Friday July 26, 2024 @02:47PM
(#41081179)
When I was optimizing GPU kernels a few years back, Nvidia's own kernels were getting those last few percent of performance by making use of hardware-specific features (I remember operand caches being one) that are not available through PTX.
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by ap4 on Friday July 26, 2024 @02:49PM
(#41081197)
Beating CuBLAS is easy. I wrote a 30-line kernel that multiplies two 4096² matrices 14% faster than CuBLAS on a 4090.
The question is how to earn money on that.
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by coffeeaddict1 on Friday July 26, 2024 @03:02PM
(#41081286)
I'm a little sceptical of your claims. Care to share the kernel you wrote?
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by david-gpu on Friday July 26, 2024 @03:33PM
(#41081537)
I would bet actual money that they are not doing an apples-to-apples comparison.
I have seen how those high-performance libraries are made and I'm still in awe at the quality and quantity of the staffing involved. Those were the smartest and most knowledgeable engineers I met in my career.
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by ap4 on Saturday July 27, 2024 @12:15PM
(#41087477)
It is apples-to-apples. The code from the article runs a chosen kernel 50 times in a loop. In the case of CuBLAS it runs the function cublasGemmEx() 50 times.
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by david-gpu on Sunday July 28, 2024 @07:52AM
(#41092636)
Does your implementation support the complete feature set of cublasGemmEx()? E.g. varying matrix sizes, etc. Does your implementation beat cublas on a wide range of cases, or only cherry-picked examples? Does the code from the article follow the best coding practices for cublas in the first place?
Generalizing from a micro benchmark is typically hubris.
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by ap4 on Wednesday July 31, 2024 @09:08PM
(#41125254)
I shared the code: https://github.com/arekpaterek/Faster_SGEMM_CUDA
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by ladberg on Friday July 26, 2024 @04:33PM
(#41082029)
If this were true (and I highly doubt it) it's obvious how to make money from it: collect a 7 figure paycheck from Nvidia, AMD, or any FAANG.
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by ap4 on Friday July 26, 2024 @04:44PM
(#41082129)
I swapped one of the kernels in the code from the article to my kernel, and left only the multiplication of matrices of size 4096².
On average over 20 runs:
CuBLAS (./sgemm 0) has 50.9 TFLOPS.
My kernel has 61.8 TFLOPS, so it's actually +21% speedup in this benchmark.
How do I collect my paycheck?
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by aaa370 on Friday July 26, 2024 @05:17PM
(#41082354)
I gotta see it to believe it ;)
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by ap4 on Saturday July 27, 2024 @03:45PM
(#41088847)
Believe it or not.
On a 4090 gpu, average of 20 runs of SGEMM_CUDA:
size tflops_cublas tflops_my diff
4096² 50.8-50.9 61.8 +21%
8192² 56.3-56.4 67.1 +19%
16384² 53.6 66.7 +24%
I guess the right thing to do now would be to hire a B2B salesman and figure out, which company needs it.
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by ap4 on Wednesday July 31, 2024 @09:10PM
(#41125271)
For all doubters, I open-sourced it:
https://github.com/arekpaterek/Faster_SGEMM_CUDA
Re:How to Optimize a CUDA Matmul Kernel for CuBLAS-Like Perf...
by JuanJohnJames on Friday July 26, 2024 @05:01PM
(#41082238)
Post the code and your curriculum
Targeting nvidia GPUs? Or in general? For whom?
Building a performant BLAS library is hard but certainly not impossible. The tricks discussed in this post are hardly anything new either. Now, making a BLAS competitive with Nvidia's on its own GPUs is bound to be tough. But not technically unfeasible (after all, you can drop down to PTX if needed).