Google TPUs are built around a 128×128 systolic array of multiply-accumulate (MAC) units. Trainium 1, Trainium 2, and Inferentia 2 also feature a 128x128 systolic array.

You learn something every day. Today, it was the term "systolic array": A systolic array is a specialized grid of simple, interconnected processing units designed to execute parallel data operations—like matrix multiplication—by rhythmically passing data directly from cell to neighboring cell without writing intermediate results back to main memory.

The term comes from the biological word systole (the contraction of the heart pumping blood through the body). In a systolic array, data "pulses" through a network of processing elements on every clock cycle, driven by a global clock beat.

Just because scales are grouped by 128x128 tiles, does not mean you need a single compute tile that large. It works completely fine to process it with multiple smaller tiles that get given the same scales, like how this works on Hopper and Blackwell today

Huawei Ascend NPU

Blackwell. You could do it with wgmma on hooper but then you'll only be able to run 1 CTA per SM. There are cases where this is OK, but most commonly 128x128 mmas are primitive in tcgen05.mma i.e blackwell. The fundamental reasons is that the systolic array accelerator (TMA) until blackwell wrote to the cuda core registers themselves so you were limited by the register file size of the SM[SP]. In blackwell there's a separate TMEM where the mma unit stores it's output. You can even do higher 256x256 and such using a hardware feature called 2-CTA MMA, this is essentially them letting neighbouring pairs of SMs co-operate and access each other's memory.

As for sibling comments, huawei ascend is more of an NPU-style architecture where you can easily have much bigger MMAs as primitive. But you usually don't anyways for many reasons.

Huawei Ascend presumably?