News for hackers. Stuff that matters.
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)
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