iree-org · adedespirlet · Dec 5, 2025
diff --git a/tests/kernel/wave/moe/gemm_test.py b/tests/kernel/wave/moe/gemm_test.py
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_0_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_0_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_10_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_10_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_11_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_11_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_12_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_12_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_13_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_13_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_14_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_14_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_15_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_15_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_16_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_16_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_17_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_17_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_18_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_18_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_19_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_19_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_1_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_1_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_20_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_20_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_21_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_21_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_22_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_22_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_23_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_23_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_24_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_24_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_25_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_25_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_26_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_26_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_27_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_27_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_28_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_28_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_29_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_29_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_2_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_2_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_30_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_30_3.att
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_31_3.att b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_31_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_3_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_3_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_4_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_4_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_5_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_5_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_6_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_6_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_7_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_7_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_8_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_8_3.att
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_9_3.att b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_32617_shader_engine_9_3.att
diff --git a/...56x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_agent_info.csv b/...56x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_agent_info.csv
@@ -0,0 +1,11 @@
+"Node_Id","Logical_Node_Id","Agent_Type","Cpu_Cores_Count","Simd_Count","Cpu_Core_Id_Base","Simd_Id_Base","Max_Waves_Per_Simd","Lds_Size_In_Kb","Gds_Size_In_Kb","Num_Gws","Wave_Front_Size","Num_Xcc","Cu_Count","Array_Count","Num_Shader_Banks","Simd_Arrays_Per_Engine","Cu_Per_Simd_Array","Simd_Per_Cu","Max_Slots_Scratch_Cu","Gfx_Target_Version","Vendor_Id","Device_Id","Location_Id","Domain","Drm_Render_Minor","Num_Sdma_Engines","Num_Sdma_Xgmi_Engines","Num_Sdma_Queues_Per_Engine","Num_Cp_Queues","Max_Engine_Clk_Ccompute","Max_Engine_Clk_Fcompute","Sdma_Fw_Version","Fw_Version","Capability","Cu_Per_Engine","Max_Waves_Per_Cu","Family_Id","Workgroup_Max_Size","Grid_Max_Size","Local_Mem_Size","Hive_Id","Gpu_Id","Workgroup_Max_Dim_X","Workgroup_Max_Dim_Y","Workgroup_Max_Dim_Z","Grid_Max_Dim_X","Grid_Max_Dim_Y","Grid_Max_Dim_Z","Name","Vendor_Name","Product_Name","Model_Name"
+0,0,"CPU",128,0,0,0,0,0,0,0,0,1,128,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,3300,0,0,0,0,0,0,26,0,0,0,0,0,0,0,0,0,0,0,"AMD EPYC 9575F 64-Core Processor","CPU","AMD EPYC 9575F 64-Core Processor",""
+1,1,"CPU",128,0,128,0,0,0,0,0,0,1,128,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,3300,0,0,0,0,0,0,26,0,0,0,0,0,0,0,0,0,0,0,"AMD EPYC 9575F 64-Core Processor","CPU","AMD EPYC 9575F 64-Core Processor",""
+2,2,"GPU",0,1024,0,2147487744,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,29952,0,128,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,32762,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
+3,3,"GPU",0,1024,0,2147487780,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,1280,0,136,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,39903,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
+4,4,"GPU",0,1024,0,2147487816,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,25856,0,144,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,46071,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
+5,5,"GPU",0,1024,0,2147487852,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,5376,0,152,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,22482,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
+6,6,"GPU",0,1024,0,2147487888,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,62720,0,160,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,57243,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
+7,7,"GPU",0,1024,0,2147487924,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,34048,0,168,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,15294,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
+8,8,"GPU",0,1024,0,2147487960,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,58624,0,176,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,5014,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
+9,9,"GPU",0,1024,0,2147487996,8,160,0,64,64,8,256,32,32,1,9,4,32,90500,4098,30112,38144,0,184,2,14,8,24,3300,2200,12,30,2889327232,8,32,141,1024,4294967295,0,4409598424993686390,63411,1024,1024,1024,2147483647,65535,65535,"gfx950","AMD","AMD Radeon Graphics","ip discovery"
diff --git a/...x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_domain_stats.csv b/...x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_domain_stats.csv
@@ -0,0 +1,2 @@
+"Name","Calls","TotalDurationNs","AverageNs","Percentage","MinNs","MaxNs","StdDev"
+"KERNEL_DISPATCH",1040,4927545,4738.024038,100.00,640,23400,1058.615401
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_1.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_1.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_10.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_10.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_11.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_11.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_12.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_12.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_13.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_13.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_14.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_14.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_15.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_15.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_16.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_16.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_17.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_17.out
diff --git a/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_18.out b/...4_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_18.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_2.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_2.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_3.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_3.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_4.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_4.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_5.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_5.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_6.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_6.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_7.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_7.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_8.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_8.out
diff --git a/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_9.out b/...64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_gfx950_code_object_id_9.out
diff --git a/...x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_kernel_stats.csv b/...x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_kernel_stats.csv
diff --git a/...x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_kernel_trace.csv b/...x232_64x32x64_mfma_16x16x32_triplebuffer2/smci350-zts-gtu-b20-05/3785403_kernel_trace.csv
diff --git a/...6x256x232_64x32x64_mfma_16x16x32_triplebuffer2/stats_ui_output_agent_32617_dispatch_3.csv b/...6x256x232_64x32x64_mfma_16x16x32_triplebuffer2/stats_ui_output_agent_32617_dispatch_3.csv
@@ -0,0 +1,287 @@
+"CodeObj","Vaddr","Instruction","Hitcount","Latency","Stall","Idle","Source"
+3,5632,"; gemm_prefetch",0,0,0,0,""
+3,5632,"s_load_dwordx2 s[2:3], s[0:1], 0x0",0,0,0,0,""
+3,5640,"s_load_dwordx4 s[4:7], s[0:1], 0x8",0,0,0,0,""
+3,5648,"s_waitcnt lgkmcnt(0)",0,0,0,0,""
+3,5652,"s_branch 58",0,0,0,0,""
+3,5656,"s_nop 0",0,0,0,0,""
+3,5660,"s_nop 0",0,0,0,0,""
+3,5664,"s_nop 0",0,0,0,0,""
+3,5668,"s_nop 0",0,0,0,0,""
+3,5672,"s_nop 0",0,0,0,0,""
+3,5676,"s_nop 0",0,0,0,0,""
+3,5680,"s_nop 0",0,0,0,0,""
+3,5684,"s_nop 0",0,0,0,0,""
+3,5688,"s_nop 0",0,0,0,0,""
+3,5692,"s_nop 0",0,0,0,0,""
+3,5696,"s_nop 0",0,0,0,0,""
+3,5700,"s_nop 0",0,0,0,0,""
+3,5704,"s_nop 0",0,0,0,0,""
+3,5708,"s_nop 0",0,0,0,0,""
+3,5712,"s_nop 0",0,0,0,0,""
+3,5716,"s_nop 0",0,0,0,0,""
+3,5720,"s_nop 0",0,0,0,0,""
+3,5724,"s_nop 0",0,0,0,0,""
+3,5728,"s_nop 0",0,0,0,0,""
+3,5732,"s_nop 0",0,0,0,0,""
+3,5736,"s_nop 0",0,0,0,0,""
+3,5740,"s_nop 0",0,0,0,0,""
+3,5744,"s_nop 0",0,0,0,0,""
+3,5748,"s_nop 0",0,0,0,0,""
+3,5752,"s_nop 0",0,0,0,0,""
+3,5756,"s_nop 0",0,0,0,0,""
+3,5760,"s_nop 0",0,0,0,0,""
+3,5764,"s_nop 0",0,0,0,0,""
+3,5768,"s_nop 0",0,0,0,0,""
+3,5772,"s_nop 0",0,0,0,0,""
+3,5776,"s_nop 0",0,0,0,0,""
+3,5780,"s_nop 0",0,0,0,0,""
+3,5784,"s_nop 0",0,0,0,0,""
+3,5788,"s_nop 0",0,0,0,0,""
+3,5792,"s_nop 0",0,0,0,0,""
+3,5796,"s_nop 0",0,0,0,0,""
+3,5800,"s_nop 0",0,0,0,0,""
+3,5804,"s_nop 0",0,0,0,0,""
+3,5808,"s_nop 0",0,0,0,0,""
+3,5812,"s_nop 0",0,0,0,0,""
+3,5816,"s_nop 0",0,0,0,0,""
+3,5820,"s_nop 0",0,0,0,0,""
+3,5824,"s_nop 0",0,0,0,0,""
+3,5828,"s_nop 0",0,0,0,0,""
+3,5832,"s_nop 0",0,0,0,0,""
+3,5836,"s_nop 0",0,0,0,0,""
+3,5840,"s_nop 0",0,0,0,0,""
+3,5844,"s_nop 0",0,0,0,0,""
+3,5848,"s_nop 0",0,0,0,0,""
+3,5852,"s_nop 0",0,0,0,0,""
+3,5856,"s_nop 0",0,0,0,0,""
+3,5860,"s_nop 0",0,0,0,0,""
+3,5864,"s_nop 0",0,0,0,0,""
+3,5868,"s_nop 0",0,0,0,0,""
+3,5872,"s_nop 0",0,0,0,0,""
+3,5876,"s_nop 0",0,0,0,0,""
+3,5880,"s_nop 0",0,0,0,0,""
+3,5884,"s_nop 0",0,0,0,0,""
+3,5888,"v_and_b32_e32 v8, 0x3ff, v0",16,440,368,14396,""
+3,5896,"v_bfe_u32 v7, v0, 10, 10",16,220,156,0,""
+3,5904,"v_lshrrev_b32_e32 v0, 3, v8",16,216,152,0,""
+3,5908,"v_lshlrev_b32_e32 v1, 5, v7",16,336,272,0,""
+3,5912,"s_lshl_b32 s8, s8, 6",16,64,0,0,""
+3,5916,"v_or3_b32 v2, v0, v1, s8",16,140,76,0,""
+3,5924,"v_xor_b32_e32 v0, v0, v8",16,64,0,0,""
+3,5928,"v_lshlrev_b32_e32 v0, 3, v0",16,96,32,0,""
+3,5932,"v_and_b32_e32 v0, 56, v0",16,80,16,0,""
+3,5936,"v_lshrrev_b32_e32 v4, 6, v8",16,160,96,0,""
+3,5940,"v_lshl_or_b32 v6, v4, 3, v1",16,72,8,0,""
+3,5948,"v_mul_u32_u24_e32 v1, 0xe8, v2",16,88,24,0,""
+3,5956,"v_lshlrev_b32_e32 v2, 1, v0",16,124,60,0,""
+3,5960,"v_lshl_add_u32 v5, v1, 1, v2",16,184,120,0,""
+3,5968,"v_lshrrev_b32_e32 v2, 5, v8",16,64,0,0,""
+3,5972,"v_lshlrev_b32_e32 v3, 3, v7",16,88,24,0,""
+3,5976,"s_lshl_b32 s9, s9, 5",16,96,32,0,""
+3,5980,"v_or3_b32 v12, v2, v3, s9",16,136,72,0,""
+3,5988,"v_lshlrev_b32_e32 v9, 1, v8",16,64,0,0,""
+3,5992,"v_lshlrev_b32_e32 v2, 6, v2",16,64,0,0,""
+3,5996,"s_mov_b64 s[12:13], s[2:3]",16,80,16,0,""
+3,6000,"v_readfirstlane_b32 s2, v6",16,64,0,0,""
+3,6004,"s_movk_i32 s3, 0xe8",16,76,12,256,""
+3,6008,"v_sub_u32_e32 v2, v9, v2",16,64,0,0,""
+3,6012,"v_lshl_or_b32 v10, v4, 1, v3",16,68,4,0,""
+3,6020,"s_mov_b64 s[0:1], s[6:7]",16,64,0,0,""
+3,6024,"s_and_b32 s6, s13, 0xffff",16,124,60,0,""
+3,6032,"s_lshl_b32 s2, s2, 7",16,96,32,0,""
+3,6036,"v_readfirstlane_b32 s10, v10",16,80,16,0,""
+3,6040,"v_mad_u32_u24 v9, v12, s3, v2",16,68,4,0,""
+3,6048,"v_not_b32_e32 v3, -2.0",16,64,0,0,""
+3,6052,"v_cmp_gt_i32_e32 vcc, s3, v2",16,76,8,0,""
+3,6056,"s_or_b32 s13, s6, 0x40e80000",16,68,4,252,""
+3,6064,"s_mov_b32 s15, 0x27000",16,64,0,0,""
+3,6072,"s_mov_b32 s14, 0x7ffffffd",16,84,20,0,""
+3,6080,"s_add_i32 m0, s2, 0x2000",16,108,44,0,""
+3,6088,"s_and_b32 s5, s5, 0xffff",16,72,8,0,""
+3,6096,"v_cndmask_b32_e32 v9, v3, v9, vcc",16,76,8,0,""
+3,6100,"s_lshl_b32 s11, s10, 7",16,88,16,0,""
+3,6104,"buffer_load_dwordx4 v5, s[12:15], 0 offen lds",16,416,352,56,""
+3,6112,"s_or_b32 s5, s5, 0x40e80000",16,76,12,0,""
+3,6120,"s_mov_b32 s6, s14",16,64,0,0,""
+3,6124,"s_mov_b32 s7, s15",16,64,0,0,""
+3,6128,"v_lshlrev_b32_e32 v9, 1, v9",16,64,0,0,""
+3,6132,"s_add_i32 m0, s11, 0x5000",16,64,0,0,""
+3,6140,"v_or_b32_e32 v11, 16, v10",16,64,0,0,""
+3,6144,"buffer_load_dword v9, s[4:7], 0 offen lds",16,468,404,128,""
+3,6152,"v_mov_b32_e32 v9, 0xe80",16,64,0,0,""
+3,6160,"v_mad_u32_u24 v9, v12, s3, v9",16,148,80,0,""
+3,6168,"v_readfirstlane_b32 s10, v11",16,168,104,0,""
+3,6172,"v_add_u32_e32 v13, v9, v2",16,64,0,0,""
+3,6176,"v_cndmask_b32_e32 v13, v3, v13, vcc",16,68,0,0,""
+3,6180,"s_lshl_b32 s16, s10, 7",16,68,4,124,""
+3,6184,"v_lshlrev_b32_e32 v13, 1, v13",16,72,8,0,""
+3,6188,"s_add_i32 m0, s16, 0x5000",16,64,0,0,""
+3,6196,"s_movk_i32 s10, 0xa8",16,64,0,0,""
+3,6200,"buffer_load_dword v13, s[4:7], 0 offen lds",16,204,140,64,""
+3,6208,"v_add_u32_e32 v13, 0x80, v5",16,64,0,0,""
+3,6216,"s_mov_b32 m0, s2",16,68,0,0,""
+3,6220,"v_cmp_gt_i32_e32 vcc, s10, v2",16,200,116,0,""
+3,6224,"buffer_load_dwordx4 v13, s[12:15], 0 offen lds",16,320,256,108,""
+3,6232,"v_add_u32_e32 v13, 64, v2",16,64,0,0,""
+3,6236,"v_mad_u32_u24 v14, v12, s3, v13",16,72,8,0,""
+3,6244,"v_cndmask_b32_e32 v14, v3, v14, vcc",16,64,0,0,""
+3,6248,"v_add_u32_e32 v13, v9, v13",16,64,0,0,""
+3,6252,"v_lshlrev_b32_e32 v14, 1, v14",16,64,0,0,""
+3,6256,"s_add_i32 m0, s11, 0x4000",16,64,0,0,""
+3,6264,"v_cndmask_b32_e32 v13, v3, v13, vcc",16,80,0,0,""
+3,6268,"buffer_load_dword v14, s[4:7], 0 offen lds",16,184,120,112,""
+3,6276,"v_lshlrev_b32_e32 v13, 1, v13",16,64,0,0,""
+3,6280,"s_add_i32 m0, s16, 0x4000",16,64,0,0,""
+3,6288,"v_cmp_eq_u32_e64 s[2:3], 0, v7",16,80,0,0,""
+3,6296,"buffer_load_dword v13, s[4:7], 0 offen lds",16,212,148,112,""
+3,6304,"v_mul_u32_u24_e32 v16, 0xe8, v12",16,64,0,0,""
+3,6312,"s_and_b64 vcc, exec, s[2:3]",16,76,0,0,""
+3,6316,"s_cbranch_vccnz 1",16,68,0,52,""
+3,6320,"s_barrier",8,6416,6416,0,""
+3,6324,"v_readfirstlane_b32 s6, v6",16,444,380,124,""
+3,6328,"v_lshlrev_b32_e32 v6, 4, v4",16,64,0,0,""
+3,6332,"v_and_b32_e32 v4, 15, v8",16,112,44,0,""
+3,6336,"v_and_b32_e32 v12, 7, v8",16,236,172,28,""
+3,6340,"v_readfirstlane_b32 s7, v10",16,376,312,0,""
+3,6344,"v_or_b32_e32 v15, v6, v4",16,64,0,0,""
+3,6348,"v_lshrrev_b32_e32 v10, 4, v8",16,64,0,0,""
+3,6352,"v_lshlrev_b32_e32 v7, 4, v7",16,136,72,0,""
+3,6356,"v_bfe_u32 v8, v8, 4, 2",16,96,32,0,""
+3,6364,"v_bitop3_b32 v17, v10, v12, 3 bitop3:0x6c",16,64,0,0,""
+3,6372,"v_or_b32_e32 v26, v4, v7",16,64,0,0,""
+3,6376,"v_lshlrev_b32_e32 v28, 7, v15",16,160,96,0,""
+3,6380,"v_lshlrev_b32_e32 v10, 3, v17",16,80,16,0,""
+3,6384,"v_bitop3_b32 v27, v8, v12, 4 bitop3:0x36",16,64,0,0,""
+3,6392,"v_lshl_or_b32 v17, v17, 4, v28",16,72,8,0,""
+3,6400,"v_lshlrev_b32_e32 v12, 7, v26",16,112,48,0,""
+3,6404,"s_waitcnt vmcnt(0)",16,8420,8420,0,""
+3,6408,"s_barrier",16,464,464,0,""
+3,6412,"v_lshl_or_b32 v12, v8, 4, v12",16,280,200,0,""
+3,6420,"ds_read_b128 v[18:21], v17 offset:8192",16,152,84,112,""
+3,6428,"ds_read_b128 v[22:25], v12 offset:20480",16,152,84,0,""
+3,6436,"v_readfirstlane_b32 s11, v11",16,284,212,0,""
+3,6440,"v_lshlrev_b32_e32 v14, 3, v8",16,196,132,0,""
+3,6444,"v_lshlrev_b32_e32 v11, 3, v27",16,104,40,0,""
+3,6448,"v_lshlrev_b32_e32 v13, 6, v15",16,248,184,0,""
+3,6452,"v_lshlrev_b32_e32 v15, 6, v26",16,188,124,0,""
+3,6456,"s_lshl_b32 s16, s6, 7",16,64,0,0,""
+3,6460,"s_lshl_b32 s17, s7, 7",16,64,0,0,""
+3,6464,"s_lshl_b32 s11, s11, 7",16,64,0,0,""
+3,6468,"v_add_u32_e32 v5, 0x100, v5",16,208,144,0,""
+3,6476,"s_add_i32 m0, s16, 0x7000",16,64,0,0,""
+3,6484,"s_movk_i32 s6, 0x68",16,64,0,0,""
+3,6488,"buffer_load_dwordx4 v5, s[12:15], 0 offen lds",16,68,4,64,""
+3,6496,"v_add_u32_e32 v5, 0x80, v2",16,136,72,0,""
+3,6504,"v_add_u32_e32 v26, v5, v16",16,64,0,0,""
+3,6508,"v_cmp_gt_i32_e32 vcc, s6, v2",16,64,0,0,""
+3,6512,"v_add_u32_e32 v5, v5, v9",16,64,0,0,""
+3,6516,"s_add_i32 m0, s17, 0x6000",16,64,0,192,""
+3,6524,"v_cndmask_b32_e32 v26, v3, v26, vcc",16,64,0,0,""
+3,6528,"v_lshlrev_b32_e32 v26, 1, v26",16,64,0,0,""
+3,6532,"s_mov_b32 s6, s14",16,64,0,0,""
+3,6536,"s_mov_b32 s7, s15",16,64,0,0,""
+3,6540,"v_cndmask_b32_e32 v5, v3, v5, vcc",16,64,0,0,""
+3,6544,"buffer_load_dword v26, s[4:7], 0 offen lds",16,88,24,128,""
+3,6552,"v_lshlrev_b32_e32 v5, 1, v5",16,64,0,0,""
+3,6556,"s_add_i32 m0, s11, 0x6000",16,80,0,0,""
+3,6564,"s_nop 0",16,64,64,0,""
+3,6568,"buffer_load_dword v5, s[4:7], 0 offen lds",16,76,12,0,""
+3,6576,"s_barrier",16,1368,1368,0,""
+3,6580,"s_setprio 1",16,64,64,0,""
+3,6584,"v_mov_b32_e32 v5, 0",16,216,152,0,""
+3,6588,"s_waitcnt lgkmcnt(0)",16,64,64,0,""
+3,6592,"v_mfma_f32_16x16x32_bf16 v[18:21], v[18:21], v[22:25], 0",16,200,136,0,""
+3,6600,"s_setprio 0",16,64,64,0,""
+3,6604,"s_waitcnt vmcnt(3)",16,64,64,0,""
+3,6608,"s_barrier",16,3276,3276,0,""
+3,6612,"v_lshl_or_b32 v30, v27, 4, v28",16,128,36,0,""
+3,6620,"s_waitcnt vmcnt(0)",16,4324,4324,0,""
+3,6624,"ds_read_b128 v[22:25], v30 offset:8192",16,84,16,0,""
+3,6632,"ds_read_b128 v[26:29], v12 offset:20544",16,68,4,0,""
+3,6640,"s_waitcnt vmcnt(0)",16,64,64,0,""
+3,6644,"s_barrier",16,420,420,0,""
+3,6648,"s_setprio 1",16,64,64,0,""
+3,6652,"s_waitcnt lgkmcnt(0)",16,652,652,0,""
+3,6656,"v_mfma_f32_16x16x32_bf16 v[18:21], v[22:25], v[26:29], v[18:21]",16,92,28,0,""
+3,6664,"s_setprio 0",16,64,64,0,""
+3,6668,"s_waitcnt vmcnt(3)",16,64,64,0,""
+3,6672,"s_barrier",16,3984,3984,0,""
+3,6676,"ds_read_b128 v[22:25], v17",16,108,32,0,""
+3,6684,"ds_read_b128 v[26:29], v12 offset:16384",16,196,132,0,""
+3,6692,"v_or_b32_e32 v0, 0x80, v0",16,200,136,0,""
+3,6700,"v_add3_u32 v1, v1, v0, 64",16,112,48,0,""
+3,6708,"v_cmp_gt_u32_e32 vcc, s10, v0",16,200,136,0,""
+3,6712,"s_add_i32 m0, s16, 0x2000",16,64,0,256,""
+3,6720,"s_nop 0",16,64,64,0,""
+3,6724,"v_cndmask_b32_e32 v0, v3, v1, vcc",16,100,36,0,""
+3,6728,"v_lshlrev_b32_e32 v0, 1, v0",16,140,76,0,""
+3,6732,"buffer_load_dwordx4 v0, s[12:15], 0 offen lds",16,92,24,128,""
+3,6740,"v_add_u32_e32 v0, 0xc0, v2",16,248,180,0,""
+3,6748,"v_add_u32_e32 v1, v0, v16",16,152,88,0,""
+3,6752,"v_cmp_gt_i32_e32 vcc, 40, v2",16,124,60,0,""
+3,6756,"v_add_u32_e32 v0, v0, v9",16,176,112,0,""
+3,6760,"s_add_i32 m0, s17, 0x5000",16,68,0,152,""
+3,6768,"v_cndmask_b32_e32 v1, v3, v1, vcc",16,132,68,0,""
+3,6772,"v_lshlrev_b32_e32 v1, 1, v1",16,64,0,0,""
+3,6776,"v_cndmask_b32_e32 v0, v3, v0, vcc",16,64,0,0,""
+3,6780,"buffer_load_dword v1, s[4:7], 0 offen lds",16,76,12,128,""
+3,6788,"v_lshlrev_b32_e32 v0, 1, v0",16,64,0,0,""
+3,6792,"s_add_i32 m0, s11, 0x5000",16,80,0,0,""
+3,6800,"s_nop 0",16,64,64,0,""
+3,6804,"buffer_load_dword v0, s[4:7], 0 offen lds",16,84,20,0,""
+3,6812,"s_barrier",16,360,360,0,""
+3,6816,"s_setprio 1",16,64,64,0,""
+3,6820,"s_waitcnt lgkmcnt(0)",16,64,64,0,""
+3,6824,"v_mfma_f32_16x16x32_bf16 v[0:3], v[22:25], v[26:29], v[18:21]",16,80,16,0,""
+3,6832,"s_setprio 0",16,64,64,0,""
+3,6836,"s_waitcnt vmcnt(3)",16,64,64,0,""
+3,6840,"s_barrier",16,1824,1824,0,""
+3,6844,"s_waitcnt vmcnt(0)",16,1236,1236,0,""
+3,6848,"ds_read_b128 v[16:19], v30",16,96,24,0,""
+3,6856,"ds_read_b128 v[20:23], v12 offset:16448",16,128,64,0,""
+3,6864,"s_waitcnt vmcnt(0)",16,64,64,0,""
+3,6868,"s_barrier",16,300,300,0,""
+3,6872,"s_setprio 1",16,64,64,0,""
+3,6876,"s_waitcnt lgkmcnt(0)",16,780,780,0,""
+3,6880,"v_mfma_f32_16x16x32_bf16 v[0:3], v[16:19], v[20:23], v[0:3]",16,64,0,0,""
+3,6888,"s_setprio 0",16,64,64,0,""
+3,6892,"s_andn2_b64 vcc, exec, s[2:3]",16,72,0,0,""
+3,6896,"s_cbranch_vccnz 1",16,68,4,56,""
+3,6900,"s_barrier",8,832,832,0,""
+3,6904,"v_or_b32_e32 v10, v10, v13",16,100,36,128,""
+3,6908,"v_lshlrev_b32_e32 v10, 1, v10",16,112,48,0,""
+3,6912,"s_waitcnt vmcnt(3)",16,128,128,0,""
+3,6916,"s_barrier",16,1188,1188,0,""
+3,6920,"v_or_b32_e32 v9, v14, v15",16,80,16,0,""
+3,6924,"ds_read_b128 v[14:17], v10 offset:28672",16,128,64,128,""
+3,6932,"v_or_b32_e32 v11, v11, v13",16,232,168,0,""
+3,6936,"v_lshlrev_b32_e32 v11, 1, v11",16,224,160,0,""
+3,6940,"v_lshlrev_b32_e32 v9, 1, v9",16,132,68,0,""
+3,6944,"ds_read_b128 v[18:21], v11 offset:28672",16,148,84,128,""
+3,6952,"ds_read_b128 v[22:25], v9 offset:24576",16,184,120,0,""
+3,6960,"ds_read_b128 v[26:29], v12 offset:24640",16,236,172,0,""
+3,6968,"s_waitcnt lgkmcnt(1)",16,2064,2064,0,""
+3,6972,"v_mfma_f32_16x16x32_bf16 v[0:3], v[14:17], v[22:25], v[0:3]",16,64,0,0,""
+3,6980,"s_waitcnt vmcnt(0)",16,64,64,0,""
+3,6984,"ds_read_b128 v[14:17], v10 offset:8192",16,636,572,0,""
+3,6992,"v_lshl_or_b32 v8, v8, 2, s8",16,64,0,0,""
+3,7000,"s_waitcnt lgkmcnt(1)",16,168,168,0,""
+3,7004,"v_mfma_f32_16x16x32_bf16 v[0:3], v[18:21], v[26:29], v[0:3]",16,64,0,0,""
+3,7012,"ds_read_b128 v[18:21], v11 offset:8192",16,400,336,0,""
+3,7020,"ds_read_b128 v[22:25], v9 offset:20480",16,424,360,0,""
+3,7028,"ds_read_b128 v[10:13], v12 offset:20544",16,512,448,0,""
+3,7036,"v_or_b32_e32 v6, v8, v6",16,96,32,0,""
+3,7040,"v_lshl_or_b32 v6, v6, 8, s9",16,88,24,0,""
+3,7048,"s_waitcnt lgkmcnt(1)",16,2352,2352,0,""
+3,7052,"v_mfma_f32_16x16x32_bf16 v[0:3], v[14:17], v[22:25], v[0:3]",16,76,12,0,""
+3,7060,"v_or3_b32 v4, v6, v4, v7",16,180,116,0,""
+3,7068,"v_lshl_add_u64 v[4:5], v[4:5], 2, s[0:1]",16,64,0,0,""
+3,7076,"s_waitcnt lgkmcnt(0)",16,296,296,0,""
+3,7080,"v_mfma_f32_16x16x32_bf16 v[0:3], v[18:21], v[10:13], v[0:3]",16,64,0,0,""
+3,7088,"s_nop 7",16,512,512,0,""
+3,7092,"global_store_dword v[4:5], v0, off",16,168,88,0,""
+3,7100,"global_store_dword v[4:5], v1, off offset:1024",16,144,64,0,""
+3,7108,"global_store_dword v[4:5], v2, off offset:2048",16,136,56,0,""
+3,7116,"global_store_dword v[4:5], v3, off offset:3072",16,128,48,0,""
+3,7124,"s_endpgm",16,4076,0,0,""
Original file line number	Diff line number	Diff line change
		@@ -0,0 +1,2 @@
		"Name","Calls","TotalDurationNs","AverageNs","Percentage","MinNs","MaxNs","StdDev"
		"KERNEL_DISPATCH",1040,4927545,4738.024038,100.00,640,23400,1058.615401