perf(rocm): profile gfx1151 decode per kernel and rank the #1814 ports

benchmarks/rocm_profiles/gfx1151_929c80ab/Meta-Llama-3.1-8B-Instruct-4bit_greedy_bench.log

@@ -0,0 +1,16 @@
+[mlx-rocm] bound HIP device 0: gfx1151 (AMD Radeon 8060S Graphics) cus=20 warp=32 lds=64KB
+[Load] wall: 0.165 s  MLX peak: 0.00 GB
+[long-prompt] target=512 tokens -> using 512 tokens (max_context=Some(131072), reserved 128 for generation)
+[phase] warmup_start monotonic_ns=169763750669641 boottime_ns=169763750669691
+[mlx-rocm] gfx1151 native_wmma=1 cus=20 warp=32 low_cu=0 host_WARP_SIZE=32
+[hipBLASLt caps] device 0 (gfx1151): bf16=1 fp8_e4m3=0 fp8_e5m2=0 int8=1
+[long-prompt] target=512 tokens -> using 512 tokens (max_context=Some(131072), reserved 128 for generation)
+[phase] measured_start monotonic_ns=169765813493615 boottime_ns=169765813493645
+[phase] decode_start monotonic_ns=169766321703540 boottime_ns=169766321703560
+[phase] measured_end monotonic_ns=169769539602196 boottime_ns=169769539602216
+[Profile Results]
+  Prompt tokens:    512
+  Generated tokens: 128
+  Prefill:          478.76 ms (1069.42 tok/s)
+  Decode:           3217.90 ms (39.78 tok/s)
+  MLX peak memory:  20.60 GB

benchmarks/rocm_profiles/gfx1151_929c80ab/Meta-Llama-3.1-8B-Instruct-4bit_greedy_decode_kernels.csv

@@ -0,0 +1,20 @@
+kernel,class,roles,port_units,calls,calls_per_token,total_ns,avg_ns,share_pct
+"void mlx::core::rocm::qmv_wide_kernel<__half, __half, 64, 4>(__half const*, unsigned char const*, __half const*, __half const*, __half*, int, int, int, bool)",qmv,,,20447,159.74,2832888679,138548,94.812
+"void mlx::core::rocm::kernel_sdpav_1pass<__half, false, 128>(__half const*, __half const*, __half const*, __half*, __half const*, mlx::core::rocm::AttnParams)",sdpa,,,4064,31.75,63928915,15731,2.140
+"void mlx::core::rocm::rms_norm_kernel<__half, 256, 4>(__half const*, __half const*, __half*, float, unsigned int, long)",rms_norm,add_rms_join;add_rms_join_post_attn,2063,8255,64.49,34238309,4148,1.146
+"void mlx::core::rocm::binary_vv<mlx::core::rocm::Add, __half, __half, unsigned int, 4>(__half const*, __half const*, __half*, unsigned int)",binary,add_rms_join;add_rms_join_post_attn,2063,8128,63.5,16540628,2035,0.554
+"void mlx::core::rocm::copy_gg_byval<__half, __half, long>(__half const*, __half*, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int)",copy,rope_append,2063,8256,64.5,15280116,1851,0.511
+"void mlx::core::rocm::rope_single_freqs_1d<__half, false, true>(__half const*, __half*, int const*, float const*, float, long, unsigned int, unsigned int, long)",rope,rope_append,2063,8128,63.5,12891275,1586,0.431
+"void mlx::core::rocm::Cf2ISigmoidADf2IBroadcastACEf2IBroadcastCAFf2IMultiplyDEGf2IBroadcastFBHf2IBroadcastBFIf2OMultiplyGH_VV_f2f2_6142509188972423790_contiguous<unsigned int, 16>(__half const*, __half const*, __half*, unsigned int)",compiled,,,4064,31.75,8863779,2181,0.297
+"void mlx::core::rocm::arg_reduce_partial<__half, mlx::core::rocm::ArgMax<__half>, 256>(__half const*, __half*, unsigned int*, unsigned long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, int, long, int, int)",arg_reduce,sampler_tail,2064,127,0.99,713374,5617,0.024
+"void mlx::core::rocm::copy_v<__half, __half, unsigned int, 4>(__half const*, __half*, unsigned int)",copy,sampler_tail,2064,127,0.99,388288,3057,0.013
+"void mlx::core::rocm::gather_rows_kernel<__half, int>(__half const*, int const*, __half*, long, unsigned int, int)",gather_scatter,,,254,1.98,370163,1457,0.012
+"void mlx::core::rocm::gather_axis_kernel<__half, int, 1, true, true>(__half const*, int const*, __half*, long, long, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int, int, long, long)",gather_scatter,sampler_tail,2064,127,0.99,276394,2176,0.009
+"void mlx::core::rocm::arg_reduce_final<__half, mlx::core::rocm::ArgMax<__half>, 256>(__half const*, unsigned int const*, unsigned int*, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, int, unsigned long, int)",arg_reduce,sampler_tail,2064,127,0.99,245578,1934,0.008
+"void mlx::core::rocm::gather_rows_kernel<unsigned int, int>(unsigned int const*, int const*, unsigned int*, long, unsigned int, int)",gather_scatter,,,127,0.99,243372,1916,0.008
+"void mlx::core::rocm::scatter_axis_kernel<__half, int, false, 1, true, true>(__half const*, int const*, __half*, long, long, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int, int, long, long, long)",gather_scatter,sampler_tail,2064,127,0.99,237324,1869,0.008
+"void mlx::core::rocm::binary_vv<mlx::core::rocm::Add, float, float, unsigned int, 4>(float const*, float const*, float*, unsigned int)",binary,sampler_tail,2064,127,0.99,199572,1571,0.007
+"void mlx::core::rocm::copy_v<__half, float, unsigned int, 4>(__half const*, float*, unsigned int)",copy,sampler_tail,2064,127,0.99,187838,1479,0.006
+"void mlx::core::rocm::affine_dequantize_packed_kernel<__half, 4>(unsigned char const*, __half const*, __half const*, __half*, unsigned long, int)",dequantize,,,127,0.99,144035,1134,0.005
+"void mlx::core::rocm::copy_s<__half, __half, unsigned int, 4>(__half const*, __half*, unsigned int)",copy,,,64,0.5,132975,2078,0.004
+"void mlx::core::rocm::copy_v<float, __half, unsigned int, 4>(float const*, __half*, unsigned int)",copy,sampler_tail,2064,127,0.99,124713,982,0.004

benchmarks/rocm_profiles/gfx1151_929c80ab/Meta-Llama-3.1-8B-Instruct-4bit_greedy_kernel_stats.csv

@@ -0,0 +1,27 @@
+"Name","Calls","TotalDurationNs","AverageNs","Percentage","MinNs","MaxNs","StdDev"
+"void mlx::core::rocm::qmv_wide_kernel<__half, __half, 64, 4>(__half const*, unsigned char const*, __half const*, __half const*, __half*, int, int, int, bool)",23508,3251875430,138330.586609,73.19,57067,2458585,103690.064508
+"Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT80x128x64_MI16x16x1_SN_LDSB0_AFC1_AG0_AGGSUA0_AGNTAB0_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR0_CLS0_CADS0_DTLA0_DTLB0_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_HPLR0_ISA1151_ICIW0_IU1_K1_LDSTI0_LBSPPA128_LBSPPB128_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA8_LPB8_LPMXSA0_LPMXSB0_LPM0_LRVW16_LWPMn1_MIAV1_MIWT5_2_MXLIBL_MXSFNS_MO40_MGRIPM1_NTn1_NTA0_NTB0_NTC0_NTD0_NTE0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS0_NLCA1_NLCB1_ONLL1_PAP0_PGL0_PGR2_PLR0_PKA0_SGROB0_SIA1_SS1_SPO0_SRVW0_SSO0_SVW1_SK0_SKFTR0_SKFDPO0_SKWS0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TLDS1_TLDSM1_ULSGRO0_USL1_USLMX0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS32_WG16_8_1_WGMXCC1",192,393932333,2051730.901042,8.87,1199246,2792190,490966.644453
+"Cijk_Alik_Bljk_HHS_BH_Bias_HA_S_SAV_UserArgs_MT64x128x64_MI16x16x1_SN_LDSB0_AFC1_AG0_AGGSUA0_AGNTAB0_AFEM1_AFEM1_ASEM1_BL1_BS1_CD1_1_CLR0_CLS0_CADS0_DTLA0_DTLB0_DTLM0_DTVA0_DTVB0_DTVMXSA0_DTVMXSB0_DTVSM0_DPLB0_EPS0_ELFLR0_EMLLn1_FDSI0_GRPM1_GRVWA8_GRVWB8_GSUAMB_GLS0_HPLR0_ISA1151_ICIW0_IU1_K1_LDSTI0_LBSPPA128_LBSPPB128_LBSPPMXSA0_LBSPPMXSB0_LBSPPM0_LPA8_LPB8_LPMXSA0_LPMXSB0_LPM0_LRVW16_LWPMn1_MIAV1_MIWT4_2_MXLIBL_MXSFNS_MO40_MGRIPM1_NTn1_NTA0_NTB0_NTC0_NTD0_NTE0_NTMXSA0_NTMXSB0_NTM0_NTWS0_NVn1_NVA0_NVB0_NVC0_NVD0_NVE0_NVMXSA0_NVMXSB0_NVM0_NVWS0_NEPBS0_NLCA1_NLCB1_ONLL1_PAP0_PGL0_PGR2_PLR0_PKA0_SGROB0_SIA1_SS1_SPO0_SRVW0_SSO0_SVW1_SK0_SKFTR0_SKFDPO0_SKWS0_SKXCCM0_SNLL0_SIP1_SGRO0_TDMI0_TDMIM0_TDMS0_TIN0_THn1_THA0_THB0_THC0_THD0_THE0_THMXSA0_THMXSB0_THM0_THWS0_TLDS1_TLDSM1_ULSGRO0_USL1_USLMX0_UDFMAC0_UIOFGRO0_UPLRP0_USFGROn1_USI0_VSn1_VWA1_VWB1_WSGRA0_WSGRB0_WS32_WG16_8_1_WGMXCC1",128,236053820,1844170.468750,5.31,895597,3243274,837111.458130
+"void mlx::core::rocm::affine_dequantize_packed_kernel<__half, 4>(unsigned char const*, __half const*, __half const*, __half*, unsigned long, int)",468,235582978,503382.431624,5.30,961,1783020,483086.513258
+"void mlx::core::rocm::kernel_sdpa_flash_wmma<__half, true, 128, 64, 64>(__half const*, __half const*, __half const*, __half*, float*, mlx::core::rocm::FAWmmaParams)",64,95527733,1492620.828125,2.15,1450638,2002912,67198.957956
+"void mlx::core::rocm::kernel_sdpav_1pass<__half, false, 128>(__half const*, __half const*, __half const*, __half*, __half const*, mlx::core::rocm::AttnParams)",4672,72747549,15570.965111,1.64,8255,32421,1698.779358
+"void mlx::core::rocm::rms_norm_kernel<__half, 256, 4>(__half const*, __half const*, __half*, float, unsigned int, long)",9620,46824909,4867.454158,1.05,3647,139662,8129.922147
+"void mlx::core::rocm::binary_vv<mlx::core::rocm::Add, __half, __half, unsigned int, 4>(__half const*, __half const*, __half*, unsigned int)",9472,26370407,2784.037901,0.5935,1562,306414,7041.504417
+"void mlx::core::rocm::copy_gg_byval<__half, __half, long>(__half const*, __half*, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int)",9792,23867678,2437.467116,0.5372,1001,176971,7721.667921
+"void mlx::core::rocm::Cf2ISigmoidADf2IBroadcastACEf2IBroadcastCAFf2IMultiplyDEGf2IBroadcastFBHf2IBroadcastBFIf2OMultiplyGH_VV_f2f2_6142509188972423790_contiguous<unsigned int, 16>(__half const*, __half const*, __half*, unsigned int)",4736,19325787,4080.613809,0.4350,1683,241412,16195.769640
+"void mlx::core::rocm::rope_single_freqs_1d<__half, false, true>(__half const*, __half*, int const*, float const*, float, long, unsigned int, unsigned int, long)",9344,15007453,1606.105843,0.3378,1082,17553,1017.044898
+"void mlx::core::rocm::copy_gg_byval<unsigned int, unsigned int, long>(unsigned int const*, unsigned int*, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int)",96,11075268,115367.375000,0.2493,27051,573434,107719.186459
+"void mlx::core::rocm::rope_freqs<__half, false, true>(__half const*, __half*, int const*, float const*, float, float, mlx::core::rocm::hip_array<long, 3>, mlx::core::rocm::hip_array<long, 3>, long, int, HIP_vector_type<unsigned int, 3u>, long)",128,6346774,49584.171875,0.1428,9739,99587,38078.550437
+"void mlx::core::rocm::copy_g_byval<__half, __half, long>(__half const*, __half*, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, int)",64,4396070,68688.593750,0.0989,63399,310862,30857.517482
+"void mlx::core::rocm::arg_reduce_partial<__half, mlx::core::rocm::ArgMax<__half>, 256>(__half const*, __half*, unsigned int*, unsigned long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, int, long, int, int)",148,848108,5730.459459,0.0191,5130,10700,795.267314
+"void mlx::core::rocm::gather_rows_kernel<__half, int>(__half const*, int const*, __half*, long, unsigned int, int)",296,480889,1624.625000,0.0108,962,27973,1923.671979
+"void mlx::core::rocm::copy_v<__half, __half, unsigned int, 4>(__half const*, __half*, unsigned int)",148,456336,3083.351351,0.0103,2364,19516,1697.485839
+"void mlx::core::rocm::gather_axis_kernel<__half, int, 1, true, true>(__half const*, int const*, __half*, long, long, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int, int, long, long)",148,336504,2273.675676,7.574e-03,1522,6853,1183.689017
+"void mlx::core::rocm::gather_rows_kernel<unsigned int, int>(unsigned int const*, int const*, unsigned int*, long, unsigned int, int)",148,328052,2216.567568,7.384e-03,1282,25288,2340.444749
+"void mlx::core::rocm::scatter_axis_kernel<__half, int, false, 1, true, true>(__half const*, int const*, __half*, long, long, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int, int, long, long, long)",148,284974,1925.500000,6.414e-03,1322,10019,1369.125897
+"void mlx::core::rocm::arg_reduce_final<__half, mlx::core::rocm::ArgMax<__half>, 256>(__half const*, unsigned int const*, unsigned int*, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, int, unsigned long, int)",148,283365,1914.628378,6.378e-03,1643,10339,898.730286
+"void mlx::core::rocm::copy_s<__half, __half, unsigned int, 4>(__half const*, __half*, unsigned int)",128,251754,1966.828125,5.666e-03,1563,5330,556.919227
+"void mlx::core::rocm::binary_vv<mlx::core::rocm::Add, float, float, unsigned int, 4>(float const*, float const*, float*, unsigned int)",148,248064,1676.108108,5.583e-03,881,6452,1437.128032
+"void mlx::core::rocm::copy_v<__half, float, unsigned int, 4>(__half const*, float*, unsigned int)",148,222987,1506.668919,5.019e-03,1042,8296,991.654252
+"__amd_rocclr_fillBufferUnAligned",1,167714,167714.000000,3.775e-03,167714,167714,0.00000000e+00
+"void mlx::core::rocm::copy_v<float, __half, unsigned int, 4>(float const*, __half*, unsigned int)",148,158217,1069.033784,3.561e-03,801,6372,576.180957

benchmarks/rocm_profiles/gfx1151_929c80ab/Meta-Llama-3.1-8B-Instruct-4bit_greedy_plain_bench.log

@@ -0,0 +1,12 @@
+[mlx-rocm] bound HIP device 0: gfx1151 (AMD Radeon 8060S Graphics) cus=20 warp=32 lds=64KB
+[Load] wall: 0.200 s  MLX peak: 0.00 GB
+[long-prompt] target=512 tokens -> using 512 tokens (max_context=Some(131072), reserved 128 for generation)
+[mlx-rocm] gfx1151 native_wmma=1 cus=20 warp=32 low_cu=0 host_WARP_SIZE=32
+[hipBLASLt caps] device 0 (gfx1151): bf16=1 fp8_e4m3=0 fp8_e5m2=0 int8=1
+[long-prompt] target=512 tokens -> using 512 tokens (max_context=Some(131072), reserved 128 for generation)
+[Profile Results]
+  Prompt tokens:    512
+  Generated tokens: 128
+  Prefill:          475.90 ms (1075.86 tok/s)
+  Decode:           3452.34 ms (37.08 tok/s)
+  MLX peak memory:  20.60 GB

benchmarks/rocm_profiles/gfx1151_929c80ab/Meta-Llama-3.1-8B-Instruct-4bit_greedy_summary.json

@@ -0,0 +1,175 @@
+{
+  "name": "Meta-Llama-3.1-8B-Instruct-4bit_greedy",
+  "model_type": "llama",
+  "temperature": 0.0,
+  "clock": "boottime",
+  "prompt_tokens": 512,
+  "generated_tokens": 128,
+  "profiled_decode_tok_s": 39.78,
+  "profiled_prefill_tok_s": 1069.42,
+  "plain_decode_tok_s": 37.08,
+  "plain_prefill_tok_s": 1075.86,
+  "profiler_decode_slowdown_pct": -6.79,
+  "decode_dispatches": 62930,
+  "dispatches_per_token": 491.6,
+  "decode_wall_ms": 3217.899,
+  "decode_gpu_sum_ms": 2987.895,
+  "decode_gpu_busy_ms": 2987.895,
+  "gpu_ms_per_token": 23.3429,
+  "host_gap_ms_per_token": 1.7969,
+  "host_gap_pct_of_wall": 7.15,
+  "plain_wall_ms_per_token": 26.9687,
+  "plain_host_gap_ms_per_token_est": 3.6258,
+  "checks": {
+    "kernels_straddling_decode_start": 0,
+    "idle_gap_before_first_decode_dispatch_us": 30510.9,
+    "last_dispatch_before_decode": "void mlx::core::rocm::arg_reduce_final<__half, mlx::core::rocm::ArgMax<__half>, 256>(__half const*, unsigned int const*,"
+  },
+  "top_kernels": [
+    {
+      "kernel": "void mlx::core::rocm::qmv_wide_kernel<__half, __half, 64, 4>(__half const*, unsigned char const*, __half const*, __half const*, __half*, int, int, int, bool)",
+      "class": "qmv",
+      "calls_per_token": 159.74,
+      "share_pct": 94.81
+    },
+    {
+      "kernel": "void mlx::core::rocm::kernel_sdpav_1pass<__half, false, 128>(__half const*, __half const*, __half const*, __half*, __half const*, mlx::core::rocm::AttnParams)",
+      "class": "sdpa",
+      "calls_per_token": 31.75,
+      "share_pct": 2.14
+    },
+    {
+      "kernel": "void mlx::core::rocm::rms_norm_kernel<__half, 256, 4>(__half const*, __half const*, __half*, float, unsigned int, long)",
+      "class": "rms_norm",
+      "calls_per_token": 64.49,
+      "share_pct": 1.15
+    },
+    {
+      "kernel": "void mlx::core::rocm::binary_vv<mlx::core::rocm::Add, __half, __half, unsigned int, 4>(__half const*, __half const*, __half*, unsigned int)",
+      "class": "binary",
+      "calls_per_token": 63.5,
+      "share_pct": 0.55
+    },
+    {
+      "kernel": "void mlx::core::rocm::copy_gg_byval<__half, __half, long>(__half const*, __half*, long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, mlx::core::rocm::hip_array<long, 10>, int)",
+      "class": "copy",
+      "calls_per_token": 64.5,
+      "share_pct": 0.51
+    },
+    {
+      "kernel": "void mlx::core::rocm::rope_single_freqs_1d<__half, false, true>(__half const*, __half*, int const*, float const*, float, long, unsigned int, unsigned int, long)",
+      "class": "rope",
+      "calls_per_token": 63.5,
+      "share_pct": 0.43
+    },
+    {
+      "kernel": "void mlx::core::rocm::Cf2ISigmoidADf2IBroadcastACEf2IBroadcastCAFf2IMultiplyDEGf2IBroadcastFBHf2IBroadcastBFIf2OMultiplyGH_VV_f2f2_6142509188972423790_contiguous<unsigned int, 16>(__half const*, __half const*, __half*, unsigned int)",
+      "class": "compiled",
+      "calls_per_token": 31.75,
+      "share_pct": 0.3
+    },
+    {
+      "kernel": "void mlx::core::rocm::arg_reduce_partial<__half, mlx::core::rocm::ArgMax<__half>, 256>(__half const*, __half*, unsigned int*, unsigned long, mlx::core::rocm::hip_array<int, 10>, mlx::core::rocm::hip_array<long, 10>, int, long, int, int)",
+      "class": "arg_reduce",
+      "calls_per_token": 0.99,
+      "share_pct": 0.02
+    },
+    {
+      "kernel": "void mlx::core::rocm::copy_v<__half, __half, unsigned int, 4>(__half const*, __half*, unsigned int)",
+      "class": "copy",
+      "calls_per_token": 0.99,
+      "share_pct": 0.01
+    },
+    {
+      "kernel": "void mlx::core::rocm::gather_rows_kernel<__half, int>(__half const*, int const*, __half*, long, unsigned int, int)",
+      "class": "gather_scatter",
+      "calls_per_token": 1.98,
+      "share_pct": 0.01
+    }
+  ],
+  "class_share_pct": {
+    "qmv": 94.81,
+    "sdpa": 2.14,
+    "rms_norm": 1.15,
+    "binary": 0.56,
+    "copy": 0.54,
+    "rope": 0.43,
+    "compiled": 0.3,
+    "gather_scatter": 0.04,
+    "arg_reduce": 0.03,
+    "dequantize": 0.0
+  },
+  "role_share_pct": {
+    "unattributed": 97.41,
+    "add_rms_join": 0.85,
+    "add_rms_join_post_attn": 0.83,
+    "rope_append": 0.83,
+    "sampler_tail": 0.08,
+    "moe_expert_gemv": 0.0,
+    "moe_activation": 0.0,
+    "moe_weighted_sum": 0.0,
+    "moe_gather_indices": 0.0,
+    "ssm_step": 0.0,
+    "paged_attention": 0.0
+  },
+  "role_dispatches_per_token": {
+    "unattributed": 229.2,
+    "rope_append": 127.5,
+    "add_rms_join_post_attn": 63.5,
+    "add_rms_join": 63.5,
+    "sampler_tail": 7.9,
+    "moe_expert_gemv": 0.0,
+    "moe_activation": 0.0,
+    "moe_weighted_sum": 0.0,
+    "moe_gather_indices": 0.0,
+    "ssm_step": 0.0,
+    "paged_attention": 0.0
+  },
+  "port_units": {
+    "2063": {
+      "port": "fused_add_rms_norm + fused_rope_qk_append",
+      "fallback_share_pct": 1.66,
+      "fallback_dispatches_per_token": 191.0,
+      "reached_default_share_pct": 0.0,
+      "reached_default_dispatches_per_token": 0.0,
+      "reached_optin_share_pct": 0.83,
+      "note": "both paths ship off (FUSED_ADD_RMSNORM_DEFAULT / FUSED_ROPE_APPEND_DEFAULT false); only the post-attention join calls the fused norm; rope_scaling builds a frequency table, which the RoPE kernel cannot take"
+    },
+    "2064": {
+      "port": "samplers (gumbel_max_sample, rejection_sample)",
+      "fallback_share_pct": 0.08,
+      "fallback_dispatches_per_token": 7.9,
+      "reached_default_share_pct": 0.0,
+      "reached_default_dispatches_per_token": 0.0,
+      "reached_optin_share_pct": 0.0,
+      "note": "greedy argmax dispatches neither sampler kernel"
+    },
+    "2065": {
+      "port": "fused MoE decode (moe_gateup, moe_down)",
+      "fallback_share_pct": 0.0,
+      "fallback_dispatches_per_token": 0.0,
+      "reached_default_share_pct": 0.0,
+      "reached_default_dispatches_per_token": 0.0,
+      "reached_optin_share_pct": 0.0,
+      "note": "llama does not call forward_fused_kernel"
+    },
+    "2067": {
+      "port": "SSM update (ssm_update_kernel)",
+      "fallback_share_pct": 0.0,
+      "fallback_dispatches_per_token": 0.0,
+      "reached_default_share_pct": 0.0,
+      "reached_default_dispatches_per_token": 0.0,
+      "reached_optin_share_pct": 0.0,
+      "note": "no Mamba2 layer"
+    },
+    "2068": {
+      "port": "paged attention (v1, v2, merge)",
+      "fallback_share_pct": 0.0,
+      "fallback_dispatches_per_token": 0.0,
+      "reached_default_share_pct": 0.0,
+      "reached_default_dispatches_per_token": 0.0,
+      "reached_optin_share_pct": 0.0,
+      "note": "the bench decodes into a dense KVCache; the paged path is not taken"
+    }
+  }
+}