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.