[Mosaic GPU] Add basic support for ldmatrix/stmatrix
For now we only support num=1, which doesn't yet let us reap the big
benefit of this instruction that comes from sharding the address
calculation. I'm leaving this for a follow-up to avoid overcomplicating
the change. It should already be no worse than going with regular shared
loads and stores.
PiperOrigin-RevId: 945625085
