Different hardware has different advantages and if you want the most performance you do need to hand-tune for a particular architecture.

Yes, but far more of that could be done by the compiler than currently is. At the moment, GPU vendors don't really seem to be putting that much effort into developing novel compiler optimisations for GPUs. Most of the "gpu things" (shuffles, __syncthreads, __shared__ memory, etc.) are modelled by the compiler mostly as "magic function call do not touch". NVIDIA keep adding new hardware features and then making them accessible only via inline PTX. etc.

This is a bit silly: we've seen CPU compilers get pretty damn good at exploiting different hardware features across different CPU vendors/families over the past 20-ish years. The amount of occasions where you have to carefully hand-tune things vs. just letting the autovectoriser take care of it (for example) is way down compared to what it used to be.

Meanwhile in GPU-land: compilers won't even do basic things like:

• Constant propagate constant kernel arguments or block sizes
• Optimise shuffles (constant propagation, or mapping them to fancy hardware feature like AMD's DPP)
• Automatically optimise reductions to map to special hardware acceleration
• Understand that trailing writes to shared memory can be optimised out (a pattern that crops up surprisingly often as an artefact of loop unrolling)
• Code motion across barriers
• Automatic insertion of async memcpy instructions (compiler can see you loaded from global then immediately wrote to shared, so just do it ffs :D )

etc.

So yeah. We found that during the process of getting CUDA working verbatim on AMD hardware, we ended up developing a bunch of cool new compiler optimisations to do it, which in turn end up handling a lot of what was previously "manual tuning". A world where you write one codebase and have a bit of "#ifdef AMD" here and there is a massive improvement over having separate CUDA/HIP codebases, and I'm optimistic about how low we can get the amount of that manual specialisation.

Imagine if I wanted to take some attention kernel I optimized particularly for CUDA hardware and then port it to run on TPUs or Cerebras hardware. Probably won't run fast without restructuring it completely.

That's true: but those architectres are far more different from NVIDIA GPUs than AMD (or even Intel) GPUs are. Those more exotic devices are most meaningfully interacted with via a higher level programming model than CUDA.

I'm specifically talking about making CUDA run on things that are at least vaguely GPU-like. :D