Yes and no. Read his notes about his goals and what he didn't implement, and note that his job when he wrote that was as part of a team that ports models to different hardware. That's a smaller job than writing a shippable framework and Anthropic has a team of people doing it. You can read the cuDNN docs to get an idea of some of the other stuff you'd need, there's a lot there but generally that's the point. So in one sense, yes, if you have a strong performance engineering team and a specific workload already you can make use of a variety of hardware. In another sense, no, a small team isn't likely to be able to write a direct alternative for the CUDA ecosystem that reaches devex parity. Lots of people have a pretty clear idea of what it takes to do this but none of the major vendors seem to be doing the work.

Knowing that reaching broad devex parity is very expensive I think the real win is figuring out what specific problem you have and building community and robust software support around that.

cuBLAS is not really Nvidia's moat: every competitor has a workable BLAS implementation and some are even drop-in replacements for cuBLAS.

In fact cuBLAS and CUDA are kinda orthogonal in that you're either calling a pre-built cuBLAS kernel or writing your own CUDA kernel but not really combining the two.

I'd say CUDA shines more because of stability, documentation, community support + examples, and ability to use modern C++ features in GPU code.

Of course it can be done, it's just a lot of effort. You need parametric kernels you can find optimal configurations for all hardware and input size combinations.

Then there are also numerics: being fast is not enough if your implementation accumulates a lot of rounding errors doing so. Floating point arithmetic can and will mess up your results in unexpected ways. -funsafe famously is neither fun nor safe.

Maybe tooling will catch up and make it easier. Think tinygrad with beamsearch, triton or halide.

> In other discussion here, people asserted that a CUDA replacement was unworkable because you couldn't replace Nvidia's CuBLAS implementation

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).

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.

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.

I'm a little sceptical of your claims. Care to share the kernel you wrote?

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.

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.

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.

I shared the code: https://github.com/arekpaterek/Faster_SGEMM_CUDA

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.

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?

I gotta see it to believe it ;)

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.

For all doubters, I open-sourced it: https://github.com/arekpaterek/Faster_SGEMM_CUDA

Post the code and your curriculum

The problem with AMD isn't that they are only hitting 90% of the hardware capabilities vs Nvidia hitting 98%.

It's the fact that AMD doesn't prioritize the reliability of its hardware and software stack. If I run llama.cpp on Vulkan I get a reasonable speedup, but if I raise the batch size to 512, the GPU is starting to make strange noises and shuts the PC down midway. Very cool. 98% of zero is still zero.

As reported by the article: "However, for smaller matrices, we’re doing poorly in comparison to Nvidia’s library! This happens because cuBLAS contains not one single implementation of SGEMM, but hundreds of them.", it would take considerable effort to replace CuBLAS.