tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%tmem], 128; tcgen05.ld.sync.aligned.32x32b.x64.b32 {%f0, %f1, %f2, %f3}, [%tmem];
cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes tcgen05.dealloc.cta_group::1.sync.aligned.b32 %tmem, 128;
mbarrier.try_wait.parity.shared::cta.b64 %p, [%bar], %phase; v_mfma_f32_32x32x8f16 v[0:15], v[16:17], v[18:19], v[0:15]
setp.ne.b32 %p, %r4, 0; s_waitcnt lgkmcnt(0)
tcgen05.mma.cta_group::1.kind::f16 [%d], %a, %b, %idesc, {%m0,%m1,%m2,%m3}, %p; ds_read_b128 v[20:23], v24 offset:0
tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%bar]; global_load_dwordx4 v[28:31], v[32:33], off
tcgen05.wait::st.sync.aligned; v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15]
tcgen05.ld.sync.aligned.32x32b.x64.b32 {%f0, %f1, %f2, %f3}, [%tmem]; s_setprio 1
tcgen05.dealloc.cta_group::1.sync.aligned.b32 %tmem, 128; buffer_load_dwordx4 v[36:39], v40, s[0:3], 0 offen
v_mfma_f32_32x32x8f16 v[0:15], v[16:17], v[18:19], v[0:15] v_pk_add_f32 v[42:43], v[44:45], v[46:47]
s_waitcnt lgkmcnt(0) s_barrier
ds_read_b128 v[20:23], v24 offset:0 tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%tmem], 128;
global_load_dwordx4 v[28:31], v[32:33], off cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] mbarrier.try_wait.parity.shared::cta.b64 %p, [%bar], %phase;
s_setprio 1 setp.ne.b32 %p, %r4, 0;
buffer_load_dwordx4 v[36:39], v40, s[0:3], 0 offen tcgen05.mma.cta_group::1.kind::f16 [%d], %a, %b, %idesc, {%m0,%m1,%m2,%m3}, %p;
v_pk_add_f32 v[42:43], v[44:45], v[46:47] tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%bar];
s_barrier tcgen05.wait::st.sync.aligned;
tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%tmem], 128; tcgen05.ld.sync.aligned.32x32b.x64.b32 {%f0, %f1, %f2, %f3}, [%tmem];
cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes tcgen05.dealloc.cta_group::1.sync.aligned.b32 %tmem, 128;
mbarrier.try_wait.parity.shared::cta.b64 %p, [%bar], %phase; v_mfma_f32_32x32x8f16 v[0:15], v[16:17], v[18:19], v[0:15]
setp.ne.b32 %p, %r4, 0; s_waitcnt lgkmcnt(0)
tcgen05.mma.cta_group::1.kind::f16 [%d], %a, %b, %idesc, {%m0,%m1,%m2,%m3}, %p; ds_read_b128 v[20:23], v24 offset:0
tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%bar]; global_load_dwordx4 v[28:31], v[32:33], off
tcgen05.wait::st.sync.aligned; v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15]
tcgen05.ld.sync.aligned.32x32b.x64.b32 {%f0, %f1, %f2, %f3}, [%tmem]; s_setprio 1
tcgen05.dealloc.cta_group::1.sync.aligned.b32 %tmem, 128; buffer_load_dwordx4 v[36:39], v40, s[0:3], 0 offen
v_mfma_f32_32x32x8f16 v[0:15], v[16:17], v[18:19], v[0:15] v_pk_add_f32 v[42:43], v[44:45], v[46:47]
s_waitcnt lgkmcnt(0) s_barrier
ds_read_b128 v[20:23], v24 offset:0 tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%tmem], 128;
global_load_dwordx4 v[28:31], v[32:33], off cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] mbarrier.try_wait.parity.shared::cta.b64 %p, [%bar], %phase;
s_setprio 1 setp.ne.b32 %p, %r4, 0;
buffer_load_dwordx4 v[36:39], v40, s[0:3], 0 offen tcgen05.mma.cta_group::1.kind::f16 [%d], %a, %b, %idesc, {%m0,%m1,%m2,%m3}, %p;
v_pk_add_f32 v[42:43], v[44:45], v[46:47] tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%bar];
s_barrier tcgen05.wait::st.sync.aligned;
tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%tmem], 128; tcgen05.ld.sync.aligned.32x32b.x64.b32 {%f0, %f1, %f2, %f3}, [%tmem];
cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes tcgen05.dealloc.cta_group::1.sync.aligned.b32 %tmem, 128;
mbarrier.try_wait.parity.shared::cta.b64 %p, [%bar], %phase; v_mfma_f32_32x32x8f16 v[0:15], v[16:17], v[18:19], v[0:15]
setp.ne.b32 %p, %r4, 0; s_waitcnt lgkmcnt(0)
tcgen05.mma.cta_group::1.kind::f16 [%d], %a, %b, %idesc, {%m0,%m1,%m2,%m3}, %p; ds_read_b128 v[20:23], v24 offset:0
tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%bar]; global_load_dwordx4 v[28:31], v[32:33], off
tcgen05.wait::st.sync.aligned; v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15]
tcgen05.ld.sync.aligned.32x32b.x64.b32 {%f0, %f1, %f2, %f3}, [%tmem]; s_setprio 1
tcgen05.dealloc.cta_group::1.sync.aligned.b32 %tmem, 128; buffer_load_dwordx4 v[36:39], v40, s[0:3], 0 offen
v_mfma_f32_32x32x8f16 v[0:15], v[16:17], v[18:19], v[0:15] v_pk_add_f32 v[42:43], v[44:45], v[46:47]
s_waitcnt lgkmcnt(0) s_barrier
ds_read_b128 v[20:23], v24 offset:0 tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%tmem], 128;
global_load_dwordx4 v[28:31], v[32:33], off cp.async.bulk.tensor.2d.shared::cluster.global.mbarrier::complete_tx::bytes
v_mfma_scale_f32_32x32x64_f8f6f4 v[0:15], v[16:23], v[24:31], v[0:15] mbarrier.try_wait.parity.shared::cta.b64 %p, [%bar], %phase;
s_setprio 1 setp.ne.b32 %p, %r4, 0;
buffer_load_dwordx4 v[36:39], v40, s[0:3], 0 offen tcgen05.mma.cta_group::1.kind::f16 [%d], %a, %b, %idesc, {%m0,%m1,%m2,%m3}, %p;
v_pk_add_f32 v[42:43], v[44:45], v[46:47] tcgen05.commit.cta_group::1.mbarrier::arrive::one.shared::cluster.b64 [%bar];
s_barrier tcgen05.wait::st.sync.aligned; Anand Pratap Singh
I make inference fast in Mojo at Modular. Attention, matmul, quantization and collectives, across accelerators, and whatever else stands between a kernel win and an end-to-end one. Before that, turbulence models learned from data, and a PhD in aerospace at Michigan.
Every kernel I touch lives somewhere under this. The job is working out which line is holding it down, and what it costs to move.