diff --git a/aiter/configs/a8w8_blockscale_group32_tuned_gemm.csv b/aiter/configs/a8w8_blockscale_group32_tuned_gemm.csv new file mode 100644 index 00000000000..01f66480f90 --- /dev/null +++ b/aiter/configs/a8w8_blockscale_group32_tuned_gemm.csv @@ -0,0 +1 @@ +gfx,cu_num,M,N,K,libtype,kernelId,splitK,us,kernelName,tflops,bw,errRatio diff --git a/aiter/configs/a8w8_blockscale_group32_untuned_gemm.csv b/aiter/configs/a8w8_blockscale_group32_untuned_gemm.csv new file mode 100644 index 00000000000..a1b223a5c73 --- /dev/null +++ b/aiter/configs/a8w8_blockscale_group32_untuned_gemm.csv @@ -0,0 +1 @@ +M,N,K diff --git a/aiter/configs/model_configs/dsv41_a8w8_blockscale_group32_tuned_gemm.csv b/aiter/configs/model_configs/dsv41_a8w8_blockscale_group32_tuned_gemm.csv new file mode 100644 index 00000000000..52bf0adc068 --- /dev/null +++ b/aiter/configs/model_configs/dsv41_a8w8_blockscale_group32_tuned_gemm.csv @@ -0,0 +1,1225 @@ +gfx,cu_num,M,N,K,libtype,kernelId,splitK,us,kernelName,tflops,bw,errRatio +gfx950,256,1,1792,5120,triton,0,0,5.772394,triton,3.18,1590.98,0.0 +gfx950,256,2,1792,5120,triton,0,0,6.28401,triton,5.84,1462.83,0.0 +gfx950,256,4,1792,5120,triton,0,0,6.2415,triton,11.76,1475.58,0.00013950892857142856 +gfx950,256,8,1792,5120,triton,0,0,6.504469,triton,22.57,1421.28,6.975446428571428e-05 +gfx950,256,16,1792,5120,triton,0,0,6.18385,triton,47.48,1506.23,0.00027901785714285713 +gfx950,256,32,1792,5120,triton,0,0,6.34947,triton,92.48,1488.88,0.0001220703125 +gfx950,256,48,1792,5120,triton,0,0,7.37666,triton,119.4,1300.43,8.138020833333333e-05 +gfx950,256,64,1792,5120,triton,0,0,7.542091,triton,155.71,1290.37,0.0001220703125 +gfx950,256,80,1792,5120,triton,0,0,8.44626,triton,173.81,1168.73,0.0001255580357142857 +gfx950,256,96,1792,5120,triton,0,0,8.887032,triton,198.22,1126.43,0.00013369605654761906 +gfx950,256,112,1792,5120,triton,0,0,9.682392,triton,212.26,1048.28,0.00014449139030612245 +gfx950,256,128,1792,5120,triton,0,0,9.853224,triton,238.38,1044.24,0.00014822823660714285 +gfx950,256,144,1792,5120,triton,0,0,10.331372,triton,255.77,1009.39,0.0001666356646825397 +gfx950,256,160,1792,5120,triton,0,0,10.695155,triton,274.52,988.08,0.00010811941964285714 +gfx950,256,176,1792,5120,triton,0,0,10.992571,triton,293.8,974.02,0.00012999695616883116 +gfx950,256,192,1792,5120,triton,0,0,11.06039,triton,318.54,980.64,0.00013950892857142856 +gfx950,256,208,1792,5120,triton,0,0,11.768394,triton,324.33,933.47,0.00013950892857142856 +gfx950,256,224,1792,5120,triton,0,0,12.08551,triton,340.11,920.5,0.00013203523596938777 +gfx950,256,240,1792,5120,triton,0,0,12.363163,triton,356.22,911.09,0.00013020833333333333 +gfx950,256,256,1792,5120,triton,0,0,12.331529,triton,380.94,924.72,0.000152587890625 +gfx950,256,288,1792,5120,triton,0,0,14.121152,triton,374.25,827.25,0.00015113467261904762 +gfx950,256,320,1792,5120,triton,0,0,14.727031,triton,398.72,812.13,0.000146484375 +gfx950,256,352,1792,5120,triton,0,0,15.550671,triton,415.37,787.03,0.00014267958603896103 +gfx950,256,384,1792,5120,triton,0,0,15.723719,triton,448.14,796.08,0.00014096214657738094 +gfx950,256,416,1792,5120,triton,0,0,16.908234,triton,451.47,756.79,0.00017572759271978022 +gfx950,256,448,1792,5120,triton,0,0,18.04957,triton,455.46,724.36,0.00015196508290816327 +gfx950,256,480,1792,5120,triton,0,0,17.461969,triton,504.41,764.69,0.00015694754464285715 +gfx950,256,512,1792,5120,triton,0,0,17.664244,triton,531.88,771.7,0.00015585763113839285 +gfx950,256,544,1792,5120,triton,0,0,17.079139,triton,584.48,814.44,0.00013950892857142856 +gfx950,256,576,1792,5120,triton,0,0,17.764215,triton,595.0,798.71,0.00015307229662698413 +gfx950,256,608,1792,5120,triton,0,0,23.36515,triton,477.5,619.17,0.00013859110667293234 +gfx950,256,640,1792,5120,triton,0,0,24.129,triton,486.72,611.12,0.00015433175223214285 +gfx950,256,672,1792,5120,triton,0,0,25.42071,triton,485.09,591.02,0.0001635908269557823 +gfx950,256,704,1792,5120,triton,0,0,25.700417,triton,502.66,595.42,0.00014664290787337662 +gfx950,256,736,1792,5120,triton,0,0,24.940213,triton,541.52,624.74,0.0001660459530279503 +gfx950,256,768,1792,5120,triton,0,0,25.404585,triton,554.74,624.29,0.00013151622953869047 +gfx950,256,800,1792,5120,triton,0,0,25.851619,triton,567.86,624.26,0.00014020647321428571 +gfx950,256,832,1792,5120,triton,0,0,26.97682,triton,565.94,608.55,0.00014152107657967034 +gfx950,256,864,1792,5120,triton,0,0,25.709439,triton,616.68,649.38,0.0001453218005952381 +gfx950,256,896,1792,5120,triton,0,0,26.392725,triton,622.96,643.13,0.00015009665975765306 +gfx950,256,928,1792,5120,triton,0,0,26.295793,triton,647.59,656.09,0.0001515355603448276 +gfx950,256,960,1792,5120,triton,0,0,25.982914,triton,677.99,674.71,0.00015404110863095238 +gfx950,256,992,1792,5120,triton,0,0,26.672213,triton,682.48,667.71,0.00015019711261520736 +gfx950,256,1024,1792,5120,triton,0,0,26.83306,triton,700.27,674.09,0.00014005388532366072 +gfx950,256,1088,1792,5120,triton,0,0,27.32162,triton,730.74,682.43,0.00015079273897058824 +gfx950,256,1152,1792,5120,triton,0,0,27.738827,triton,762.08,692.24,0.00014919704861111112 +gfx950,256,1216,1792,5120,triton,0,0,33.265929,triton,670.77,593.97,0.00013950892857142856 +gfx950,256,1280,1792,5120,triton,0,0,35.520265,triton,661.26,571.96,0.0001338413783482143 +gfx950,256,1344,1792,5120,triton,0,0,33.944212,triton,726.56,614.93,0.00014283056972789115 +gfx950,256,1408,1792,5120,triton,0,0,36.188506,triton,713.95,592.18,0.00015338055499188312 +gfx950,256,1472,1792,5120,triton,0,0,34.456657,triton,783.92,638.12,0.00013003141983695653 +gfx950,256,1536,1792,5120,triton,0,0,35.872143,triton,785.73,628.46,0.00015077136811755953 +gfx950,256,1600,1792,5120,triton,0,0,34.901253,triton,841.23,661.91,0.00014508928571428572 +gfx950,256,1664,1792,5120,triton,0,0,36.905216,triton,827.38,641.06,0.00014185643458104397 +gfx950,256,1728,1792,5120,triton,0,0,35.7737,triton,886.38,676.91,0.00014241536458333334 +gfx950,256,1792,1792,5120,triton,0,0,37.172783,triton,884.61,666.42,0.00014822823660714285 +gfx950,256,1856,1792,5120,triton,0,0,36.7077,triton,927.81,690.04,0.00013860693118842365 +gfx950,256,1920,1792,5120,triton,0,0,38.79932,triton,908.06,667.2,0.00014706566220238094 +gfx950,256,1984,1792,5120,triton,0,0,36.7245,triton,991.34,720.06,0.00015722881264400922 +gfx950,256,2048,1792,5120,triton,0,0,38.12491,triton,985.73,708.22,0.00014986310686383928 +gfx950,256,4096,1792,5120,triton,0,0,63.13672,triton,1190.46,709.99,0.00014032636369977678 +gfx950,256,6144,1792,5120,triton,0,0,99.52143,triton,1132.85,629.54,0.00013978140694754465 +gfx950,256,8192,1792,5120,triton,0,0,119.48265,triton,1258.12,673.56,0.0001452990940638951 +gfx950,256,10240,1792,5120,triton,0,0,152.63348,triton,1231.09,644.05,0.0001502445765904018 +gfx950,256,12288,1792,5120,triton,0,0,169.4947,triton,1330.34,685.15,0.00013991764613560269 +gfx950,256,14336,1792,5120,triton,0,0,209.373368,triton,1256.45,639.79,0.00014842286401865434 +gfx950,256,16384,1792,5120,triton,0,0,228.434388,triton,1316.12,664.44,0.00014730862208775112 +gfx950,256,32768,1792,5120,triton,0,0,433.672842,triton,1386.52,678.82,0.00014674663543701172 +gfx950,256,1,4096,1280,triton,0,0,3.814374,triton,2.75,1376.99,0.0 +gfx950,256,2,4096,1280,triton,0,0,3.829663,triton,5.48,1373.97,0.0001220703125 +gfx950,256,4,4096,1280,triton,0,0,3.8326,triton,10.94,1377.86,0.0 +gfx950,256,8,4096,1280,triton,0,0,3.96221,triton,21.17,1342.35,9.1552734375e-05 +gfx950,256,16,4096,1280,triton,0,0,4.10144,triton,40.91,1315.25,1.52587890625e-05 +gfx950,256,32,4096,1280,triton,0,0,5.181263,triton,64.76,1070.39,2.288818359375e-05 +gfx950,256,48,4096,1280,triton,0,0,5.370061,triton,93.73,1060.98,2.5431315104166668e-05 +gfx950,256,64,4096,1280,triton,0,0,5.405525,triton,124.15,1082.06,6.103515625e-05 +gfx950,256,80,4096,1280,triton,0,0,5.906889,triton,142.01,1015.87,3.662109375e-05 +gfx950,256,96,4096,1280,triton,0,0,5.92265,triton,169.96,1038.76,2.7974446614583332e-05 +gfx950,256,112,4096,1280,triton,0,0,5.901242,triton,199.01,1068.21,3.705705915178571e-05 +gfx950,256,128,4096,1280,triton,0,0,5.968041,triton,224.89,1081.64,4.38690185546875e-05 +gfx950,256,144,4096,1280,triton,0,0,6.99386,triton,215.9,944.66,4.4080946180555555e-05 +gfx950,256,160,4096,1280,triton,0,0,7.576848,triton,221.43,891.98,5.4931640625e-05 +gfx950,256,176,4096,1280,triton,0,0,7.637837,triton,241.63,904.7,4.022771661931818e-05 +gfx950,256,192,4096,1280,triton,0,0,7.608384,triton,264.61,928.12,4.1961669921875e-05 +gfx950,256,208,4096,1280,triton,0,0,7.097175,triton,307.31,1016.33,3.990760216346154e-05 +gfx950,256,224,4096,1280,triton,0,0,7.166333,triton,327.76,1027.67,3.378731863839286e-05 +gfx950,256,240,4096,1280,triton,0,0,7.172365,triton,350.87,1047.93,2.9500325520833334e-05 +gfx950,256,256,4096,1280,triton,0,0,7.238273,triton,370.86,1059.33,2.86102294921875e-05 +gfx950,256,288,4096,1280,triton,0,0,8.589374,triton,351.59,927.99,2.8822157118055555e-05 +gfx950,256,320,4096,1280,triton,0,0,8.787695,triton,381.83,941.53,3.204345703125e-05 +gfx950,256,352,4096,1280,triton,0,0,9.049547,triton,407.86,947.78,4.022771661931818e-05 +gfx950,256,384,4096,1280,triton,0,0,8.95907,triton,449.44,991.19,3.3060709635416664e-05 +gfx950,256,416,4096,1280,triton,0,0,8.553188,triton,509.99,1073.66,3.579946664663462e-05 +gfx950,256,448,4096,1280,triton,0,0,8.657408,triton,542.61,1095.75,3.705705915178571e-05 +gfx950,256,480,4096,1280,triton,0,0,8.740191,triton,575.86,1120.05,3.763834635416667e-05 +gfx950,256,512,4096,1280,triton,0,0,8.89803,triton,603.36,1134.24,3.3855438232421875e-05 +gfx950,256,544,4096,1280,triton,0,0,11.320409,triton,503.89,918.31,3.141515395220588e-05 +gfx950,256,576,4096,1280,triton,0,0,11.408352,triton,529.42,937.8,3.814697265625e-05 +gfx950,256,608,4096,1280,triton,0,0,11.652618,triton,547.12,944.15,3.4934596011513155e-05 +gfx950,256,640,4096,1280,triton,0,0,11.694172,triton,573.87,966.72,3.62396240234375e-05 +gfx950,256,672,4096,1280,triton,0,0,11.645214,triton,605.09,996.81,3.9963495163690477e-05 +gfx950,256,704,4096,1280,triton,0,0,11.820032,triton,624.53,1007.71,3.502585671164773e-05 +gfx950,256,736,4096,1280,triton,0,0,11.371208,triton,678.69,1074.14,3.781525985054348e-05 +gfx950,256,768,4096,1280,triton,0,0,11.583454,triton,695.22,1080.63,3.4332275390625e-05 +gfx950,256,800,4096,1280,triton,0,0,11.662385,triton,719.29,1099.3,3.11279296875e-05 +gfx950,256,832,4096,1280,triton,0,0,12.225204,triton,713.62,1073.49,3.814697265625e-05 +gfx950,256,864,4096,1280,triton,0,0,12.423433,triton,729.24,1080.76,3.9842393663194445e-05 +gfx950,256,896,4096,1280,triton,0,0,12.362333,triton,759.99,1110.61,3.62396240234375e-05 +gfx950,256,928,4096,1280,triton,0,0,11.917926,triton,816.48,1177.46,4.0514715786637934e-05 +gfx950,256,960,4096,1280,triton,0,0,12.140525,triton,829.15,1180.84,3.8909912109375e-05 +gfx950,256,992,4096,1280,triton,0,0,12.104552,triton,859.34,1209.39,3.6178096648185486e-05 +gfx950,256,1024,4096,1280,triton,0,0,12.336411,triton,870.38,1211.23,3.933906555175781e-05 +gfx950,256,1088,4096,1280,triton,0,0,16.75749,triton,680.8,927.85,3.837136661305147e-05 +gfx950,256,1152,4096,1280,triton,0,0,17.541884,triton,688.61,920.92,3.4968058268229164e-05 +gfx950,256,1216,4096,1280,triton,0,0,17.721283,triton,719.51,945.8,3.834774619654605e-05 +gfx950,256,1280,4096,1280,triton,0,0,17.861969,triton,751.42,972.29,3.92913818359375e-05 +gfx950,256,1344,4096,1280,triton,0,0,18.02202,triton,781.98,997.29,3.796532040550595e-05 +gfx950,256,1408,4096,1280,triton,0,0,18.243299,triton,809.28,1018.43,3.485246138139205e-05 +gfx950,256,1472,4096,1280,triton,0,0,18.475545,triton,835.43,1038.44,3.7483547044836954e-05 +gfx950,256,1536,4096,1280,triton,0,0,18.69876,triton,861.35,1058.46,3.910064697265625e-05 +gfx950,256,1600,4096,1280,triton,0,0,18.919563,triton,886.77,1078.15,3.9215087890625e-05 +gfx950,256,1664,4096,1280,triton,0,0,19.407774,triton,899.04,1082.26,3.5506028395432695e-05 +gfx950,256,1728,4096,1280,triton,0,0,19.844165,triton,913.08,1089.01,3.9842393663194445e-05 +gfx950,256,1792,4096,1280,triton,0,0,19.333258,triton,971.93,1149.14,3.937312534877232e-05 +gfx950,256,1856,4096,1280,triton,0,0,19.352608,triton,1005.63,1179.32,4.130396349676724e-05 +gfx950,256,1920,4096,1280,triton,0,0,19.575333,triton,1028.47,1196.87,3.7892659505208334e-05 +gfx950,256,1984,4096,1280,triton,0,0,19.679184,triton,1057.14,1221.36,3.888530115927419e-05 +gfx950,256,2048,4096,1280,triton,0,0,20.237566,triton,1061.14,1217.61,3.719329833984375e-05 +gfx950,256,4096,4096,1280,triton,0,0,37.16531,triton,1155.64,1184.98,3.403425216674805e-05 +gfx950,256,6144,4096,1280,triton,0,0,57.02587,triton,1129.74,1112.46,3.838539123535156e-05 +gfx950,256,8192,4096,1280,triton,0,0,78.83527,triton,1089.61,1050.77,3.695487976074219e-05 +gfx950,256,10240,4096,1280,triton,0,0,90.54016,triton,1185.93,1129.18,3.478527069091797e-05 +gfx950,256,12288,4096,1280,triton,0,0,109.34511,triton,1178.37,1112.39,3.820657730102539e-05 +gfx950,256,14336,4096,1280,triton,0,0,126.682135,triton,1186.62,1113.29,3.8351331438337055e-05 +gfx950,256,16384,4096,1280,triton,0,0,146.63558,triton,1171.6,1094.09,3.9264559745788574e-05 +gfx950,256,32768,4096,1280,triton,0,0,286.04948,triton,1201.18,1103.38,3.746896982192993e-05 +gfx950,256,1,8192,1280,triton,0,0,4.466886,triton,4.69,2351.4,0.0001220703125 +gfx950,256,2,8192,1280,triton,0,0,4.914773,triton,8.53,2140.71,0.0 +gfx950,256,4,8192,1280,triton,0,0,5.219857,triton,16.07,2022.36,0.0 +gfx950,256,8,8192,1280,triton,0,0,5.116134,triton,32.79,2077.17,3.0517578125e-05 +gfx950,256,16,8192,1280,triton,0,0,5.112576,triton,65.63,2106.25,6.866455078125e-05 +gfx950,256,32,8192,1280,triton,0,0,5.514659,triton,121.69,2003.93,4.1961669921875e-05 +gfx950,256,48,8192,1280,triton,0,0,6.053976,triton,166.28,1872.1,2.7974446614583332e-05 +gfx950,256,64,8192,1280,triton,0,0,6.132369,triton,218.87,1894.25,3.62396240234375e-05 +gfx950,256,80,8192,1280,triton,0,0,7.357326,triton,228.03,1617.28,2.899169921875e-05 +gfx950,256,96,8192,1280,triton,0,0,7.901419,triton,254.8,1541.69,3.9418538411458336e-05 +gfx950,256,112,8192,1280,triton,0,0,7.975747,triton,294.49,1562.75,3.814697265625e-05 +gfx950,256,128,8192,1280,triton,0,0,7.990957,triton,335.92,1595.15,3.4332275390625e-05 +gfx950,256,144,8192,1280,triton,0,0,8.709891,triton,346.72,1495.93,4.0690104166666664e-05 +gfx950,256,160,8192,1280,triton,0,0,8.93948,triton,375.35,1489.12,3.4332275390625e-05 +gfx950,256,176,8192,1280,triton,0,0,8.856716,triton,416.74,1534.95,3.259832208806818e-05 +gfx950,256,192,8192,1280,triton,0,0,8.979303,triton,448.42,1545.47,3.4968058268229164e-05 +gfx950,256,208,8192,1280,triton,0,0,8.997783,triton,484.79,1573.71,4.2841984675480766e-05 +gfx950,256,224,8192,1280,triton,0,0,9.589702,triton,489.86,1506.04,3.215244838169643e-05 +gfx950,256,240,8192,1280,triton,0,0,9.682161,triton,519.84,1520.85,3.712972005208333e-05 +gfx950,256,256,8192,1280,triton,0,0,9.757515,triton,550.21,1538.07,3.6716461181640625e-05 +gfx950,256,288,8192,1280,triton,0,0,11.39811,triton,529.89,1366.28,3.4332275390625e-05 +gfx950,256,320,8192,1280,triton,0,0,11.752802,triton,571.0,1373.14,3.204345703125e-05 +gfx950,256,352,8192,1280,triton,0,0,12.124062,triton,608.87,1377.71,3.537264737215909e-05 +gfx950,256,384,8192,1280,triton,0,0,11.824882,triton,681.03,1460.37,4.4186909993489586e-05 +gfx950,256,416,8192,1280,triton,0,0,11.96771,triton,728.97,1490.18,3.1104454627403844e-05 +gfx950,256,448,8192,1280,triton,0,0,12.560333,triton,748.01,1464.87,3.841945103236607e-05 +gfx950,256,480,8192,1280,triton,0,0,12.784473,triton,787.39,1483.4,3.636678059895833e-05 +gfx950,256,512,8192,1280,triton,0,0,12.831011,triton,836.83,1522.07,3.62396240234375e-05 +gfx950,256,544,8192,1280,triton,0,0,16.761012,triton,680.66,1198.91,3.792257869944853e-05 +gfx950,256,576,8192,1280,triton,0,0,17.134058,triton,705.0,1205.8,3.6451551649305555e-05 +gfx950,256,608,8192,1280,triton,0,0,17.38286,triton,733.52,1221.06,3.39307283100329e-05 +gfx950,256,640,8192,1280,triton,0,0,17.888067,triton,750.32,1218.17,4.08172607421875e-05 +gfx950,256,672,8192,1280,triton,0,0,17.473281,triton,806.54,1279.44,3.6148797898065477e-05 +gfx950,256,704,8192,1280,triton,0,0,17.96878,triton,821.64,1275.61,3.866715864701705e-05 +gfx950,256,736,8192,1280,triton,0,0,17.952604,triton,859.77,1308.25,3.466398819633152e-05 +gfx950,256,768,8192,1280,triton,0,0,19.141926,triton,841.41,1256.49,4.212061564127604e-05 +gfx950,256,800,8192,1280,triton,0,0,17.833423,triton,940.77,1380.38,3.8909912109375e-05 +gfx950,256,832,8192,1280,triton,0,0,18.219582,triton,957.67,1382.15,4.049447866586538e-05 +gfx950,256,864,8192,1280,triton,0,0,19.486466,triton,929.85,1321.3,3.334327980324074e-05 +gfx950,256,896,8192,1280,triton,0,0,18.321713,triton,1025.59,1436.15,3.528594970703125e-05 +gfx950,256,928,8192,1280,triton,0,0,19.536736,triton,996.15,1375.76,3.9593926791487066e-05 +gfx950,256,960,8192,1280,triton,0,0,18.938713,triton,1063.04,1449.05,3.47137451171875e-05 +gfx950,256,992,8192,1280,triton,0,0,19.198695,triton,1083.6,1458.87,3.85161369077621e-05 +gfx950,256,1024,8192,1280,triton,0,0,20.404518,triton,1052.45,1400.36,3.802776336669922e-05 +gfx950,256,1088,8192,1280,triton,0,0,29.23165,triton,780.56,1016.17,3.343469956341912e-05 +gfx950,256,1152,8192,1280,triton,0,0,29.75485,triton,811.94,1036.29,3.7511189778645836e-05 +gfx950,256,1216,8192,1280,triton,0,0,29.859041,triton,854.06,1070.54,3.5034982781661184e-05 +gfx950,256,1280,8192,1280,triton,0,0,30.235606,triton,887.81,1094.59,3.47137451171875e-05 +gfx950,256,1344,8192,1280,triton,0,0,31.41326,triton,897.26,1089.55,3.3696492513020836e-05 +gfx950,256,1408,8192,1280,triton,0,0,31.44566,triton,939.01,1124.37,4.014101895419034e-05 +gfx950,256,1472,8192,1280,triton,0,0,32.77167,triton,941.97,1113.38,4.055189049762228e-05 +gfx950,256,1536,8192,1280,triton,0,0,32.16887,triton,1001.35,1169.38,3.711382548014323e-05 +gfx950,256,1600,8192,1280,triton,0,0,34.96173,triton,959.75,1108.3,3.570556640625e-05 +gfx950,256,1664,8192,1280,triton,0,0,35.9093,triton,971.8,1110.54,3.513923058143029e-05 +gfx950,256,1728,8192,1280,triton,0,0,36.55251,triton,991.42,1121.92,3.8782755533854164e-05 +gfx950,256,1792,8192,1280,triton,0,0,37.60011,triton,999.49,1120.73,3.7533896309988836e-05 +gfx950,256,1856,8192,1280,triton,0,0,38.04692,triton,1023.03,1137.28,3.716041301858836e-05 +gfx950,256,1920,8192,1280,triton,0,0,38.53256,triton,1044.97,1152.29,3.7892659505208334e-05 +gfx950,256,1984,8192,1280,triton,0,0,39.92578,triton,1042.12,1140.39,3.814697265625e-05 +gfx950,256,2048,8192,1280,triton,0,0,39.48057,triton,1087.87,1181.89,3.695487976074219e-05 +gfx950,256,4096,8192,1280,triton,0,0,74.69009,triton,1150.08,1109.08,3.713369369506836e-05 +gfx950,256,6144,8192,1280,triton,0,0,110.76125,triton,1163.3,1074.5,3.4749507904052734e-05 +gfx950,256,8192,8192,1280,triton,0,0,146.48743,triton,1172.79,1059.4,3.7282705307006836e-05 +gfx950,256,10240,8192,1280,triton,0,0,173.563724,triton,1237.29,1102.56,3.601312637329102e-05 +gfx950,256,12288,8192,1280,triton,0,0,209.773124,triton,1228.46,1084.7,3.7044286727905273e-05 +gfx950,256,14336,8192,1280,triton,0,0,244.567081,triton,1229.31,1078.3,3.5490308489118305e-05 +gfx950,256,16384,8192,1280,triton,0,0,284.815122,triton,1206.39,1052.94,3.6597251892089844e-05 +gfx950,256,32768,8192,1280,triton,0,0,528.989444,triton,1299.07,1114.01,3.542378544807434e-05 +gfx950,256,1,16384,1280,triton,0,0,5.975314,triton,7.02,3515.39,6.103515625e-05 +gfx950,256,2,16384,1280,triton,0,0,5.978816,triton,14.03,3519.03,6.103515625e-05 +gfx950,256,4,16384,1280,triton,0,0,5.93108,triton,28.29,3558.83,1.52587890625e-05 +gfx950,256,8,16384,1280,triton,0,0,6.064045,triton,55.33,3503.26,2.288818359375e-05 +gfx950,256,16,16384,1280,triton,0,0,6.133465,triton,109.41,3508.02,3.0517578125e-05 +gfx950,256,32,16384,1280,triton,0,0,6.85536,triton,195.79,3218.07,4.1961669921875e-05 +gfx950,256,48,16384,1280,triton,0,0,8.080359,triton,249.16,2797.63,3.5603841145833336e-05 +gfx950,256,64,16384,1280,triton,0,0,8.171929,triton,328.48,2832.94,2.288818359375e-05 +gfx950,256,80,16384,1280,triton,0,0,8.841872,triton,379.49,2679.9,3.7384033203125e-05 +gfx950,256,96,16384,1280,triton,0,0,9.113369,triton,441.83,2659.84,4.5140584309895836e-05 +gfx950,256,112,16384,1280,triton,0,0,9.539557,triton,492.44,2598.12,2.7792794363839285e-05 +gfx950,256,128,16384,1280,triton,0,0,9.9666,triton,538.67,2541.45,3.8623809814453125e-05 +gfx950,256,144,16384,1280,triton,0,0,11.858663,triton,509.32,2181.9,3.857082790798611e-05 +gfx950,256,160,16384,1280,triton,0,0,12.445386,triton,539.23,2122.81,2.899169921875e-05 +gfx950,256,176,16384,1280,triton,0,0,12.562384,triton,587.63,2146.41,3.884055397727273e-05 +gfx950,256,192,16384,1280,triton,0,0,12.469663,triton,645.81,2206.05,3.6557515462239586e-05 +gfx950,256,208,16384,1280,triton,0,0,12.227935,triton,713.46,2294.21,3.1104454627403844e-05 +gfx950,256,224,16384,1280,triton,0,0,12.382461,triton,758.75,2309.58,3.542218889508929e-05 +gfx950,256,240,16384,1280,triton,0,0,12.594796,triton,799.25,2313.9,3.4586588541666666e-05 +gfx950,256,256,16384,1280,triton,0,0,12.779949,triton,840.18,2323.0,3.838539123535156e-05 +gfx950,256,288,16384,1280,triton,0,0,17.848584,triton,676.78,1724.36,3.4544203016493055e-05 +gfx950,256,320,16384,1280,triton,0,0,18.574881,triton,722.58,1715.59,3.24249267578125e-05 +gfx950,256,352,16384,1280,triton,0,0,18.771178,triton,786.52,1755.69,4.109469327059659e-05 +gfx950,256,384,16384,1280,triton,0,0,18.599196,triton,865.96,1830.51,3.925959269205729e-05 +gfx950,256,416,16384,1280,triton,0,0,19.214888,triton,908.06,1828.56,3.9467444786658656e-05 +gfx950,256,448,16384,1280,triton,0,0,18.502488,triton,1015.57,1957.85,3.732953752790179e-05 +gfx950,256,480,16384,1280,triton,0,0,18.62917,triton,1080.71,2003.02,3.85284423828125e-05 +gfx950,256,512,16384,1280,triton,0,0,21.086155,triton,1018.43,1821.29,3.6716461181640625e-05 +gfx950,256,544,16384,1280,triton,0,0,29.2115,triton,781.1,1351.99,3.5454245174632356e-05 +gfx950,256,576,16384,1280,triton,0,0,31.16366,triton,775.24,1302.26,3.772311740451389e-05 +gfx950,256,608,16384,1280,triton,0,0,31.26487,triton,815.66,1332.89,3.644039756373355e-05 +gfx950,256,640,16384,1280,triton,0,0,31.966535,triton,839.74,1337.72,3.604888916015625e-05 +gfx950,256,672,16384,1280,triton,0,0,31.98406,triton,881.24,1371.05,3.760201590401786e-05 +gfx950,256,704,16384,1280,triton,0,0,32.73167,triton,902.12,1373.02,3.372539173473012e-05 +gfx950,256,736,16384,1280,triton,0,0,32.34807,triton,954.31,1422.99,4.212752632472826e-05 +gfx950,256,768,16384,1280,triton,0,0,32.37847,triton,994.87,1455.3,3.790855407714844e-05 +gfx950,256,800,16384,1280,triton,0,0,37.39611,triton,897.27,1289.17,3.5400390625e-05 +gfx950,256,832,16384,1280,triton,0,0,39.19572,triton,890.32,1257.78,3.653306227463942e-05 +gfx950,256,864,16384,1280,triton,0,0,38.52372,triton,940.69,1308.0,3.539191351996528e-05 +gfx950,256,896,16384,1280,triton,0,0,37.88573,triton,991.96,1358.78,3.7942613874162945e-05 +gfx950,256,928,16384,1280,triton,0,0,39.46853,triton,986.18,1331.9,3.466112860317888e-05 +gfx950,256,960,16384,1280,triton,0,0,39.76337,triton,1012.62,1349.42,3.7511189778645836e-05 +gfx950,256,992,16384,1280,triton,0,0,40.64054,triton,1023.79,1347.11,3.568587764616936e-05 +gfx950,256,1024,16384,1280,triton,0,0,40.78774,triton,1053.0,1368.96,3.7550926208496094e-05 +gfx950,256,1088,16384,1280,triton,0,0,48.0042,triton,950.63,1208.56,3.5117654239430144e-05 +gfx950,256,1152,16384,1280,triton,0,0,48.90261,triton,988.05,1230.91,3.7776099310980906e-05 +gfx950,256,1216,16384,1280,triton,0,0,51.02143,triton,999.63,1222.5,3.538633647717928e-05 +gfx950,256,1280,16384,1280,triton,0,0,50.95982,triton,1053.52,1266.74,3.6096572875976564e-05 +gfx950,256,1344,16384,1280,triton,0,0,55.01026,triton,1024.74,1213.08,3.742036365327381e-05 +gfx950,256,1408,16384,1280,triton,0,0,56.60747,triton,1043.25,1217.35,3.64563681862571e-05 +gfx950,256,1472,16384,1280,triton,0,0,56.99032,triton,1083.34,1247.41,3.424934718919837e-05 +gfx950,256,1536,16384,1280,triton,0,0,57.84669,triton,1113.71,1266.61,3.858407338460287e-05 +gfx950,256,1600,16384,1280,triton,0,0,64.59154,triton,1038.97,1168.08,3.7689208984375e-05 +gfx950,256,1664,16384,1280,triton,0,0,66.2888,triton,1052.87,1171.05,3.5616067739633414e-05 +gfx950,256,1728,16384,1280,triton,0,0,69.01198,triton,1050.22,1156.41,3.4932736997251154e-05 +gfx950,256,1792,16384,1280,triton,0,0,67.47961,triton,1113.85,1214.97,3.818103245326451e-05 +gfx950,256,1856,16384,1280,triton,0,0,73.664241,triton,1056.77,1142.54,3.4529587318157325e-05 +gfx950,256,1920,16384,1280,triton,0,0,74.28602,triton,1084.06,1162.31,3.445943196614583e-05 +gfx950,256,1984,16384,1280,triton,0,0,75.28003,triton,1105.41,1175.91,3.639344246156754e-05 +gfx950,256,2048,16384,1280,triton,0,0,76.50764,triton,1122.76,1185.53,3.6776065826416016e-05 +gfx950,256,4096,16384,1280,triton,0,0,145.22341,triton,1183.0,1104.73,3.6209821701049805e-05 +gfx950,256,6144,16384,1280,triton,0,0,216.90242,triton,1188.08,1061.13,3.673632939656576e-05 +gfx950,256,8192,16384,1280,triton,0,0,281.890337,triton,1218.9,1063.86,3.693997859954834e-05 +gfx950,256,10240,16384,1280,triton,0,0,334.018711,triton,1285.85,1106.59,3.706812858581543e-05 +gfx950,256,12288,16384,1280,triton,0,0,405.895859,triton,1269.77,1082.43,3.511210282643636e-05 +gfx950,256,14336,16384,1280,triton,0,0,460.187745,triton,1306.63,1106.25,3.7355082375662665e-05 +gfx950,256,16384,16384,1280,triton,0,0,515.854196,triton,1332.15,1122.05,3.6716461181640625e-05 +gfx950,256,32768,16384,1280,triton,0,0,1034.189168,triton,1328.95,1099.08,3.772415220737457e-05 +gfx950,256,1,32768,1280,triton,0,0,8.890484,triton,9.44,4725.26,9.1552734375e-05 +gfx950,256,2,32768,1280,triton,0,0,9.549418,triton,17.57,4406.2,7.62939453125e-05 +gfx950,256,4,32768,1280,triton,0,0,9.634681,triton,34.83,4381.08,5.340576171875e-05 +gfx950,256,8,32768,1280,triton,0,0,9.581011,triton,70.04,4433.52,4.57763671875e-05 +gfx950,256,16,32768,1280,triton,0,0,9.168,triton,146.4,4691.55,2.288818359375e-05 +gfx950,256,32,32768,1280,triton,0,0,10.086355,triton,266.14,4370.37,2.86102294921875e-05 +gfx950,256,48,32768,1280,triton,0,0,10.488229,triton,383.91,4304.85,3.814697265625e-05 +gfx950,256,64,32768,1280,triton,0,0,10.961849,triton,489.76,4216.37,3.24249267578125e-05 +gfx950,256,80,32768,1280,triton,0,0,12.835372,triton,522.84,3684.22,4.57763671875e-05 +gfx950,256,96,32768,1280,triton,0,0,13.401239,triton,600.92,3608.43,3.719329833984375e-05 +gfx950,256,112,32768,1280,triton,0,0,13.877678,triton,677.0,3561.58,3.378731863839286e-05 +gfx950,256,128,32768,1280,triton,0,0,14.322587,triton,749.68,3525.58,3.552436828613281e-05 +gfx950,256,144,32768,1280,triton,0,0,17.549958,triton,688.3,2938.16,4.280938042534722e-05 +gfx950,256,160,32768,1280,triton,0,0,18.506846,triton,725.23,2844.01,2.899169921875e-05 +gfx950,256,176,32768,1280,triton,0,0,18.133693,triton,814.17,2961.48,3.918734463778409e-05 +gfx950,256,192,32768,1280,triton,0,0,18.661495,triton,863.07,2935.01,4.0531158447265625e-05 +gfx950,256,208,32768,1280,triton,0,0,21.213302,triton,822.52,2632.35,3.7413377028245195e-05 +gfx950,256,224,32768,1280,triton,0,0,21.624167,triton,868.96,2631.77,3.569466727120536e-05 +gfx950,256,240,32768,1280,triton,0,0,21.943622,triton,917.47,2642.17,3.3696492513020836e-05 +gfx950,256,256,32768,1280,triton,0,0,21.855527,triton,982.58,2701.74,4.00543212890625e-05 +gfx950,256,288,32768,1280,triton,0,0,32.6633,triton,739.64,1873.24,3.8358900282118055e-05 +gfx950,256,320,32768,1280,triton,0,0,33.34799,triton,804.95,1898.89,3.747940063476562e-05 +gfx950,256,352,32768,1280,triton,0,0,35.06129,triton,842.18,1867.08,3.875385631214488e-05 +gfx950,256,384,32768,1280,triton,0,0,35.27252,triton,913.24,1916.52,3.6716461181640625e-05 +gfx950,256,416,32768,1280,triton,0,0,40.53016,triton,861.0,1720.66,3.7266657902644234e-05 +gfx950,256,448,32768,1280,triton,0,0,41.76055,triton,899.92,1721.16,3.7125178745814734e-05 +gfx950,256,480,32768,1280,triton,0,0,44.59817,triton,902.85,1659.59,3.579457600911458e-05 +gfx950,256,512,32768,1280,triton,0,0,42.48776,triton,1010.87,1792.35,3.618001937866211e-05 +gfx950,256,544,32768,1280,triton,0,0,51.46863,triton,886.64,1521.14,3.741769229664522e-05 +gfx950,256,576,32768,1280,triton,0,0,52.42624,triton,921.65,1534.14,3.7246280246310765e-05 +gfx950,256,608,32768,1280,triton,0,0,53.11989,triton,960.14,1554.36,3.794619911595395e-05 +gfx950,256,640,32768,1280,triton,0,0,53.00384,triton,1012.89,1598.1,3.690719604492188e-05 +gfx950,256,672,32768,1280,triton,0,0,59.43909,triton,948.39,1461.05,3.814697265625e-05 +gfx950,256,704,32768,1280,triton,0,0,62.40392,triton,946.35,1425.9,3.294511274857955e-05 +gfx950,256,736,32768,1280,triton,0,0,63.62852,triton,970.32,1432.06,3.603230351987092e-05 +gfx950,256,768,32768,1280,triton,0,0,61.22915,triton,1052.19,1523.09,3.6438306172688804e-05 +gfx950,256,800,32768,1280,triton,0,0,68.55037,triton,978.97,1391.62,3.810882568359375e-05 +gfx950,256,832,32768,1280,triton,0,0,70.60044,triton,988.57,1381.49,3.7780174842247594e-05 +gfx950,256,864,32768,1280,triton,0,0,69.99438,triton,1035.48,1424.0,3.761715359157986e-05 +gfx950,256,896,32768,1280,triton,0,0,72.80001,triton,1032.44,1398.49,3.6988939557756695e-05 +gfx950,256,928,32768,1280,triton,0,0,77.28044,triton,1007.32,1345.08,3.660136255724677e-05 +gfx950,256,960,32768,1280,triton,0,0,79.23366,triton,1016.37,1338.91,3.8623809814453125e-05 +gfx950,256,992,32768,1280,triton,0,0,80.91768,triton,1028.39,1337.46,3.6116569272933464e-05 +gfx950,256,1024,32768,1280,triton,0,0,80.656598,triton,1065.0,1368.3,3.644824028015137e-05 +gfx950,256,1088,32768,1280,triton,0,0,87.52134,triton,1042.81,1309.84,3.5622540642233455e-05 +gfx950,256,1152,32768,1280,triton,0,0,89.95935,triton,1074.23,1321.88,3.759066263834635e-05 +gfx950,256,1216,32768,1280,triton,0,0,95.96827,triton,1062.91,1283.67,3.669136448910362e-05 +gfx950,256,1280,32768,1280,triton,0,0,99.55423,triton,1078.55,1280.38,3.781318664550781e-05 +gfx950,256,1344,32768,1280,triton,0,0,105.06953,triton,1073.03,1253.87,3.6829993838355656e-05 +gfx950,256,1408,32768,1280,triton,0,0,112.494714,triton,1049.93,1209.12,3.554604270241477e-05 +gfx950,256,1472,32768,1280,triton,0,0,117.65679,triton,1049.5,1192.42,3.8084776505180024e-05 +gfx950,256,1536,32768,1280,triton,0,0,116.5072,triton,1105.93,1240.89,3.5881996154785156e-05 +gfx950,256,1600,32768,1280,triton,0,0,123.07403,triton,1090.54,1209.42,3.65447998046875e-05 +gfx950,256,1664,32768,1280,triton,0,0,129.218667,triton,1080.23,1185.01,3.60562251164363e-05 +gfx950,256,1728,32768,1280,triton,0,0,131.0321,triton,1106.26,1201.24,3.7104995162398724e-05 +gfx950,256,1792,32768,1280,triton,0,0,134.82173,triton,1114.98,1199.19,3.593308585030692e-05 +gfx950,256,1856,32768,1280,triton,0,0,142.09419,triton,1095.7,1167.91,3.873890843884698e-05 +gfx950,256,1920,32768,1280,triton,0,0,149.62225,triton,1076.45,1137.73,3.5715103149414065e-05 +gfx950,256,1984,32768,1280,triton,0,0,155.29118,triton,1071.73,1123.73,3.417845695249496e-05 +gfx950,256,2048,32768,1280,triton,0,0,151.75707,triton,1132.06,1178.08,3.750622272491455e-05 +gfx950,256,4096,32768,1280,triton,0,0,276.260918,triton,1243.74,1142.48,3.6403536796569824e-05 +gfx950,256,6144,32768,1280,triton,0,0,396.13111,triton,1301.07,1142.2,3.680586814880371e-05 +gfx950,256,8192,32768,1280,triton,0,0,534.939922,triton,1284.62,1101.62,3.76775860786438e-05 +gfx950,256,10240,32768,1280,triton,0,0,642.302766,triton,1337.37,1130.52,3.6910176277160645e-05 +gfx950,256,12288,32768,1280,triton,0,0,740.03204,triton,1392.9,1166.14,3.759562969207764e-05 +gfx950,256,14336,32768,1280,triton,0,0,924.907011,triton,1300.23,1080.99,3.775741372789655e-05 +gfx950,256,16384,32768,1280,triton,0,0,1084.219571,triton,1267.63,1048.36,3.637373447418213e-05 +gfx950,256,32768,32768,1280,triton,0,0,2166.11598,triton,1268.99,1030.12,3.69194895029068e-05 +gfx950,256,1,5120,2048,triton,0,0,4.721645,triton,4.44,2223.39,0.0 +gfx950,256,2,5120,2048,triton,0,0,4.806481,triton,8.73,2186.7,9.765625e-05 +gfx950,256,4,5120,2048,triton,0,0,4.806312,triton,17.45,2191.89,0.000146484375 +gfx950,256,8,5120,2048,triton,0,0,5.350725,triton,31.36,1978.06,0.0 +gfx950,256,16,5120,2048,triton,0,0,5.28936,triton,63.44,2019.6,7.32421875e-05 +gfx950,256,32,5120,2048,triton,0,0,6.260505,triton,107.19,1737.72,4.8828125e-05 +gfx950,256,48,5120,2048,triton,0,0,7.28789,triton,138.12,1519.72,8.138020833333333e-05 +gfx950,256,64,5120,2048,triton,0,0,7.369137,triton,182.13,1529.65,5.79833984375e-05 +gfx950,256,80,5120,2048,triton,0,0,7.287607,triton,230.22,1573.74,5.615234375e-05 +gfx950,256,96,5120,2048,triton,0,0,7.410684,triton,271.67,1574.13,4.679361979166667e-05 +gfx950,256,112,5120,2048,triton,0,0,8.183784,triton,287.01,1449.45,5.7547433035714285e-05 +gfx950,256,128,5120,2048,triton,0,0,8.238663,triton,325.82,1463.66,6.103515625e-05 +gfx950,256,144,5120,2048,triton,0,0,8.495378,triton,355.48,1442.58,6.239149305555555e-05 +gfx950,256,160,5120,2048,triton,0,0,8.988105,triton,373.32,1385.37,5.9814453125e-05 +gfx950,256,176,5120,2048,triton,0,0,9.136667,triton,403.98,1384.36,4.549893465909091e-05 +gfx950,256,192,5120,2048,triton,0,0,9.088768,triton,443.02,1413.29,6.001790364583333e-05 +gfx950,256,208,5120,2048,triton,0,0,11.219554,triton,388.79,1162.4,5.915715144230769e-05 +gfx950,256,224,5120,2048,triton,0,0,11.437011,triton,410.74,1157.49,5.057198660714286e-05 +gfx950,256,240,5120,2048,triton,0,0,11.567823,triton,435.1,1161.4,6.429036458333333e-05 +gfx950,256,256,5120,2048,triton,0,0,11.758208,triton,456.59,1159.32,5.8746337890625e-05 +gfx950,256,288,5120,2048,triton,0,0,11.749109,triton,514.06,1193.68,6.442599826388889e-05 +gfx950,256,320,5120,2048,triton,0,0,12.345875,triton,543.57,1167.83,6.16455078125e-05 +gfx950,256,352,5120,2048,triton,0,0,12.571795,triton,587.19,1178.12,6.491921164772728e-05 +gfx950,256,384,5120,2048,triton,0,0,12.700299,triton,634.08,1197.16,5.900065104166667e-05 +gfx950,256,416,5120,2048,triton,0,0,14.38808,triton,606.35,1084.06,5.258413461538462e-05 +gfx950,256,448,5120,2048,triton,0,0,14.544421,triton,645.97,1099.44,5.449567522321429e-05 +gfx950,256,480,5120,2048,triton,0,0,14.749708,triton,682.48,1110.8,5.126953125e-05 +gfx950,256,512,5120,2048,triton,0,0,15.027638,triton,714.51,1116.42,5.68389892578125e-05 +gfx950,256,544,5120,2048,triton,0,0,14.760899,triton,772.89,1163.24,5.4931640625e-05 +gfx950,256,576,5120,2048,triton,0,0,15.43957,triton,782.38,1137.57,5.3575303819444444e-05 +gfx950,256,608,5120,2048,triton,0,0,15.548193,triton,820.07,1154.92,5.6216591282894735e-05 +gfx950,256,640,5120,2048,triton,0,0,15.534114,triton,864.02,1181.28,5.157470703125e-05 +gfx950,256,672,5120,2048,triton,0,0,15.169989,triton,929.0,1235.55,5.841936383928572e-05 +gfx950,256,704,5120,2048,triton,0,0,15.32386,triton,963.46,1248.8,6.658380681818181e-05 +gfx950,256,736,5120,2048,triton,0,0,15.558711,triton,992.05,1255.23,5.705460258152174e-05 +gfx950,256,768,5120,2048,triton,0,0,15.714582,triton,1024.92,1267.8,5.620320638020833e-05 +gfx950,256,800,5120,2048,triton,0,0,23.314594,triton,719.6,871.39,6.6650390625e-05 +gfx950,256,832,5120,2048,triton,0,0,23.910407,triton,729.74,866.12,6.103515625e-05 +gfx950,256,864,5120,2048,triton,0,0,24.12714,triton,751.0,874.64,6.510416666666667e-05 +gfx950,256,896,5120,2048,triton,0,0,24.301235,triton,773.23,884.56,6.692068917410715e-05 +gfx950,256,928,5120,2048,triton,0,0,24.470323,triton,795.31,894.51,5.724676724137931e-05 +gfx950,256,960,5120,2048,triton,0,0,24.521826,triton,821.01,908.67,5.391438802083333e-05 +gfx950,256,992,5120,2048,triton,0,0,24.490479,triton,849.46,925.89,5.414409022177419e-05 +gfx950,256,1024,5120,2048,triton,0,0,25.085724,triton,856.06,919.59,6.44683837890625e-05 +gfx950,256,1088,5120,2048,triton,0,0,25.156327,triton,907.01,948.27,5.7265337775735294e-05 +gfx950,256,1152,5120,2048,triton,0,0,26.087347,triton,926.09,944.58,5.8492024739583336e-05 +gfx950,256,1216,5120,2048,triton,0,0,27.33243,triton,933.01,930.32,5.573473478618421e-05 +gfx950,256,1280,5120,2048,triton,0,0,26.574485,triton,1010.12,986.45,5.7525634765625e-05 +gfx950,256,1344,5120,2048,triton,0,0,27.37684,triton,1029.55,986.27,5.769275483630952e-05 +gfx950,256,1408,5120,2048,triton,0,0,27.54125,triton,1072.13,1008.93,5.770596590909091e-05 +gfx950,256,1472,5120,2048,triton,0,0,27.40684,triton,1126.36,1042.58,6.090247112771739e-05 +gfx950,256,1536,5120,2048,triton,0,0,28.58405,triton,1126.93,1027.15,6.052652994791667e-05 +gfx950,256,1600,5120,2048,triton,0,0,37.406182,triton,897.03,805.92,5.92041015625e-05 +gfx950,256,1664,5120,2048,triton,0,0,40.17334,triton,868.65,769.99,5.939190204326923e-05 +gfx950,256,1728,5120,2048,triton,0,0,42.08895,triton,861.0,753.63,5.809642650462963e-05 +gfx950,256,1792,5120,2048,triton,0,0,40.76334,triton,921.93,797.43,6.125313895089286e-05 +gfx950,256,1856,5120,2048,triton,0,0,41.85215,triton,930.02,795.47,6.471831223060345e-05 +gfx950,256,1920,5120,2048,triton,0,0,43.49857,triton,925.67,783.44,6.001790364583333e-05 +gfx950,256,1984,5120,2048,triton,0,0,43.34656,triton,959.88,804.33,5.8770948840725806e-05 +gfx950,256,2048,5120,2048,triton,0,0,44.59618,triton,963.08,799.43,5.588531494140625e-05 +gfx950,256,4096,5120,2048,triton,0,0,72.39921,triton,1186.47,840.03,6.012916564941406e-05 +gfx950,256,6144,5120,2048,triton,0,0,105.82749,triton,1217.54,812.48,5.613962809244792e-05 +gfx950,256,8192,5120,2048,triton,0,0,136.82399,triton,1255.62,812.35,5.9533119201660154e-05 +gfx950,256,10240,5120,2048,triton,0,0,164.21318,triton,1307.74,830.11,5.870819091796875e-05 +gfx950,256,12288,5120,2048,triton,0,0,189.15469,triton,1362.37,853.7,5.656878153483073e-05 +gfx950,256,14336,5120,2048,triton,0,0,220.767866,triton,1361.83,845.44,5.896432059151786e-05 +gfx950,256,16384,5120,2048,triton,0,0,241.068776,triton,1425.31,878.64,5.778074264526367e-05 +gfx950,256,32768,5120,2048,triton,0,0,452.761969,triton,1517.78,912.49,5.9604644775390625e-05 +gfx950,256,1,5120,4096,triton,0,0,6.900466,triton,6.08,3041.22,0.000390625 +gfx950,256,2,5120,4096,triton,0,0,7.477207,triton,11.22,2808.56,9.765625e-05 +gfx950,256,4,5120,4096,triton,0,0,7.499708,triton,22.37,2803.96,0.000146484375 +gfx950,256,8,5120,4096,triton,0,0,7.910079,triton,42.42,2665.74,9.765625e-05 +gfx950,256,16,5120,4096,triton,0,0,7.339966,triton,91.43,2888.42,0.00018310546875 +gfx950,256,32,5120,4096,triton,0,0,7.994584,triton,167.89,2680.6,0.000146484375 +gfx950,256,48,5120,4096,triton,0,0,9.321156,triton,215.99,2323.71,8.951822916666666e-05 +gfx950,256,64,5120,4096,triton,0,0,9.676126,triton,277.42,2262.17,0.0001007080078125 +gfx950,256,80,5120,4096,triton,0,0,12.008459,triton,279.42,1841.9,0.00010986328125 +gfx950,256,96,5120,4096,triton,0,0,13.041524,triton,308.75,1713.59,0.00012003580729166666 +gfx950,256,112,5120,4096,triton,0,0,14.015837,triton,335.17,1610.83,0.0001220703125 +gfx950,256,128,5120,4096,triton,0,0,14.265652,triton,376.34,1598.7,0.00010223388671875 +gfx950,256,144,5120,4096,triton,0,0,14.252437,triton,423.77,1616.28,0.00012071397569444445 +gfx950,256,160,5120,4096,triton,0,0,14.414865,triton,465.55,1613.98,0.000128173828125 +gfx950,256,176,5120,4096,triton,0,0,14.334545,triton,514.98,1639.02,0.00011208274147727273 +gfx950,256,192,5120,4096,triton,0,0,14.474575,triton,556.36,1639.01,0.00011698404947916667 +gfx950,256,208,5120,4096,triton,0,0,19.049182,triton,457.98,1257.45,0.00011737530048076923 +gfx950,256,224,5120,4096,triton,0,0,19.722447,triton,476.37,1226.16,0.00012381417410714285 +gfx950,256,240,5120,4096,triton,0,0,20.226593,triton,497.68,1206.93,0.00011800130208333333 +gfx950,256,256,5120,4096,triton,0,0,20.286787,triton,529.28,1214.66,0.0001129150390625 +gfx950,256,288,5120,4096,triton,0,0,20.102237,triton,600.91,1248.63,0.00010918511284722223 +gfx950,256,320,5120,4096,triton,0,0,19.827368,triton,676.93,1289.08,0.000118408203125 +gfx950,256,352,5120,4096,triton,0,0,21.176884,triton,697.17,1228.59,0.00011818625710227272 +gfx950,256,384,5120,4096,triton,0,0,21.69875,triton,742.26,1220.19,0.00011749267578125 +gfx950,256,416,5120,4096,triton,0,0,25.243073,triton,691.21,1067.04,0.00013380784254807694 +gfx950,256,448,5120,4096,triton,0,0,25.891719,triton,725.73,1058.02,0.00011814662388392857 +gfx950,256,480,5120,4096,triton,0,0,26.139908,triton,770.19,1065.53,0.00011881510416666667 +gfx950,256,512,5120,4096,triton,0,0,26.276765,triton,817.26,1077.44,0.0001247406005859375 +gfx950,256,544,5120,4096,triton,0,0,27.25484,triton,837.17,1055.6,0.00011740291819852941 +gfx950,256,576,5120,4096,triton,0,0,27.44964,triton,880.13,1064.82,0.00011020236545138889 +gfx950,256,608,5120,4096,triton,0,0,26.890929,triton,948.33,1104.01,0.00012303402549342104 +gfx950,256,640,5120,4096,triton,0,0,27.88922,triton,962.51,1080.94,0.00012115478515625 +gfx950,256,672,5120,4096,triton,0,0,27.88484,triton,1010.79,1097.56,0.00013108026413690476 +gfx950,256,704,5120,4096,triton,0,0,28.54365,triton,1034.48,1088.3,0.00012123801491477273 +gfx950,256,736,5120,4096,triton,0,0,29.379901,triton,1050.72,1072.94,0.00011941661005434783 +gfx950,256,768,5120,4096,triton,0,0,29.558753,triton,1089.77,1081.97,0.000115203857421875 +gfx950,256,800,5120,4096,triton,0,0,41.598011,triton,806.64,779.85,0.000113525390625 +gfx950,256,832,5120,4096,triton,0,0,42.42602,triton,822.53,775.45,0.0001190185546875 +gfx950,256,864,5120,4096,triton,0,0,42.86376,triton,845.44,778.23,0.00013359917534722224 +gfx950,256,896,5120,4096,triton,0,0,43.356052,triton,866.8,779.97,0.00011444091796875 +gfx950,256,928,5120,4096,triton,0,0,44.19697,triton,880.67,775.51,0.00011680866109913793 +gfx950,256,960,5120,4096,triton,0,0,45.39098,triton,887.08,765.22,0.00011454264322916667 +gfx950,256,992,5120,4096,triton,0,0,45.82578,triton,907.95,767.97,0.00010651619203629032 +gfx950,256,1024,5120,4096,triton,0,0,45.24018,triton,949.37,788.05,0.00012912750244140626 +gfx950,256,1088,5120,4096,triton,0,0,47.01939,triton,970.54,777.74,0.0001235064338235294 +gfx950,256,1152,5120,4096,triton,0,0,47.7166,triton,1012.61,785.61,0.00010867648654513888 +gfx950,256,1216,5120,4096,triton,0,0,48.25421,triton,1056.96,795.87,0.00011805484169407894 +gfx950,256,1280,5120,4096,triton,0,0,49.28301,triton,1089.36,797.87,0.0001171875 +gfx950,256,1344,5120,4096,triton,0,0,49.39181,triton,1141.31,814.69,0.00012483142671130951 +gfx950,256,1408,5120,4096,triton,0,0,52.49104,triton,1125.06,784.07,0.00012109929865056818 +gfx950,256,1472,5120,4096,triton,0,0,53.63905,triton,1151.03,784.39,0.00011556874150815218 +gfx950,256,1536,5120,4096,triton,0,0,52.41824,triton,1229.05,820.17,0.00011990865071614583 +gfx950,256,1600,5120,4096,triton,0,0,69.43758,triton,966.46,632.35,0.00012451171875 +gfx950,256,1664,5120,4096,triton,0,0,70.8684,triton,984.83,632.53,0.00011491041917067308 +gfx950,256,1728,5120,4096,triton,0,0,72.98102,triton,993.1,626.79,0.0001182273582175926 +gfx950,256,1792,5120,4096,triton,0,0,78.12366,triton,962.09,597.28,0.00011629377092633928 +gfx950,256,1856,5120,4096,triton,0,0,76.23924,triton,1021.08,624.08,0.00011817669046336206 +gfx950,256,1920,5120,4096,triton,0,0,80.40928,triton,1001.51,603.12,0.00011952718098958333 +gfx950,256,1984,5120,4096,triton,0,0,81.52808,triton,1020.69,606.1,0.00012128276209677419 +gfx950,256,2048,5120,4096,triton,0,0,79.16527,triton,1085.06,635.78,0.00011930465698242187 +gfx950,256,4096,5120,4096,triton,0,0,130.4073,triton,1317.4,611.1,0.00011982917785644532 +gfx950,256,6144,5120,4096,triton,0,0,182.60417,triton,1411.24,597.2,0.00011695226033528646 +gfx950,256,8192,5120,4096,triton,0,0,231.164316,triton,1486.38,598.76,0.00012171268463134766 +gfx950,256,10240,5120,4096,triton,0,0,291.341303,triton,1474.2,575.86,0.00011877059936523438 +gfx950,256,12288,5120,4096,triton,0,0,337.850788,triton,1525.51,583.49,0.0001128991444905599 +gfx950,256,14336,5120,4096,triton,0,0,404.222674,triton,1487.54,560.32,0.00011516298566545759 +gfx950,256,16384,5120,4096,triton,0,0,455.817396,triton,1507.61,561.3,0.00011864900588989257 +gfx950,256,32768,5120,4096,triton,0,0,854.649072,triton,1608.13,574.19,0.00011526942253112794 +gfx950,256,1,5120,8192,triton,0,0,9.874798,triton,8.49,4249.35,0.0 +gfx950,256,2,5120,8192,triton,0,0,9.866151,triton,17.0,4254.94,0.0001953125 +gfx950,256,4,5120,8192,triton,0,0,10.192691,triton,32.92,4122.24,0.000341796875 +gfx950,256,8,5120,8192,triton,0,0,10.243944,triton,65.51,4108.82,0.000146484375 +gfx950,256,16,5120,8192,triton,0,0,10.946352,triton,122.61,3858.63,0.00025634765625 +gfx950,256,32,5120,8192,triton,0,0,12.510892,triton,214.56,3399.67,0.0002197265625 +gfx950,256,48,5120,8192,triton,0,0,14.005449,triton,287.5,3057.94,0.00023193359375 +gfx950,256,64,5120,8192,triton,0,0,14.511283,triton,369.97,2971.67,0.000177001953125 +gfx950,256,80,5120,8192,triton,0,0,19.902326,triton,337.19,2181.53,0.00022705078125 +gfx950,256,96,5120,8192,triton,0,0,19.852511,triton,405.64,2201.86,0.0002197265625 +gfx950,256,112,5120,8192,triton,0,0,20.367355,triton,461.29,2160.68,0.0002232142857142857 +gfx950,256,128,5120,8192,triton,0,0,20.448106,triton,525.11,2166.57,0.0002197265625 +gfx950,256,144,5120,8192,triton,0,0,24.787101,triton,487.33,1799.21,0.00026041666666666666 +gfx950,256,160,5120,8192,triton,0,0,26.738616,triton,501.96,1678.93,0.000230712890625 +gfx950,256,176,5120,8192,triton,0,0,25.291,triton,583.76,1786.69,0.00021528764204545455 +gfx950,256,192,5120,8192,triton,0,0,24.758372,triton,650.53,1837.03,0.00022481282552083332 +gfx950,256,208,5120,8192,triton,0,0,32.412054,triton,538.33,1412.34,0.00021690955528846153 +gfx950,256,224,5120,8192,triton,0,0,33.338733,triton,563.62,1381.93,0.0002197265625 +gfx950,256,240,5120,8192,triton,0,0,33.891912,triton,594.03,1368.08,0.00019612630208333334 +gfx950,256,256,5120,8192,triton,0,0,34.692396,triton,619.01,1345.01,0.000202178955078125 +gfx950,256,288,5120,8192,triton,0,0,32.627202,triton,740.46,1448.22,0.00021226671006944445 +gfx950,256,320,5120,8192,triton,0,0,33.81783,triton,793.77,1414.68,0.000213623046875 +gfx950,256,352,5120,8192,triton,0,0,35.455406,triton,832.82,1365.97,0.00020252574573863637 +gfx950,256,384,5120,8192,triton,0,0,34.892217,triton,923.19,1404.92,0.00022074381510416667 +gfx950,256,416,5120,8192,triton,0,0,47.521543,triton,734.33,1043.96,0.00020329402043269232 +gfx950,256,448,5120,8192,triton,0,0,48.750053,triton,770.89,1029.75,0.00021275111607142857 +gfx950,256,480,5120,8192,triton,0,0,50.126763,triton,803.27,1013.24,0.000203857421875 +gfx950,256,512,5120,8192,triton,0,0,50.227592,triton,855.1,1022.95,0.0002109527587890625 +gfx950,256,544,5120,8192,triton,0,0,50.07343,triton,911.34,1037.88,0.0002251120174632353 +gfx950,256,576,5120,8192,triton,0,0,52.53824,triton,919.68,1000.41,0.00019700792100694445 +gfx950,256,608,5120,8192,triton,0,0,53.631619,triton,950.98,991.01,0.00021715666118421052 +gfx950,256,640,5120,8192,triton,0,0,51.360702,triton,1045.3,1046.32,0.0002154541015625 +gfx950,256,672,5120,8192,triton,0,0,53.590264,triton,1051.9,1013.79,0.00021740141369047618 +gfx950,256,704,5120,8192,triton,0,0,54.07548,triton,1092.1,1015.6,0.00019725452769886364 +gfx950,256,736,5120,8192,triton,0,0,54.76266,triton,1127.41,1013.63,0.000213623046875 +gfx950,256,768,5120,8192,triton,0,0,54.227138,triton,1188.05,1034.52,0.00021158854166666667 +gfx950,256,800,5120,8192,triton,0,0,77.267979,triton,868.52,733.66,0.00020654296875 +gfx950,256,832,5120,8192,triton,0,0,80.56607,triton,866.29,710.95,0.00020845853365384615 +gfx950,256,864,5120,8192,triton,0,0,80.59008,triton,899.34,718.06,0.00021701388888888888 +gfx950,256,896,5120,8192,triton,0,0,78.802959,triton,953.8,741.83,0.00021231515066964286 +gfx950,256,928,5120,8192,triton,0,0,81.10048,triton,959.87,728.08,0.00022056842672413794 +gfx950,256,960,5120,8192,triton,0,0,82.85289,triton,971.97,719.8,0.0002099609375 +gfx950,256,992,5120,8192,triton,0,0,81.40371,triton,1022.25,739.86,0.00022543630292338709 +gfx950,256,1024,5120,8192,triton,0,0,82.47929,triton,1041.47,737.37,0.00021686553955078126 +gfx950,256,1088,5120,8192,triton,0,0,84.57011,triton,1079.2,733.08,0.00021164838005514707 +gfx950,256,1152,5120,8192,triton,0,0,88.63854,triton,1090.23,712.75,0.00021718343098958332 +gfx950,256,1216,5120,8192,triton,0,0,92.08137,triton,1107.78,698.91,0.0002094469572368421 +gfx950,256,1280,5120,8192,triton,0,0,89.87696,triton,1194.68,729.17,0.000214996337890625 +gfx950,256,1344,5120,8192,triton,0,0,91.309761,triton,1234.73,730.65,0.00020315987723214287 +gfx950,256,1408,5120,8192,triton,0,0,93.85023,triton,1258.51,723.44,0.00021722966974431818 +gfx950,256,1472,5120,8192,triton,0,0,93.88299,triton,1315.26,735.76,0.00021043860394021738 +gfx950,256,1536,5120,8192,triton,0,0,95.6106,triton,1347.64,734.8,0.0002124786376953125 +gfx950,256,1600,5120,8192,triton,0,0,127.939443,triton,1049.07,558.34,0.0002156982421875 +gfx950,256,1664,5120,8192,triton,0,0,133.24818,triton,1047.57,544.95,0.0002224261944110577 +gfx950,256,1728,5120,8192,triton,0,0,139.22401,triton,1041.16,530.03,0.00022017867476851853 +gfx950,256,1792,5120,8192,triton,0,0,139.84099,triton,1074.96,536.13,0.00021634783063616072 +gfx950,256,1856,5120,8192,triton,0,0,139.399357,triton,1116.88,546.29,0.00021625387257543104 +gfx950,256,1920,5120,8192,triton,0,0,146.799474,triton,1097.15,526.79,0.00020751953125 +gfx950,256,1984,5120,8192,triton,0,0,145.765919,triton,1141.76,538.62,0.00021411526587701614 +gfx950,256,2048,5120,8192,triton,0,0,148.685909,triton,1155.45,535.97,0.0002185821533203125 +gfx950,256,4096,5120,8192,triton,0,0,238.24599,triton,1442.2,492.94,0.00021147727966308594 +gfx950,256,6144,5120,8192,triton,0,0,344.614885,triton,1495.57,450.33,0.0002136866251627604 +gfx950,256,8192,5120,8192,triton,0,0,443.211948,triton,1550.49,435.32,0.0002140522003173828 +gfx950,256,10240,5120,8192,triton,0,0,535.048083,triton,1605.45,431.15,0.00021739959716796876 +gfx950,256,12288,5120,8192,triton,0,0,622.368086,triton,1656.24,431.31,0.00021486282348632814 +gfx950,256,14336,5120,8192,triton,0,0,739.354643,triton,1626.54,414.12,0.00020900453839983258 +gfx950,256,16384,5120,8192,triton,0,0,833.843163,triton,1648.26,412.47,0.00021250247955322265 +gfx950,256,32768,5120,8192,triton,0,0,1651.778919,triton,1664.13,391.05,0.0002121567726135254 +gfx950,256,1,1152,5120,triton,0,0,4.93701,triton,2.39,1196.2,0.0 +gfx950,256,2,1152,5120,triton,0,0,5.197818,triton,4.54,1137.61,0.0008680555555555555 +gfx950,256,4,1152,5120,triton,0,0,5.34347,triton,8.83,1109.38,0.0 +gfx950,256,8,1152,5120,triton,0,0,5.58624,triton,16.89,1066.48,0.00010850694444444444 +gfx950,256,16,1152,5120,triton,0,0,5.36942,triton,35.15,1120.61,0.00016276041666666666 +gfx950,256,32,1152,5120,triton,0,0,5.338387,triton,70.71,1149.37,0.00013563368055555556 +gfx950,256,48,1152,5120,triton,0,0,6.54226,triton,86.55,956.03,0.0003074363425925926 +gfx950,256,64,1152,5120,triton,0,0,6.155326,triton,122.65,1035.42,0.00018988715277777778 +gfx950,256,80,1152,5120,triton,0,0,7.431296,triton,126.99,873.62,0.00017361111111111112 +gfx950,256,96,1152,5120,triton,0,0,7.566062,triton,149.68,873.76,0.00010850694444444444 +gfx950,256,112,1152,5120,triton,0,0,8.539263,triton,154.72,788.09,8.525545634920634e-05 +gfx950,256,128,1152,5120,triton,0,0,8.507374,triton,177.49,805.01,0.00014919704861111112 +gfx950,256,144,1152,5120,triton,0,0,9.02905,triton,188.14,771.65,0.0002290702160493827 +gfx950,256,160,1152,5120,triton,0,0,9.12387,triton,206.87,776.65,0.00013020833333333333 +gfx950,256,176,1152,5120,triton,0,0,9.387546,triton,221.16,767.49,0.0001282354797979798 +gfx950,256,192,1152,5120,triton,0,0,9.36266,triton,241.91,782.22,0.00015823929398148148 +gfx950,256,208,1152,5120,triton,0,0,9.65946,triton,254.02,770.48,0.00015858707264957266 +gfx950,256,224,1152,5120,triton,0,0,9.63466,triton,274.26,784.79,0.00015888516865079366 +gfx950,256,240,1152,5120,triton,0,0,10.918701,triton,259.29,703.38,0.00013020833333333333 +gfx950,256,256,1152,5120,triton,0,0,10.97429,triton,275.18,710.64,0.00016615125868055556 +gfx950,256,288,1152,5120,triton,0,0,12.17951,triton,278.94,659.83,0.00014166184413580247 +gfx950,256,320,1152,5120,triton,0,0,12.608808,triton,299.38,656.2,0.00016004774305555556 +gfx950,256,352,1152,5120,triton,0,0,12.47433,triton,332.87,682.32,0.00013809974747474747 +gfx950,256,384,1152,5120,triton,0,0,12.766354,triton,354.83,685.32,0.0001243308738425926 +gfx950,256,416,1152,5120,triton,0,0,12.814612,triton,382.95,701.28,0.00015024038461538462 +gfx950,256,448,1152,5120,triton,0,0,13.047596,triton,405.04,706.96,0.00013757130456349206 +gfx950,256,480,1152,5120,triton,0,0,15.543031,triton,364.3,608.75,0.000146484375 +gfx950,256,512,1152,5120,triton,0,0,15.691253,triton,384.91,618.14,0.00015428331163194444 +gfx950,256,544,1152,5120,triton,0,0,16.041592,triton,400.04,619.45,0.00013403799019607844 +gfx950,256,576,1152,5120,triton,0,0,16.109337,triton,421.79,631.59,0.00015823929398148148 +gfx950,256,608,1152,5120,triton,0,0,17.529242,triton,409.16,593.98,0.00014848318713450292 +gfx950,256,640,1152,5120,triton,0,0,18.33437,triton,411.78,580.85,0.00013292100694444443 +gfx950,256,672,1152,5120,triton,0,0,17.626082,triton,449.74,617.67,0.00016146866732804233 +gfx950,256,704,1152,5120,triton,0,0,18.338505,triton,452.86,606.63,0.00014919704861111112 +gfx950,256,736,1152,5120,triton,0,0,17.819778,triton,487.22,637.62,0.00014271022041062803 +gfx950,256,768,1152,5120,triton,0,0,17.937,triton,505.08,646.7,0.00013450339988425926 +gfx950,256,800,1152,5120,triton,0,0,17.542532,triton,537.96,674.79,0.00013671875 +gfx950,256,832,1152,5120,triton,0,0,18.521071,triton,529.92,651.96,0.00012937366452991453 +gfx950,256,864,1152,5120,triton,0,0,17.834026,triton,571.5,690.4,0.00013563368055555556 +gfx950,256,896,1152,5120,triton,0,0,18.128395,triton,583.04,692.29,0.0001433841765873016 +gfx950,256,928,1152,5120,triton,0,0,23.820626,triton,459.57,536.83,0.00013563368055555556 +gfx950,256,960,1152,5120,triton,0,0,24.629495,triton,459.8,528.85,0.00011845341435185185 +gfx950,256,992,1152,5120,triton,0,0,25.385073,triton,460.98,522.47,0.0001548849126344086 +gfx950,256,1024,1152,5120,triton,0,0,25.860479,triton,467.11,522.05,0.00017632378472222222 +gfx950,256,1088,1152,5120,triton,0,0,26.014851,triton,493.36,537.21,0.00014680351307189542 +gfx950,256,1152,1152,5120,triton,0,0,25.700731,triton,528.76,562.27,0.00013563368055555556 +gfx950,256,1216,1152,5120,triton,0,0,26.967988,triton,531.91,553.46,0.00015704952485380116 +gfx950,256,1280,1152,5120,triton,0,0,26.195842,triton,576.41,587.92,0.00014580620659722222 +gfx950,256,1344,1152,5120,triton,0,0,28.524598,triton,555.82,556.58,0.00015500992063492063 +gfx950,256,1408,1152,5120,triton,0,0,27.98003,triton,593.62,584.39,0.00013809974747474747 +gfx950,256,1472,1152,5120,triton,0,0,27.451526,triton,632.55,612.95,0.00016040157004830919 +gfx950,256,1536,1152,5120,triton,0,0,27.96721,triton,647.88,618.64,0.0001384593822337963 +gfx950,256,1600,1152,5120,triton,0,0,28.715874,triton,657.28,619.05,0.0001437717013888889 +gfx950,256,1664,1152,5120,triton,0,0,28.855571,triton,680.26,632.52,0.0001351120125534188 +gfx950,256,1728,1152,5120,triton,0,0,28.255456,triton,721.43,662.77,0.00015622990612139918 +gfx950,256,1792,1152,5120,triton,0,0,29.333702,triton,720.65,654.61,0.00014580620659722222 +gfx950,256,1856,1152,5120,triton,0,0,34.849058,triton,628.26,564.64,0.0001594865002394636 +gfx950,256,1920,1152,5120,triton,0,0,37.486365,triton,604.2,537.59,0.00014241536458333334 +gfx950,256,1984,1152,5120,triton,0,0,33.495558,triton,698.73,615.83,0.00014350918458781362 +gfx950,256,2048,1152,5120,triton,0,0,35.088185,triton,688.53,601.42,0.0001373291015625 +gfx950,256,4096,1152,5120,triton,0,0,54.865173,triton,880.67,661.75,0.00013139512803819444 +gfx950,256,6144,1152,5120,triton,0,0,64.01514,triton,1132.19,804.67,0.0001552723072193287 +gfx950,256,8192,1152,5120,triton,0,0,98.728978,triton,978.81,675.75,0.00014294518364800347 +gfx950,256,10240,1152,5120,triton,0,0,102.78226,triton,1175.26,797.02,0.00014046563042534722 +gfx950,256,12288,1152,5120,triton,0,0,118.47399,triton,1223.52,819.79,0.00014538235134548612 +gfx950,256,14336,1152,5120,triton,0,0,133.04096,triton,1271.14,844.32,0.00014156765407986112 +gfx950,256,16384,1152,5120,triton,0,0,163.963758,triton,1178.76,777.81,0.0001404020521375868 +gfx950,256,32768,1152,5120,triton,0,0,297.09949,triton,1301.07,838.67,0.00014930301242404513 +gfx950,256,1,2304,5120,triton,0,0,5.722451,triton,4.12,2063.14,0.0 +gfx950,256,2,2304,5120,triton,0,0,5.690383,triton,8.29,2076.47,0.0 +gfx950,256,4,2304,5120,triton,0,0,5.701765,triton,16.55,2075.74,0.00010850694444444444 +gfx950,256,8,2304,5120,triton,0,0,5.985,triton,31.54,1984.01,0.00010850694444444444 +gfx950,256,16,2304,5120,triton,0,0,6.314685,triton,59.78,1892.75,0.0003255208333333333 +gfx950,256,32,2304,5120,triton,0,0,6.907643,triton,109.3,1752.81,8.138020833333333e-05 +gfx950,256,48,2304,5120,triton,0,0,8.655082,triton,130.84,1416.9,0.0001537181712962963 +gfx950,256,64,2304,5120,triton,0,0,8.569404,triton,176.2,1449.23,0.00020345052083333334 +gfx950,256,80,2304,5120,triton,0,0,9.240608,triton,204.25,1360.81,8.680555555555556e-05 +gfx950,256,96,2304,5120,triton,0,0,9.433747,triton,240.09,1349.45,9.946469907407407e-05 +gfx950,256,112,2304,5120,triton,0,0,10.75516,triton,245.69,1198.12,0.00017826140873015872 +gfx950,256,128,2304,5120,triton,0,0,10.856823,triton,278.16,1201.24,0.00016276041666666666 +gfx950,256,144,2304,5120,triton,0,0,11.616075,triton,292.47,1136.12,0.00017783082561728395 +gfx950,256,160,2304,5120,triton,0,0,11.864895,triton,318.15,1125.42,0.00013834635416666666 +gfx950,256,176,2304,5120,triton,0,0,12.154532,triton,341.63,1111.41,0.000152896148989899 +gfx950,256,192,2304,5120,triton,0,0,12.292759,triton,368.5,1111.57,0.00013563368055555556 +gfx950,256,208,2304,5120,triton,0,0,14.062113,triton,348.98,982.78,0.00011476696047008547 +gfx950,256,224,2304,5120,triton,0,0,14.28005,triton,370.08,978.68,0.0001472594246031746 +gfx950,256,240,2304,5120,triton,0,0,14.744827,triton,384.02,958.38,0.00011574074074074075 +gfx950,256,256,2304,5120,triton,0,0,14.821024,triton,407.52,963.96,0.00012885199652777778 +gfx950,256,288,2304,5120,triton,0,0,16.699641,triton,406.88,874.16,0.0001431688850308642 +gfx950,256,320,2304,5120,triton,0,0,17.48011,triton,431.91,852.94,0.00012885199652777778 +gfx950,256,352,2304,5120,triton,0,0,17.046013,triton,487.19,892.92,0.00013809974747474747 +gfx950,256,384,2304,5120,triton,0,0,17.940548,triton,504.98,865.75,0.00012659143518518518 +gfx950,256,416,2304,5120,triton,0,0,17.04965,triton,575.65,929.25,0.00013980702457264956 +gfx950,256,448,2304,5120,triton,0,0,18.240602,triton,579.46,885.64,0.0001366024925595238 +gfx950,256,480,2304,5120,triton,0,0,22.036175,triton,513.91,747.22,0.00014919704861111112 +gfx950,256,512,2304,5120,triton,0,0,22.387963,triton,539.56,749.39,0.00014495849609375 +gfx950,256,544,2304,5120,triton,0,0,25.294405,triton,507.41,675.58,0.0001452078227124183 +gfx950,256,576,2304,5120,triton,0,0,25.338978,triton,536.31,686.68,0.0001243308738425926 +gfx950,256,608,2304,5120,triton,0,0,25.328086,triton,566.35,699.27,0.00014063071089181288 +gfx950,256,640,2304,5120,triton,0,0,25.653687,triton,588.59,702.53,0.0001614040798611111 +gfx950,256,672,2304,5120,triton,0,0,25.77706,triton,615.06,711.24,0.00015048879794973545 +gfx950,256,704,2304,5120,triton,0,0,25.662583,triton,647.22,726.54,0.00015906131628787878 +gfx950,256,736,2304,5120,triton,0,0,25.910372,triton,670.17,731.61,0.00014565877868357488 +gfx950,256,768,2304,5120,triton,0,0,26.622247,triton,680.61,723.74,0.00015936957465277778 +gfx950,256,800,2304,5120,triton,0,0,26.915113,triton,701.26,727.43,0.0001486545138888889 +gfx950,256,832,2304,5120,triton,0,0,26.479163,triton,741.31,751.16,0.00014919704861111112 +gfx950,256,864,2304,5120,triton,0,0,26.72306,triton,762.8,755.96,0.0001371407214506173 +gfx950,256,896,2304,5120,triton,0,0,27.075977,triton,780.74,757.6,0.0001346648685515873 +gfx950,256,928,2304,5120,triton,0,0,32.880238,triton,665.88,633.33,0.00015480947677203066 +gfx950,256,960,2304,5120,triton,0,0,33.33953,triton,679.35,633.94,0.00013970269097222221 +gfx950,256,992,2304,5120,triton,0,0,34.725084,triton,673.99,617.61,0.00014263412858422938 +gfx950,256,1024,2304,5120,triton,0,0,35.759795,triton,675.6,608.45,0.000133514404296875 +gfx950,256,1088,2304,5120,triton,0,0,33.971372,triton,755.61,658.81,0.00014042075163398692 +gfx950,256,1152,2304,5120,triton,0,0,36.296382,triton,748.81,633.76,0.0001488202883873457 +gfx950,256,1216,2304,5120,triton,0,0,35.40129,triton,810.4,667.37,0.00014170150310672513 +gfx950,256,1280,2304,5120,triton,0,0,36.5842,triton,825.47,662.81,0.00014682345920138888 +gfx950,256,1344,2304,5120,triton,0,0,35.7569,triton,886.79,695.56,0.0001369254298941799 +gfx950,256,1408,2304,5120,triton,0,0,37.551232,triton,884.63,678.9,0.00014364839804292928 +gfx950,256,1472,2304,5120,triton,0,0,35.9845,triton,965.11,725.76,0.0001512610394021739 +gfx950,256,1536,2304,5120,triton,0,0,38.09372,triton,951.31,701.92,0.00014552363642939814 +gfx950,256,1600,2304,5120,triton,0,0,37.07131,triton,1018.27,738.07,0.00014268663194444444 +gfx950,256,1664,2304,5120,triton,0,0,40.01893,triton,981.0,699.27,0.00014737121060363247 +gfx950,256,1728,2304,5120,triton,0,0,39.09733,triton,1042.75,731.67,0.00013764306841563786 +gfx950,256,1792,2304,5120,triton,0,0,39.72293,triton,1064.34,735.82,0.00014604840959821428 +gfx950,256,1856,2304,5120,triton,0,0,54.089141,triton,809.56,551.9,0.0001328274664750958 +gfx950,256,1920,2304,5120,triton,0,0,56.129351,triton,807.04,542.93,0.0001449019820601852 +gfx950,256,1984,2304,5120,triton,0,0,57.31908,triton,816.63,542.52,0.0001415403085797491 +gfx950,256,2048,2304,5120,triton,0,0,57.542221,triton,839.7,551.24,0.00014707777235243056 +gfx950,256,4096,2304,5120,triton,0,0,96.803179,triton,998.28,533.48,0.00014241536458333334 +gfx950,256,6144,2304,5120,triton,0,0,120.63743,triton,1201.58,593.23,0.00014615941930700232 +gfx950,256,8192,2304,5120,triton,0,0,161.36036,triton,1197.78,566.98,0.0001415146721733941 +gfx950,256,10240,2304,5120,triton,0,0,175.00967,triton,1380.45,636.6,0.00014517042371961806 +gfx950,256,12288,2304,5120,triton,0,0,210.985735,triton,1374.08,622.48,0.00014290986237702546 +gfx950,256,14336,2304,5120,triton,0,0,239.338825,triton,1413.18,631.98,0.00015040806361607144 +gfx950,256,16384,2304,5120,triton,0,0,272.917253,triton,1416.35,627.22,0.00014591217041015625 +gfx950,256,32768,2304,5120,triton,0,0,545.182216,triton,1418.05,606.34,0.00014662742614746094 +gfx950,256,1,4608,5120,triton,0,0,7.521773,triton,6.27,3138.53,0.00021701388888888888 +gfx950,256,2,4608,5120,triton,0,0,7.799511,triton,12.1,3028.6,0.00021701388888888888 +gfx950,256,4,4608,5120,triton,0,0,7.899558,triton,23.89,2993.88,0.00021701388888888888 +gfx950,256,8,4608,5120,triton,0,0,8.455489,triton,44.64,2803.82,8.138020833333333e-05 +gfx950,256,16,4608,5120,triton,0,0,8.509545,triton,88.72,2799.48,0.00018988715277777778 +gfx950,256,32,4608,5120,triton,0,0,9.268636,triton,162.91,2594.96,0.00013563368055555556 +gfx950,256,48,4608,5120,triton,0,0,10.676011,triton,212.15,2274.36,0.00010850694444444444 +gfx950,256,64,4608,5120,triton,0,0,10.889621,triton,277.32,2250.81,0.000152587890625 +gfx950,256,80,4608,5120,triton,0,0,14.425793,triton,261.68,1714.97,0.00014105902777777776 +gfx950,256,96,4608,5120,triton,0,0,14.960965,triton,302.78,1668.96,0.00014241536458333334 +gfx950,256,112,4608,5120,triton,0,0,15.19327,triton,347.84,1658.54,0.0001666356646825397 +gfx950,256,128,4608,5120,triton,0,0,15.801831,triton,382.22,1609.18,0.00013224283854166666 +gfx950,256,144,4608,5120,triton,0,0,16.901121,triton,402.03,1518.09,0.00014618296682098765 +gfx950,256,160,4608,5120,triton,0,0,17.141337,triton,440.44,1510.19,0.00014784071180555555 +gfx950,256,176,4608,5120,triton,0,0,17.486898,triton,474.91,1493.47,0.00012946851325757575 +gfx950,256,192,4608,5120,triton,0,0,17.539371,triton,516.53,1502.08,0.00014467592592592592 +gfx950,256,208,4608,5120,triton,0,0,22.791539,triton,430.63,1166.0,0.0001481537126068376 +gfx950,256,224,4608,5120,triton,0,0,23.043882,triton,458.67,1163.18,0.00014241536458333334 +gfx950,256,240,4608,5120,triton,0,0,23.569907,triton,480.47,1146.95,0.00015823929398148148 +gfx950,256,256,4608,5120,triton,0,0,23.966511,triton,504.02,1137.54,0.00012885199652777778 +gfx950,256,288,4608,5120,triton,0,0,24.534311,triton,553.9,1129.92,0.00014392240547839505 +gfx950,256,320,4608,5120,triton,0,0,24.995435,triton,604.09,1127.43,0.00014444986979166667 +gfx950,256,352,4608,5120,triton,0,0,27.19045,triton,610.86,1053.28,0.00014241536458333334 +gfx950,256,384,4608,5120,triton,0,0,27.64804,triton,655.36,1052.44,0.00012546115451388888 +gfx950,256,416,4608,5120,triton,0,0,28.66604,triton,684.76,1031.07,0.00013980702457264956 +gfx950,256,448,4608,5120,triton,0,0,29.570566,triton,714.88,1015.05,0.00014871264260912697 +gfx950,256,480,4608,5120,triton,0,0,34.669854,triton,653.28,878.98,0.00015733506944444445 +gfx950,256,512,4608,5120,triton,0,0,35.43093,triton,681.87,873.05,0.00015809800889756944 +gfx950,256,544,4608,5120,triton,0,0,33.71886,triton,761.27,930.98,0.00013762829350490196 +gfx950,256,576,4608,5120,triton,0,0,33.961044,triton,800.3,937.85,0.00014053156346450617 +gfx950,256,608,4608,5120,triton,0,0,35.500163,triton,808.14,910.12,0.00014134457236842106 +gfx950,256,640,4608,5120,triton,0,0,36.257989,triton,832.89,903.75,0.0001529269748263889 +gfx950,256,672,4608,5120,triton,0,0,35.2429,triton,899.73,942.79,0.00014596767526455027 +gfx950,256,704,4608,5120,triton,0,0,36.0573,triton,921.28,934.22,0.00015073834043560606 +gfx950,256,736,4608,5120,triton,0,0,37.30051,triton,931.06,915.38,0.0001453639228562802 +gfx950,256,768,4608,5120,triton,0,0,37.42332,triton,968.35,924.64,0.00014806676793981482 +gfx950,256,800,4608,5120,triton,0,0,36.36131,triton,1038.16,964.26,0.000146484375 +gfx950,256,832,4608,5120,triton,0,0,37.08971,triton,1058.48,957.69,0.0001385028545673077 +gfx950,256,864,4608,5120,triton,0,0,38.59173,triton,1056.41,932.31,0.00014593179333847737 +gfx950,256,896,4608,5120,triton,0,0,38.97093,triton,1084.87,935.01,0.00014774383060515872 +gfx950,256,928,4608,5120,triton,0,0,53.227093,triton,822.67,693.2,0.00013843989463601533 +gfx950,256,960,4608,5120,triton,0,0,56.37387,triton,803.54,662.64,0.00015010127314814815 +gfx950,256,992,4608,5120,triton,0,0,56.58507,triton,827.22,668.27,0.00014175907258064517 +gfx950,256,1024,4608,5120,triton,0,0,56.32308,triton,857.88,679.53,0.000148773193359375 +gfx950,256,1088,4608,5120,triton,0,0,57.54308,triton,892.17,681.06,0.0001412185968137255 +gfx950,256,1152,4608,5120,triton,0,0,59.5159,triton,913.34,673.9,0.00014015480324074074 +gfx950,256,1216,4608,5120,triton,0,0,60.53351,triton,947.87,677.73,0.00013991684941520467 +gfx950,256,1280,4608,5120,triton,0,0,61.01151,triton,989.94,687.46,0.0001361423068576389 +gfx950,256,1344,4608,5120,triton,0,0,62.30232,triton,1017.91,687.94,0.00014580620659722222 +gfx950,256,1408,4608,5120,triton,0,0,64.26834,triton,1033.76,681.18,0.00014780988596906566 +gfx950,256,1472,4608,5120,triton,0,0,65.87675,triton,1054.36,678.47,0.0001484599090428744 +gfx950,256,1536,4608,5120,triton,0,0,64.06995,triton,1131.23,711.93,0.0001373291015625 +gfx950,256,1600,4608,5120,triton,0,0,65.55197,triton,1151.72,709.83,0.00014431423611111112 +gfx950,256,1664,4608,5120,triton,0,0,66.20275,triton,1186.01,716.71,0.00014071994357638888 +gfx950,256,1728,4608,5120,triton,0,0,68.02999,triton,1198.55,710.94,0.00014844352816358026 +gfx950,256,1792,4608,5120,triton,0,0,69.34518,triton,1219.37,710.69,0.0001385401165674603 +gfx950,256,1856,4608,5120,triton,0,0,89.62055,triton,977.2,560.15,0.00014487080190373562 +gfx950,256,1920,4608,5120,triton,0,0,94.507,triton,958.63,540.89,0.0001361988208912037 +gfx950,256,1984,4608,5120,triton,0,0,96.62543,triton,968.86,538.53,0.0001475563186043907 +gfx950,256,2048,4608,5120,triton,0,0,96.06221,triton,1005.98,551.24,0.0001495149400499132 +gfx950,256,4096,4608,5120,triton,0,0,158.56153,triton,1218.92,519.12,0.0001434220208062066 +gfx950,256,6144,4608,5120,triton,0,0,215.09841,triton,1347.8,519.17,0.00014199150933159722 +gfx950,256,8192,4608,5120,triton,0,0,268.500273,triton,1439.65,525.26,0.00014442867702907987 +gfx950,256,10240,4608,5120,triton,0,0,325.346062,triton,1485.14,523.73,0.00014584859212239584 +gfx950,256,12288,4608,5120,triton,0,0,385.56424,triton,1503.82,518.08,0.0001453117088035301 +gfx950,256,14336,4608,5120,triton,0,0,471.347135,triton,1435.16,486.08,0.00014783465673053075 +gfx950,256,16384,4608,5120,triton,0,0,523.789777,triton,1475.96,493.47,0.00014506445990668403 +gfx950,256,32768,4608,5120,triton,0,0,1024.946711,triton,1508.55,481.35,0.0001411702897813585 +gfx950,256,1,5120,576,triton,0,0,3.21221,triton,1.84,921.46,0.0 +gfx950,256,2,5120,576,triton,0,0,3.24704,triton,3.63,914.91,9.765625e-05 +gfx950,256,4,5120,576,triton,0,0,3.186091,triton,7.4,939.2,0.0 +gfx950,256,8,5120,576,triton,0,0,3.218402,triton,14.66,943.22,0.0 +gfx950,256,16,5120,576,triton,0,0,3.22145,triton,29.29,969.18,0.0 +gfx950,256,32,5120,576,triton,0,0,4.159447,triton,45.38,792.23,6.103515625e-06 +gfx950,256,48,5120,576,triton,0,0,4.434505,triton,63.84,782.11,1.220703125e-05 +gfx950,256,64,5120,576,triton,0,0,4.651141,triton,81.16,782.89,6.103515625e-06 +gfx950,256,80,5120,576,triton,0,0,4.104577,triton,114.96,929.3,2.44140625e-05 +gfx950,256,96,5120,576,triton,0,0,4.179646,triton,135.47,954.02,6.103515625e-06 +gfx950,256,112,5120,576,triton,0,0,4.57465,triton,144.41,909.47,1.5694754464285713e-05 +gfx950,256,128,5120,576,triton,0,0,4.84234,triton,155.91,894.93,1.983642578125e-05 +gfx950,256,144,5120,576,triton,0,0,4.982724,triton,170.46,904.45,2.44140625e-05 +gfx950,256,160,5120,576,triton,0,0,5.565959,triton,169.55,840.77,2.0751953125e-05 +gfx950,256,176,5120,576,triton,0,0,5.592531,triton,185.62,867.72,1.6645951704545453e-05 +gfx950,256,192,5120,576,triton,0,0,5.668711,triton,199.77,886.58,1.52587890625e-05 +gfx950,256,208,5120,576,triton,0,0,5.784879,triton,212.08,898.7,1.690204326923077e-05 +gfx950,256,224,5120,576,triton,0,0,5.879043,triton,224.73,913.74,1.5694754464285713e-05 +gfx950,256,240,5120,576,triton,0,0,5.889773,triton,240.35,941.46,9.765625e-06 +gfx950,256,256,5120,576,triton,0,0,6.011825,triton,251.16,951.13,1.52587890625e-05 +gfx950,256,288,5120,576,triton,0,0,6.09707,triton,278.61,994.6,1.9666883680555557e-05 +gfx950,256,320,5120,576,triton,0,0,6.49982,triton,290.38,986.22,1.46484375e-05 +gfx950,256,352,5120,576,triton,0,0,6.85104,triton,303.05,986.18,1.4426491477272727e-05 +gfx950,256,384,5120,576,triton,0,0,6.916222,triton,327.48,1026.93,1.7293294270833333e-05 +gfx950,256,416,5120,576,triton,0,0,7.39265,triton,331.91,1007.57,1.690204326923077e-05 +gfx950,256,448,5120,576,triton,0,0,7.646737,triton,345.56,1019.35,1.9182477678571428e-05 +gfx950,256,480,5120,576,triton,0,0,7.779469,triton,363.93,1046.45,1.6276041666666666e-05 +gfx950,256,512,5120,576,triton,0,0,7.918649,triton,381.37,1071.76,1.6021728515625e-05 +gfx950,256,544,5120,576,triton,0,0,8.074474,triton,397.38,1093.94,1.579733455882353e-05 +gfx950,256,576,5120,576,triton,0,0,8.497792,triton,399.8,1080.18,1.797146267361111e-05 +gfx950,256,608,5120,576,triton,0,0,8.806189,triton,407.23,1081.65,1.606188322368421e-05 +gfx950,256,640,5120,576,triton,0,0,8.786485,triton,429.62,1123.47,1.953125e-05 +gfx950,256,672,5120,576,triton,0,0,8.30588,triton,477.21,1230.15,1.6276041666666666e-05 +gfx950,256,704,5120,576,triton,0,0,8.57428,triton,484.28,1232.01,1.5813654119318182e-05 +gfx950,256,736,5120,576,triton,0,0,8.738724,triton,496.77,1248.43,1.6452955163043478e-05 +gfx950,256,768,5120,576,triton,0,0,8.931232,triton,507.19,1260.27,1.6530354817708332e-05 +gfx950,256,800,5120,576,triton,0,0,10.46993,triton,450.68,1108.12,1.6357421875e-05 +gfx950,256,832,5120,576,triton,0,0,10.81203,triton,453.88,1105.07,1.7371544471153845e-05 +gfx950,256,864,5120,576,triton,0,0,11.202804,triton,454.89,1097.42,1.7180266203703705e-05 +gfx950,256,896,5120,576,triton,0,0,11.302461,triton,467.58,1118.36,1.5912737165178572e-05 +gfx950,256,928,5120,576,triton,0,0,11.041969,triton,495.71,1176.09,1.452215786637931e-05 +gfx950,256,960,5120,576,triton,0,0,11.321101,triton,500.16,1177.67,1.5055338541666667e-05 +gfx950,256,992,5120,576,triton,0,0,10.86229,triton,538.66,1259.27,2.0082535282258065e-05 +gfx950,256,1024,5120,576,triton,0,0,10.947758,triton,551.69,1281.06,1.811981201171875e-05 +gfx950,256,1088,5120,576,triton,0,0,11.6327,triton,551.66,1265.13,1.5438304227941176e-05 +gfx950,256,1152,5120,576,triton,0,0,10.95468,triton,620.26,1406.63,2.0005967881944445e-05 +gfx950,256,1216,5120,576,triton,0,0,11.91671,triton,601.87,1351.16,1.750745271381579e-05 +gfx950,256,1280,5120,576,triton,0,0,13.064418,triton,577.89,1285.45,1.800537109375e-05 +gfx950,256,1344,5120,576,triton,0,0,12.669735,triton,625.68,1380.13,1.613071986607143e-05 +gfx950,256,1408,5120,576,triton,0,0,12.80188,triton,648.71,1419.95,1.5952370383522727e-05 +gfx950,256,1472,5120,576,triton,0,0,12.695781,triton,683.87,1486.34,1.6187584918478262e-05 +gfx950,256,1536,5120,576,triton,0,0,12.967235,triton,698.66,1508.61,1.9327799479166666e-05 +gfx950,256,1600,5120,576,triton,0,0,14.113475,triton,668.66,1435.13,1.318359375e-05 +gfx950,256,1664,5120,576,triton,0,0,14.604768,triton,672.02,1434.25,1.690204326923077e-05 +gfx950,256,1728,5120,576,triton,0,0,16.410909,triton,621.06,1318.58,1.6841182002314814e-05 +gfx950,256,1792,5120,576,triton,0,0,15.606455,triton,677.26,1430.91,1.7002650669642858e-05 +gfx950,256,1856,5120,576,triton,0,0,15.30434,triton,715.3,1504.38,1.841577990301724e-05 +gfx950,256,1920,5120,576,triton,0,0,15.484573,triton,731.35,1531.58,1.77001953125e-05 +gfx950,256,1984,5120,576,triton,0,0,15.752612,triton,742.87,1549.46,1.6046339465725805e-05 +gfx950,256,2048,5120,576,triton,0,0,15.605235,triton,774.07,1608.45,1.52587890625e-05 +gfx950,256,4096,5120,576,triton,0,0,27.98125,triton,863.41,1688.68,1.850128173828125e-05 +gfx950,256,6144,5120,576,triton,0,0,42.78216,triton,847.05,1622.23,1.697540283203125e-05 +gfx950,256,8192,5120,576,triton,0,0,56.51707,triton,854.93,1619.93,1.6808509826660156e-05 +gfx950,256,10240,5120,576,triton,0,0,69.03598,triton,874.88,1647.04,1.483917236328125e-05 +gfx950,256,12288,5120,576,triton,0,0,81.53128,triton,888.95,1666.31,1.7595291137695312e-05 +gfx950,256,14336,5120,576,triton,0,0,95.49302,triton,885.48,1654.65,1.595360892159598e-05 +gfx950,256,16384,5120,576,triton,0,0,111.20754,triton,868.98,1620.02,1.710653305053711e-05 +gfx950,256,32768,5120,576,triton,0,0,210.181412,triton,919.56,1700.28,1.537203788757324e-05 +gfx950,256,1,5120,1152,triton,0,0,4.77544,triton,2.47,1237.51,0.0 +gfx950,256,2,5120,1152,triton,0,0,4.418265,triton,5.34,1340.12,0.0 +gfx950,256,4,5120,1152,triton,0,0,4.401051,triton,10.72,1350.54,0.0 +gfx950,256,8,5120,1152,triton,0,0,4.79026,triton,19.7,1250.32,4.8828125e-05 +gfx950,256,16,5120,1152,triton,0,0,4.53305,triton,41.64,1341.37,3.662109375e-05 +gfx950,256,32,5120,1152,triton,0,0,4.65664,triton,81.06,1344.91,4.8828125e-05 +gfx950,256,48,5120,1152,triton,0,0,5.55543,triton,101.92,1160.14,5.289713541666667e-05 +gfx950,256,64,5120,1152,triton,0,0,5.623908,triton,134.24,1178.42,3.0517578125e-05 +gfx950,256,80,5120,1152,triton,0,0,6.282283,triton,150.22,1083.94,3.173828125e-05 +gfx950,256,96,5120,1152,triton,0,0,5.84184,triton,193.85,1196.86,2.6448567708333334e-05 +gfx950,256,112,5120,1152,triton,0,0,6.503949,triton,203.14,1103.04,1.9182477678571428e-05 +gfx950,256,128,5120,1152,triton,0,0,7.188625,triton,210.05,1023.34,3.509521484375e-05 +gfx950,256,144,5120,1152,triton,0,0,7.080737,triton,239.9,1064.68,4.0690104166666664e-05 +gfx950,256,160,5120,1152,triton,0,0,7.210102,triton,261.78,1070.85,3.662109375e-05 +gfx950,256,176,5120,1152,triton,0,0,7.163515,triton,289.83,1103.26,3.107244318181818e-05 +gfx950,256,192,5120,1152,triton,0,0,7.319465,triton,309.44,1104.66,2.3396809895833335e-05 +gfx950,256,208,5120,1152,triton,0,0,8.777,triton,279.56,941.98,3.286508413461538e-05 +gfx950,256,224,5120,1152,triton,0,0,8.451146,triton,312.67,999.87,4.0108816964285714e-05 +gfx950,256,240,5120,1152,triton,0,0,8.61907,triton,328.48,1001.54,2.685546875e-05 +gfx950,256,256,5120,1152,triton,0,0,8.963122,triton,336.92,983.43,2.593994140625e-05 +gfx950,256,288,5120,1152,triton,0,0,8.992072,triton,377.82,1020.8,3.594292534722222e-05 +gfx950,256,320,5120,1152,triton,0,0,9.020558,triton,418.47,1057.99,3.173828125e-05 +gfx950,256,352,5120,1152,triton,0,0,9.232313,triton,449.76,1073.21,2.9407848011363636e-05 +gfx950,256,384,5120,1152,triton,0,0,9.167276,triton,494.13,1120.59,3.102620442708333e-05 +gfx950,256,416,5120,1152,triton,0,0,11.219701,triton,437.39,948.09,3.662109375e-05 +gfx950,256,448,5120,1152,triton,0,0,10.956742,triton,482.34,1004.12,3.444126674107143e-05 +gfx950,256,480,5120,1152,triton,0,0,11.235802,triton,503.95,1011.62,3.3772786458333334e-05 +gfx950,256,512,5120,1152,triton,0,0,11.66862,triton,517.61,1005.34,3.62396240234375e-05 +gfx950,256,544,5120,1152,triton,0,0,11.268677,triton,569.48,1073.37,3.087660845588235e-05 +gfx950,256,576,5120,1152,triton,0,0,11.53837,triton,588.88,1079.88,3.153483072916667e-05 +gfx950,256,608,5120,1152,triton,0,0,11.83668,triton,605.94,1083.46,3.212376644736842e-05 +gfx950,256,640,5120,1152,triton,0,0,11.710596,triton,644.69,1126.26,3.23486328125e-05 +gfx950,256,672,5120,1152,triton,0,0,12.405333,triton,639.02,1092.57,2.8773716517857142e-05 +gfx950,256,704,5120,1152,triton,0,0,11.867448,triton,699.79,1172.81,2.9962713068181818e-05 +gfx950,256,736,5120,1152,triton,0,0,12.19111,triton,712.18,1171.57,3.7947944972826085e-05 +gfx950,256,768,5120,1152,triton,0,0,12.596293,triton,719.24,1162.83,2.8737386067708333e-05 +gfx950,256,800,5120,1152,triton,0,0,14.220515,triton,663.63,1055.65,3.173828125e-05 +gfx950,256,832,5120,1152,triton,0,0,14.763277,triton,664.8,1041.53,3.216083233173077e-05 +gfx950,256,864,5120,1152,triton,0,0,15.016265,triton,678.74,1048.26,3.729926215277778e-05 +gfx950,256,896,5120,1152,triton,0,0,15.381633,triton,687.16,1047.06,3.095354352678571e-05 +gfx950,256,928,5120,1152,triton,0,0,16.066372,triton,681.37,1025.12,3.472689924568966e-05 +gfx950,256,960,5120,1152,triton,0,0,15.829552,triton,715.41,1063.49,3.72314453125e-05 +gfx950,256,992,5120,1152,triton,0,0,16.099106,triton,726.88,1068.33,3.425844254032258e-05 +gfx950,256,1024,5120,1152,triton,0,0,16.085802,triton,750.95,1091.87,3.108978271484375e-05 +gfx950,256,1088,5120,1152,triton,0,0,16.24432,triton,790.1,1126.1,3.0517578125e-05 +gfx950,256,1152,5120,1152,triton,0,0,16.649521,triton,816.21,1142.48,3.594292534722222e-05 +gfx950,256,1216,5120,1152,triton,0,0,16.755545,triton,856.11,1178.77,3.212376644736842e-05 +gfx950,256,1280,5120,1152,triton,0,0,16.980629,triton,889.22,1206.08,3.4332275390625e-05 +gfx950,256,1344,5120,1152,triton,0,0,17.477474,triton,907.14,1213.51,3.197079613095238e-05 +gfx950,256,1408,5120,1152,triton,0,0,17.334,triton,958.2,1265.62,3.273703835227273e-05 +gfx950,256,1472,5120,1152,triton,0,0,18.08133,triton,960.35,1253.63,3.065026324728261e-05 +gfx950,256,1536,5120,1152,triton,0,0,17.69163,triton,1024.18,1322.45,3.2806396484375e-05 +gfx950,256,1600,5120,1152,triton,0,0,22.59,triton,835.52,1067.97,3.02734375e-05 +gfx950,256,1664,5120,1152,triton,0,0,23.08965,triton,850.14,1076.44,2.711369441105769e-05 +gfx950,256,1728,5120,1152,triton,0,0,23.125704,triton,881.46,1106.28,3.2099971064814814e-05 +gfx950,256,1792,5120,1152,triton,0,0,24.243273,triton,871.97,1085.36,3.5858154296875e-05 +gfx950,256,1856,5120,1152,triton,0,0,23.958758,triton,913.83,1128.68,3.946238550646552e-05 +gfx950,256,1920,5120,1152,triton,0,0,24.279141,triton,932.87,1143.82,3.0415852864583334e-05 +gfx950,256,1984,5120,1152,triton,0,0,24.693629,triton,947.78,1154.14,3.169890372983871e-05 +gfx950,256,2048,5120,1152,triton,0,0,24.12118,triton,1001.58,1211.76,3.1566619873046876e-05 +gfx950,256,4096,5120,1152,triton,0,0,43.28618,triton,1116.25,1214.24,3.3235549926757814e-05 +gfx950,256,6144,5120,1152,triton,0,0,64.18514,triton,1129.2,1182.37,3.147125244140625e-05 +gfx950,256,8192,5120,1152,triton,0,0,83.6473,triton,1155.29,1186.19,3.170967102050781e-05 +gfx950,256,10240,5120,1152,triton,0,0,105.99469,triton,1139.64,1156.21,3.1585693359375e-05 +gfx950,256,12288,5120,1152,triton,0,0,123.95363,triton,1169.43,1176.92,3.4300486246744793e-05 +gfx950,256,14336,5120,1152,triton,0,0,148.84028,triton,1136.21,1136.88,3.0735560825892856e-05 +gfx950,256,16384,5120,1152,triton,0,0,169.73125,triton,1138.7,1134.41,3.594160079956055e-05 +gfx950,256,32768,5120,1152,triton,0,0,307.666732,triton,1256.38,1232.47,3.5423040390014646e-05 +gfx950,256,1,5120,2304,triton,0,0,5.41905,triton,4.35,2179.17,0.000390625 +gfx950,256,2,5120,2304,triton,0,0,5.484062,triton,8.6,2155.62,9.765625e-05 +gfx950,256,4,5120,2304,triton,0,0,5.474525,triton,17.24,2163.96,0.000146484375 +gfx950,256,8,5120,2304,triton,0,0,5.576012,triton,33.85,2133.57,2.44140625e-05 +gfx950,256,16,5120,2304,triton,0,0,6.308188,triton,59.84,1901.84,0.00010986328125 +gfx950,256,32,5120,2304,triton,0,0,5.99241,triton,125.99,2035.56,5.4931640625e-05 +gfx950,256,48,5120,2304,triton,0,0,7.412122,triton,152.79,1672.75,6.917317708333333e-05 +gfx950,256,64,5120,2304,triton,0,0,7.545313,triton,200.12,1669.82,8.85009765625e-05 +gfx950,256,80,5120,2304,triton,0,0,9.461675,triton,199.48,1352.83,4.8828125e-05 +gfx950,256,96,5120,2304,triton,0,0,9.703765,triton,233.41,1339.76,4.679361979166667e-05 +gfx950,256,112,5120,2304,triton,0,0,9.665708,triton,273.38,1365.8,6.452287946428572e-05 +gfx950,256,128,5120,2304,triton,0,0,9.846082,triton,306.71,1361.16,5.4931640625e-05 +gfx950,256,144,5120,2304,triton,0,0,10.423571,triton,325.93,1305.01,5.425347222222222e-05 +gfx950,256,160,5120,2304,triton,0,0,9.993309,triton,377.74,1381.28,6.7138671875e-05 +gfx950,256,176,5120,2304,triton,0,0,10.088957,triton,411.57,1388.07,7.32421875e-05 +gfx950,256,192,5120,2304,triton,0,0,10.11989,triton,447.62,1403.66,6.917317708333333e-05 +gfx950,256,208,5120,2304,triton,0,0,12.273954,triton,399.82,1173.67,6.760817307692308e-05 +gfx950,256,224,5120,2304,triton,0,0,12.379329,triton,426.91,1179.9,6.975446428571428e-05 +gfx950,256,240,5120,2304,triton,0,0,12.926349,triton,438.04,1145.49,7.242838541666667e-05 +gfx950,256,256,5120,2304,triton,0,0,13.090157,triton,461.4,1146.49,6.1798095703125e-05 +gfx950,256,288,5120,2304,triton,0,0,13.467272,triton,504.54,1144.19,7.188585069444445e-05 +gfx950,256,320,5120,2304,triton,0,0,13.389506,triton,563.86,1180.82,6.89697265625e-05 +gfx950,256,352,5120,2304,triton,0,0,14.204649,triton,584.65,1141.31,5.937056107954546e-05 +gfx950,256,384,5120,2304,triton,0,0,14.132157,triton,641.07,1175.57,5.900065104166667e-05 +gfx950,256,416,5120,2304,triton,0,0,15.955626,triton,615.12,1066.38,6.385216346153847e-05 +gfx950,256,448,5120,2304,triton,0,0,16.221742,triton,651.57,1073.63,5.7547433035714285e-05 +gfx950,256,480,5120,2304,triton,0,0,16.66111,triton,679.7,1069.41,5.8186848958333335e-05 +gfx950,256,512,5120,2304,triton,0,0,16.986247,triton,711.14,1072.57,7.171630859375e-05 +gfx950,256,544,5120,2304,triton,0,0,16.80083,triton,763.92,1108.3,7.108800551470589e-05 +gfx950,256,576,5120,2304,triton,0,0,16.478234,triton,824.7,1154.36,6.951226128472222e-05 +gfx950,256,608,5120,2304,triton,0,0,16.772927,triton,855.22,1158.01,6.553248355263158e-05 +gfx950,256,640,5120,2304,triton,0,0,16.798689,triton,898.85,1180.13,6.28662109375e-05 +gfx950,256,672,5120,2304,triton,0,0,16.697053,triton,949.54,1211.35,6.626674107142857e-05 +gfx950,256,704,5120,2304,triton,0,0,16.971726,triton,978.65,1215.4,5.964799360795455e-05 +gfx950,256,736,5120,2304,triton,0,0,17.449366,triton,995.13,1205.14,6.368885869565217e-05 +gfx950,256,768,5120,2304,triton,0,0,17.927402,triton,1010.71,1195.39,6.993611653645833e-05 +gfx950,256,800,5120,2304,triton,0,0,23.958758,triton,787.79,911.22,6.689453125e-05 +gfx950,256,832,5120,2304,triton,0,0,24.425181,triton,803.65,910.25,6.807767427884616e-05 +gfx950,256,864,5120,2304,triton,0,0,24.859177,triton,819.99,910.51,7.143373842592592e-05 +gfx950,256,896,5120,2304,triton,0,0,24.931798,triton,847.88,923.96,6.561279296875e-05 +gfx950,256,928,5120,2304,triton,0,0,26.78644,triton,817.36,874.97,6.924333243534482e-05 +gfx950,256,960,5120,2304,triton,0,0,26.188717,triton,864.85,910.27,6.40869140625e-05 +gfx950,256,992,5120,2304,triton,0,0,26.56832,triton,880.91,912.37,6.910754788306452e-05 +gfx950,256,1024,5120,2304,triton,0,0,27.00203,triton,894.72,912.58,6.427764892578126e-05 +gfx950,256,1088,5120,2304,triton,0,0,28.18401,triton,910.77,902.79,6.265079273897058e-05 +gfx950,256,1152,5120,2304,triton,0,0,28.688309,triton,947.39,914.91,7.714165581597222e-05 +gfx950,256,1216,5120,2304,triton,0,0,28.75149,triton,997.83,940.82,6.18382504111842e-05 +gfx950,256,1280,5120,2304,triton,0,0,29.33605,triton,1029.42,949.44,6.8359375e-05 +gfx950,256,1344,5120,2304,triton,0,0,29.96766,triton,1058.11,956.22,6.365094866071429e-05 +gfx950,256,1408,5120,2304,triton,0,0,31.11166,triton,1067.73,946.86,6.117387251420454e-05 +gfx950,256,1472,5120,2304,triton,0,0,30.48286,triton,1139.29,992.73,6.899626358695652e-05 +gfx950,256,1536,5120,2304,triton,0,0,31.62567,triton,1145.87,982.24,6.268819173177083e-05 +gfx950,256,1600,5120,2304,triton,0,0,42.338939,triton,891.58,752.66,6.8115234375e-05 +gfx950,256,1664,5120,2304,triton,0,0,44.67217,triton,878.82,731.32,6.67865459735577e-05 +gfx950,256,1728,5120,2304,triton,0,0,45.37138,triton,898.55,737.75,6.679958767361112e-05 +gfx950,256,1792,5120,2304,triton,0,0,45.38658,triton,931.52,755.19,6.354195731026786e-05 +gfx950,256,1856,5120,2304,triton,0,0,46.365918,triton,944.41,756.55,6.734913793103448e-05 +gfx950,256,1920,5120,2304,triton,0,0,48.92821,triton,925.82,733.34,6.479899088541667e-05 +gfx950,256,1984,5120,2304,triton,0,0,48.0546,triton,974.07,763.38,6.566201486895162e-05 +gfx950,256,2048,5120,2304,triton,0,0,49.27381,triton,980.61,760.78,6.275177001953125e-05 +gfx950,256,4096,5120,2304,triton,0,0,82.9705,triton,1164.71,761.44,6.642341613769532e-05 +gfx950,256,6144,5120,2304,triton,0,0,119.4784,triton,1213.23,743.79,6.847381591796875e-05 +gfx950,256,8192,5120,2304,triton,0,0,151.16386,triton,1278.57,757.83,6.799697875976562e-05 +gfx950,256,10240,5120,2304,triton,0,0,180.791856,triton,1336.3,775.74,6.776809692382813e-05 +gfx950,256,12288,5120,2304,triton,0,0,209.069216,triton,1386.67,793.69,6.380081176757813e-05 +gfx950,256,14336,5120,2304,triton,0,0,245.585794,triton,1377.23,780.29,6.60078866141183e-05 +gfx950,256,16384,5120,2304,triton,0,0,277.36174,triton,1393.66,783.52,6.481409072875977e-05 +gfx950,256,32768,5120,2304,triton,0,0,527.926368,triton,1464.4,800.94,6.707310676574707e-05 +gfx950,256,1,25600,6144,triton,0,0,29.240082,triton,10.76,5381.1,0.0001953125 +gfx950,256,2,25600,6144,triton,0,0,30.49806,triton,20.63,5161.02,0.0001171875 +gfx950,256,4,25600,6144,triton,0,0,30.80607,triton,40.85,5113.14,0.00013671875 +gfx950,256,8,25600,6144,triton,0,0,30.775194,triton,81.77,5125.72,0.00017578125 +gfx950,256,16,25600,6144,triton,0,0,32.632268,triton,154.24,4848.08,0.00016357421875 +gfx950,256,32,25600,6144,triton,0,0,35.253455,triton,285.54,4513.64,0.00015869140625 +gfx950,256,48,25600,6144,triton,0,0,37.373051,triton,404.02,4282.2,0.00016031901041666666 +gfx950,256,64,25600,6144,triton,0,0,38.946404,triton,516.93,4132.77,0.0001715087890625 +gfx950,256,80,25600,6144,triton,0,0,45.257556,triton,556.06,3576.73,0.00018310546875 +gfx950,256,96,25600,6144,triton,0,0,47.6542,triton,633.71,3416.1,0.00015869140625 +gfx950,256,112,25600,6144,triton,0,0,49.93502,triton,705.56,3278.44,0.00016566685267857144 +gfx950,256,128,25600,6144,triton,0,0,47.4558,triton,848.48,3469.05,0.0001666259765625 +gfx950,256,144,25600,6144,triton,0,0,68.710293,triton,659.27,2409.3,0.00016737196180555556 +gfx950,256,160,25600,6144,triton,0,0,73.57442,triton,684.09,2262.49,0.000173828125 +gfx950,256,176,25600,6144,triton,0,0,73.60362,triton,752.2,2274.06,0.00016978870738636363 +gfx950,256,192,25600,6144,triton,0,0,75.00563,triton,805.25,2243.78,0.00015604654947916667 +gfx950,256,208,25600,6144,triton,0,0,75.96885,triton,861.29,2227.41,0.0001605694110576923 +gfx950,256,224,25600,6144,triton,0,0,82.7285,triton,851.75,2056.5,0.00017333984375 +gfx950,256,240,25600,6144,triton,0,0,81.67409,triton,924.37,2094.29,0.00016422526041666666 +gfx950,256,256,25600,6144,triton,0,0,80.49689,triton,1000.42,2136.31,0.00016510009765625 +gfx950,256,288,25600,6144,triton,0,0,110.853,triton,817.27,1567.86,0.0001650661892361111 +gfx950,256,320,25600,6144,triton,0,0,116.96358,triton,860.64,1501.63,0.00016748046875 +gfx950,256,352,25600,6144,triton,0,0,117.4156,triton,943.06,1511.48,0.00016745827414772727 +gfx950,256,384,25600,6144,triton,0,0,119.873939,triton,1007.69,1495.79,0.00016998291015625 +gfx950,256,416,25600,6144,triton,0,0,136.906,triton,955.86,1323.11,0.00016948993389423076 +gfx950,256,448,25600,6144,triton,0,0,143.33463,triton,983.21,1276.57,0.00016932896205357144 +gfx950,256,480,25600,6144,triton,0,0,145.120213,triton,1040.48,1273.51,0.00016609700520833334 +gfx950,256,512,25600,6144,triton,0,0,144.77464,triton,1112.5,1289.22,0.0001676177978515625 +gfx950,256,544,25600,6144,triton,0,0,153.49672,triton,1114.86,1227.92,0.00016974954044117647 +gfx950,256,576,25600,6144,triton,0,0,156.80151,triton,1155.56,1213.74,0.00016798231336805555 +gfx950,256,608,25600,6144,triton,0,0,154.9799,triton,1234.1,1239.85,0.00016440943667763157 +gfx950,256,640,25600,6144,triton,0,0,157.18192,triton,1280.85,1234.15,0.00016656494140625 +gfx950,256,672,25600,6144,triton,0,0,191.457928,triton,1104.12,1022.79,0.0001708984375 +gfx950,256,704,25600,6144,triton,0,0,198.55347,triton,1115.36,995.48,0.00017001065340909092 +gfx950,256,736,25600,6144,triton,0,0,199.05791,triton,1163.11,1002.18,0.0001718537703804348 +gfx950,256,768,25600,6144,triton,0,0,195.032878,triton,1238.72,1032.27,0.00017130533854166667 +gfx950,256,800,25600,6144,triton,0,0,210.27879,triton,1196.78,966.15,0.000168701171875 +gfx950,256,832,25600,6144,triton,0,0,215.76807,triton,1212.99,950.08,0.00017315204326923077 +gfx950,256,864,25600,6144,triton,0,0,212.532921,triton,1278.82,973.17,0.00016877350983796296 +gfx950,256,896,25600,6144,triton,0,0,217.79763,triton,1294.12,958.08,0.00017338344029017857 +gfx950,256,928,25600,6144,triton,0,0,242.75863,triton,1202.53,867.12,0.00016357421875 +gfx950,256,960,25600,6144,triton,0,0,255.39713,triton,1182.43,831.4,0.00017313639322916665 +gfx950,256,992,25600,6144,triton,0,0,255.497747,triton,1221.37,838.25,0.00017074092741935484 +gfx950,256,1024,25600,6144,triton,0,0,251.03274,triton,1283.19,860.47,0.0001702117919921875 +gfx950,256,1088,25600,6144,triton,0,0,266.163837,triton,1285.88,825.34,0.00017215504365808825 +gfx950,256,1152,25600,6144,triton,0,0,279.994131,triton,1294.27,797.68,0.00016920301649305555 +gfx950,256,1216,25600,6144,triton,0,0,276.354388,triton,1384.17,821.47,0.0001685534025493421 +gfx950,256,1280,25600,6144,triton,0,0,282.158092,triton,1427.05,817.58,0.000171905517578125 +gfx950,256,1344,25600,6144,triton,0,0,314.55352,triton,1344.08,745.05,0.00016688755580357144 +gfx950,256,1408,25600,6144,triton,0,0,331.348739,triton,1336.71,718.36,0.0001672640713778409 +gfx950,256,1472,25600,6144,triton,0,0,339.6152,triton,1363.46,711.68,0.00016734247622282608 +gfx950,256,1536,25600,6144,triton,0,0,333.817532,triton,1447.45,735.03,0.00017514546712239584 +gfx950,256,1600,25600,6144,triton,0,0,365.799138,triton,1375.94,680.8,0.0001669189453125 +gfx950,256,1664,25600,6144,triton,0,0,381.7558,triton,1371.16,661.96,0.00016923170823317308 +gfx950,256,1728,25600,6144,triton,0,0,379.557236,triton,1432.15,675.46,0.00016612865306712962 +gfx950,256,1792,25600,6144,triton,0,0,384.295151,triton,1466.88,676.69,0.00017364501953125 +gfx950,256,1856,25600,6144,triton,0,0,410.91484,triton,1420.85,641.78,0.0001660577182112069 +gfx950,256,1920,25600,6144,triton,0,0,432.070581,triton,1397.87,618.85,0.000170654296875 +gfx950,256,1984,25600,6144,triton,0,0,448.03586,triton,1393.0,604.99,0.0001715087890625 +gfx950,256,2048,25600,6144,triton,0,0,437.631391,triton,1472.12,627.76,0.00017425537109375 +gfx950,256,4096,25600,6144,triton,0,0,805.705137,triton,1599.21,486.74,0.0001657867431640625 +gfx950,256,6144,25600,6144,triton,0,0,1198.745121,triton,1612.3,425.12,0.00016859054565429688 +gfx950,256,8192,25600,6144,triton,0,0,1557.091232,triton,1655.0,402.71,0.00016988754272460936 +gfx950,256,10240,25600,6144,triton,0,0,1945.549885,triton,1655.69,382.66,0.0001712493896484375 +gfx950,256,12288,25600,6144,triton,0,0,2313.956463,triton,1670.5,372.49,0.00016946156819661458 +gfx950,256,14336,25600,6144,triton,0,0,2822.03852,triton,1598.03,347.04,0.00016983577183314732 +gfx950,256,16384,25600,6144,triton,0,0,3214.41882,triton,1603.39,341.22,0.00017205238342285156 +gfx950,256,32768,25600,6144,triton,0,0,6258.51091,triton,1647.02,325.37,0.00017000913619995118 +gfx950,256,1,5120,15360,triton,0,0,16.247906,triton,9.68,4841.78,0.0001953125 +gfx950,256,2,5120,15360,triton,0,0,16.531789,triton,19.03,4760.19,0.000390625 +gfx950,256,4,5120,15360,triton,0,0,16.546229,triton,38.02,4759.13,0.000439453125 +gfx950,256,8,5120,15360,triton,0,0,18.139149,triton,69.37,4346.84,0.0002685546875 +gfx950,256,16,5120,15360,triton,0,0,17.782106,triton,141.52,4445.64,0.000390625 +gfx950,256,32,5120,15360,triton,0,0,18.161791,triton,277.13,4375.25,0.00040283203125 +gfx950,256,48,5120,15360,triton,0,0,20.775989,triton,363.39,3844.44,0.00036214192708333335 +gfx950,256,64,5120,15360,triton,0,0,21.075942,triton,477.62,3809.16,0.0003326416015625 +gfx950,256,80,5120,15360,triton,0,0,32.040912,triton,392.71,2518.38,0.0003076171875 +gfx950,256,96,5120,15360,triton,0,0,32.808261,triton,460.23,2471.96,0.000299072265625 +gfx950,256,112,5120,15360,triton,0,0,33.192808,triton,530.72,2455.66,0.00034005301339285713 +gfx950,256,128,5120,15360,triton,0,0,34.651828,triton,581.0,2364.09,0.00030364990234375 +gfx950,256,144,5120,15360,triton,0,0,44.311707,triton,511.13,1857.96,0.0003363715277777778 +gfx950,256,160,5120,15360,triton,0,0,46.33938,triton,543.08,1785.51,0.00036376953125 +gfx950,256,176,5120,15360,triton,0,0,46.57059,triton,594.42,1785.44,0.00033180930397727274 +gfx950,256,192,5120,15360,triton,0,0,46.372717,triton,651.22,1801.89,0.00032857259114583336 +gfx950,256,208,5120,15360,triton,0,0,52.402635,triton,624.31,1602.36,0.0003126878004807692 +gfx950,256,224,5120,15360,triton,0,0,52.59246,triton,669.91,1604.37,0.00033743722098214285 +gfx950,256,240,5120,15360,triton,0,0,52.829209,triton,714.54,1604.93,0.00034505208333333333 +gfx950,256,256,5120,15360,triton,0,0,53.203966,triton,756.81,1601.32,0.000316619873046875 +gfx950,256,288,5120,15360,triton,0,0,56.068155,triton,807.92,1534.13,0.00030042860243055557 +gfx950,256,320,5120,15360,triton,0,0,59.28448,triton,848.99,1464.72,0.0003369140625 +gfx950,256,352,5120,15360,triton,0,0,59.256263,triton,934.33,1479.24,0.0003029563210227273 +gfx950,256,384,5120,15360,triton,0,0,58.94689,triton,1024.62,1500.9,0.0003265380859375 +gfx950,256,416,5120,15360,triton,0,0,83.02392,triton,788.1,1075.51,0.0003211388221153846 +gfx950,256,448,5120,15360,triton,0,0,85.533263,triton,823.82,1053.53,0.0003348214285714286 +gfx950,256,480,5120,15360,triton,0,0,88.24601,triton,855.53,1030.43,0.000338134765625 +gfx950,256,512,5120,15360,triton,0,0,89.571283,triton,899.07,1024.33,0.0003337860107421875 +gfx950,256,544,5120,15360,triton,0,0,91.650687,triton,933.59,1010.03,0.0003364114200367647 +gfx950,256,576,5120,15360,triton,0,0,95.527081,triton,948.39,977.62,0.00030992296006944443 +gfx950,256,608,5120,15360,triton,0,0,93.142719,triton,1026.71,1011.44,0.000348864103618421 +gfx950,256,640,5120,15360,triton,0,0,91.94101,triton,1094.87,1033.57,0.0003363037109375 +gfx950,256,672,5120,15360,triton,0,0,96.92903,triton,1090.45,988.83,0.0003287179129464286 +gfx950,256,704,5120,15360,triton,0,0,99.40106,triton,1113.97,972.48,0.0003368030894886364 +gfx950,256,736,5120,15360,triton,0,0,100.564707,triton,1151.13,969.37,0.00033489724864130434 +gfx950,256,768,5120,15360,triton,0,0,96.012102,triton,1258.13,1023.87,0.0003300984700520833 +gfx950,256,800,5120,15360,triton,0,0,137.777677,triton,913.28,719.44,0.00032666015625 +gfx950,256,832,5120,15360,triton,0,0,142.836766,triton,916.17,699.7,0.0003441443810096154 +gfx950,256,864,5120,15360,triton,0,0,143.720663,triton,945.55,701.09,0.0003352412471064815 +gfx950,256,896,5120,15360,triton,0,0,143.686558,triton,980.81,706.96,0.00034005301339285713 +gfx950,256,928,5120,15360,triton,0,0,150.883303,triton,967.38,678.67,0.0003247491244612069 +gfx950,256,960,5120,15360,triton,0,0,152.39228,triton,990.83,677.33,0.0003226725260416667 +gfx950,256,992,5120,15360,triton,0,0,152.43827,triton,1023.55,682.5,0.0003433719758064516 +gfx950,256,1024,5120,15360,triton,0,0,153.412808,triton,1049.86,683.5,0.00033359527587890626 +gfx950,256,1088,5120,15360,triton,0,0,157.483152,triton,1086.64,676.24,0.00033982220818014704 +gfx950,256,1152,5120,15360,triton,0,0,163.84958,triton,1105.86,659.96,0.0003246731228298611 +gfx950,256,1216,5120,15360,triton,0,0,167.320818,triton,1143.08,656.06,0.0003403513055098684 +gfx950,256,1280,5120,15360,triton,0,0,164.86718,triton,1221.14,675.76,0.0003240966796875 +gfx950,256,1344,5120,15360,triton,0,0,168.24325,triton,1256.47,671.94,0.00033467610677083334 +gfx950,256,1408,5120,15360,triton,0,0,170.35363,triton,1300.0,673.23,0.00033888383345170457 +gfx950,256,1472,5120,15360,triton,0,0,171.50404,triton,1349.97,678.27,0.0003222921620244565 +gfx950,256,1536,5120,15360,triton,0,0,175.42127,triton,1377.21,672.47,0.0003363291422526042 +gfx950,256,1600,5120,15360,triton,0,0,232.7975,triton,1081.02,513.76,0.0003311767578125 +gfx950,256,1664,5120,15360,triton,0,0,240.960439,triton,1086.17,503.16,0.0003248948317307692 +gfx950,256,1728,5120,15360,triton,0,0,252.411181,triton,1076.78,486.82,0.0003311722366898148 +gfx950,256,1792,5120,15360,triton,0,0,250.715144,triton,1124.21,496.65,0.0003291538783482143 +gfx950,256,1856,5120,15360,triton,0,0,256.15974,triton,1139.62,492.49,0.0003258014547413793 +gfx950,256,1920,5120,15360,triton,0,0,263.087959,triton,1147.87,485.75,0.00033376057942708336 +gfx950,256,1984,5120,15360,triton,0,0,257.357131,triton,1212.54,502.93,0.00033333070816532256 +gfx950,256,2048,5120,15360,triton,0,0,266.637355,triton,1208.09,491.57,0.00034351348876953124 +gfx950,256,4096,5120,15360,triton,0,0,429.50499,triton,1499.97,427.24,0.00033769607543945315 +gfx950,256,6144,5120,15360,triton,0,0,639.601093,triton,1510.89,368.87,0.00033178329467773435 +gfx950,256,8192,5120,15360,triton,0,0,809.422146,triton,1591.86,356.25,0.00033118724822998045 +gfx950,256,10240,5120,15360,triton,0,0,1019.213097,triton,1580.25,334.36,0.00033184051513671873 +gfx950,256,12288,5120,15360,triton,0,0,1178.278758,triton,1640.3,333.72,0.00033127466837565105 +gfx950,256,14336,5120,15360,triton,0,0,1405.591645,triton,1604.21,317.05,0.0003391402108328683 +gfx950,256,16384,5120,15360,triton,0,0,1555.738326,triton,1656.44,320.15,0.0003303050994873047 +gfx950,256,32768,5120,15360,triton,0,0,3120.545495,triton,1651.62,294.02,0.0003316521644592285 +gfx950,256,1,1536,5120,triton,0,0,5.58999,triton,2.81,1408.32,0.0006510416666666666 +gfx950,256,2,1536,5120,triton,0,0,5.29143,triton,5.94,1489.33,0.0 +gfx950,256,4,1536,5120,triton,0,0,5.293724,triton,11.88,1491.78,0.00016276041666666666 +gfx950,256,8,1536,5120,triton,0,0,5.697816,triton,22.08,1391.74,0.0003255208333333333 +gfx950,256,16,1536,5120,triton,0,0,5.609371,triton,44.86,1425.36,0.0001220703125 +gfx950,256,32,1536,5120,triton,0,0,5.666649,triton,88.82,1434.09,0.000244140625 +gfx950,256,48,1536,5120,triton,0,0,6.847905,triton,110.25,1205.85,0.00014919704861111112 +gfx950,256,64,1536,5120,triton,0,0,6.68266,triton,150.63,1255.28,0.00013224283854166666 +gfx950,256,80,1536,5120,triton,0,0,8.56028,triton,146.99,995.26,0.000146484375 +gfx950,256,96,1536,5120,triton,0,0,8.255567,triton,182.9,1047.87,0.00010172526041666667 +gfx950,256,112,1536,5120,triton,0,0,8.419956,triton,209.22,1042.98,0.00015694754464285715 +gfx950,256,128,1536,5120,triton,0,0,8.573796,triton,234.82,1039.55,0.0001373291015625 +gfx950,256,144,1536,5120,triton,0,0,9.55069,triton,237.15,946.94,0.00014919704861111112 +gfx950,256,160,1536,5120,triton,0,0,9.605101,triton,262.0,955.23,0.00010579427083333334 +gfx950,256,176,1536,5120,triton,0,0,10.522063,triton,263.09,884.44,0.0001553622159090909 +gfx950,256,192,1536,5120,triton,0,0,10.55213,triton,286.19,894.34,0.000152587890625 +gfx950,256,208,1536,5120,triton,0,0,11.547854,triton,283.3,828.57,0.00017528044871794872 +gfx950,256,224,1536,5120,triton,0,0,11.2092,triton,314.31,865.3,0.00015404110863095238 +gfx950,256,240,1536,5120,triton,0,0,11.957495,triton,315.69,822.11,0.00013563368055555556 +gfx950,256,256,1536,5120,triton,0,0,12.338041,triton,326.35,807.38,0.00015513102213541666 +gfx950,256,288,1536,5120,triton,0,0,11.534869,triton,392.71,886.32,0.00014015480324074074 +gfx950,256,320,1536,5120,triton,0,0,12.075465,triton,416.81,868.35,0.000140380859375 +gfx950,256,352,1536,5120,triton,0,0,14.482926,triton,382.28,742.11,0.0001220703125 +gfx950,256,384,1536,5120,triton,0,0,14.792273,triton,408.31,744.31,0.00014919704861111112 +gfx950,256,416,1536,5120,triton,0,0,15.52196,triton,421.54,726.21,0.00014554537259615384 +gfx950,256,448,1536,5120,triton,0,0,15.445396,triton,456.22,746.78,0.00016276041666666666 +gfx950,256,480,1536,5120,triton,0,0,16.91997,triton,446.2,697.19,0.00015190972222222222 +gfx950,256,512,1536,5120,triton,0,0,17.35794,triton,463.94,694.7,0.00014495849609375 +gfx950,256,544,1536,5120,triton,0,0,16.949745,triton,504.81,726.9,0.00014600566789215686 +gfx950,256,576,1536,5120,triton,0,0,17.687726,triton,512.2,711.39,0.0001514576099537037 +gfx950,256,608,1536,5120,triton,0,0,17.068713,triton,560.27,752.55,0.0001392029879385965 +gfx950,256,640,1536,5120,triton,0,0,17.735612,triton,567.58,739.03,0.0001312255859375 +gfx950,256,672,1536,5120,triton,0,0,23.492412,triton,449.92,569.09,0.00015307229662698413 +gfx950,256,704,1536,5120,triton,0,0,22.962256,triton,482.22,593.65,0.0001581365411931818 +gfx950,256,736,1536,5120,triton,0,0,23.669684,triton,489.08,586.98,0.00015479930932971015 +gfx950,256,768,1536,5120,triton,0,0,24.230237,triton,498.53,584.22,0.00013309054904513888 +gfx950,256,800,1536,5120,triton,0,0,23.230089,triton,541.66,620.66,0.00015462239583333333 +gfx950,256,832,1536,5120,triton,0,0,23.9287,triton,546.88,613.49,0.00014632787459935897 +gfx950,256,864,1536,5120,triton,0,0,24.561654,triton,553.28,608.36,0.00014995056905864196 +gfx950,256,896,1536,5120,triton,0,0,25.062212,triton,562.32,606.66,0.00013514927455357144 +gfx950,256,928,1536,5120,triton,0,0,24.176175,triton,603.74,639.74,0.00013259361530172415 +gfx950,256,960,1536,5120,triton,0,0,24.528026,triton,615.6,641.25,0.00015326605902777778 +gfx950,256,992,1536,5120,triton,0,0,25.11925,triton,621.15,636.59,0.00014832199260752689 +gfx950,256,1024,1536,5120,triton,0,0,25.451402,triton,632.82,638.59,0.0001621246337890625 +gfx950,256,1088,1536,5120,triton,0,0,25.025205,triton,683.82,670.41,0.00014600566789215686 +gfx950,256,1152,1536,5120,triton,0,0,26.572723,triton,681.88,651.1,0.00012828685619212962 +gfx950,256,1216,1536,5120,triton,0,0,26.054543,triton,734.08,684.17,0.00013866759183114036 +gfx950,256,1280,1536,5120,triton,0,0,27.306035,triton,737.3,672.02,0.00015004475911458334 +gfx950,256,1344,1536,5120,triton,0,0,29.799707,triton,709.38,633.37,0.0001506502666170635 +gfx950,256,1408,1536,5120,triton,0,0,35.163728,triton,629.79,551.67,0.00013917865175189394 +gfx950,256,1472,1536,5120,triton,0,0,33.411035,triton,692.96,596.3,0.00013931937839673913 +gfx950,256,1536,1536,5120,triton,0,0,35.341244,triton,683.6,578.57,0.00014241536458333334 +gfx950,256,1600,1536,5120,triton,0,0,33.814353,triton,744.23,620.2,0.00014241536458333334 +gfx950,256,1664,1536,5120,triton,0,0,35.876512,triton,729.52,599.16,0.00014554537259615384 +gfx950,256,1728,1536,5120,triton,0,0,34.324284,triton,791.83,641.53,0.00014580620659722222 +gfx950,256,1792,1536,5120,triton,0,0,36.257034,triton,777.39,621.79,0.000152587890625 +gfx950,256,1856,1536,5120,triton,0,0,34.835659,triton,838.0,662.21,0.00014732623922413793 +gfx950,256,1920,1536,5120,triton,0,0,37.588286,triton,803.41,627.67,0.0001437717013888889 +gfx950,256,1984,1536,5120,triton,0,0,35.162525,triton,887.47,685.88,0.00013618059055779569 +gfx950,256,2048,1536,5120,triton,0,0,36.40567,triton,884.81,676.86,0.0001452763875325521 +gfx950,256,4096,1536,5120,triton,0,0,60.32591,triton,1067.94,686.58,0.00014130274454752603 +gfx950,256,6144,1536,5120,triton,0,0,94.91699,triton,1018.12,613.12,0.00013838873969184028 +gfx950,256,8192,1536,5120,triton,0,0,104.36907,triton,1234.55,718.35,0.00014813741048177084 +gfx950,256,10240,1536,5120,triton,0,0,123.05925,triton,1308.81,745.58,0.00014902750651041668 +gfx950,256,12288,1536,5120,triton,0,0,155.776061,triton,1240.71,696.69,0.00013913048638237847 +gfx950,256,14336,1536,5120,triton,0,0,176.25653,triton,1279.3,710.92,0.00014772869291759672 +gfx950,256,16384,1536,5120,triton,0,0,195.08902,triton,1320.93,728.29,0.00014026959737141928 +gfx950,256,32768,1536,5120,triton,0,0,371.780687,triton,1386.29,743.18,0.0001408259073893229 +gfx950,256,1,512,5120,triton,0,0,4.122704,triton,1.27,637.34,0.0 +gfx950,256,2,512,5120,triton,0,0,4.086927,triton,2.57,644.43,0.0 +gfx950,256,4,512,5120,triton,0,0,4.58342,triton,4.58,577.3,0.0 +gfx950,256,8,512,5120,triton,0,0,4.765505,triton,8.8,560.4,0.0 +gfx950,256,16,512,5120,triton,0,0,4.962032,triton,16.91,548.11,0.0001220703125 +gfx950,256,32,512,5120,triton,0,0,4.694561,triton,35.74,600.28,0.0001220703125 +gfx950,256,48,512,5120,triton,0,0,4.888113,triton,51.48,596.62,0.0 +gfx950,256,64,512,5120,triton,0,0,5.06823,triton,66.21,594.81,0.0001220703125 +gfx950,256,80,512,5120,triton,0,0,5.273788,triton,79.53,590.27,0.0001220703125 +gfx950,256,96,512,5120,triton,0,0,5.37263,triton,93.68,597.71,0.00014241536458333334 +gfx950,256,112,512,5120,triton,0,0,6.23301,triton,94.21,530.97,0.00013950892857142856 +gfx950,256,128,512,5120,triton,0,0,6.24304,triton,107.49,545.87,0.0001373291015625 +gfx950,256,144,512,5120,triton,0,0,6.303571,triton,119.77,556.22,0.00023057725694444444 +gfx950,256,160,512,5120,triton,0,0,6.238645,triton,134.46,577.77,0.00010986328125 +gfx950,256,176,512,5120,triton,0,0,6.421811,triton,143.69,576.59,0.0001553622159090909 +gfx950,256,192,512,5120,triton,0,0,6.548162,triton,153.73,580.48,0.00014241536458333334 +gfx950,256,208,512,5120,triton,0,0,7.194875,triton,151.57,541.97,9.390024038461539e-05 +gfx950,256,224,512,5120,triton,0,0,7.335347,triton,160.1,544.99,8.719308035714286e-05 +gfx950,256,240,512,5120,triton,0,0,7.77348,triton,161.87,526.92,0.00015462239583333333 +gfx950,256,256,512,5120,triton,0,0,7.863949,triton,170.67,533.36,0.00012969970703125 +gfx950,256,288,512,5120,triton,0,0,8.481879,triton,178.02,517.68,0.00019666883680555556 +gfx950,256,320,512,5120,triton,0,0,8.629165,triton,194.42,531.63,0.000128173828125 +gfx950,256,352,512,5120,triton,0,0,8.904796,triton,207.25,537.25,0.00011652166193181819 +gfx950,256,384,512,5120,triton,0,0,9.09199,triton,221.43,547.82,0.00015767415364583334 +gfx950,256,416,512,5120,triton,0,0,9.591394,triton,227.4,539.79,0.00021597055288461538 +gfx950,256,448,512,5120,triton,0,0,9.921979,triton,236.73,541.62,0.00013078962053571428 +gfx950,256,480,512,5120,triton,0,0,9.656859,triton,260.6,576.85,0.00011800130208333333 +gfx950,256,512,512,5120,triton,0,0,9.759887,triton,275.04,590.91,0.000164031982421875 +gfx950,256,544,512,5120,triton,0,0,10.910642,triton,261.41,546.6,0.00012566061580882352 +gfx950,256,576,512,5120,triton,0,0,11.258286,triton,268.24,547.19,0.00011528862847222222 +gfx950,256,608,512,5120,triton,0,0,11.441175,triton,278.61,555.62,0.0001670435855263158 +gfx950,256,640,512,5120,triton,0,0,11.669424,triton,287.54,561.6,0.0001495361328125 +gfx950,256,672,512,5120,triton,0,0,12.365162,triton,284.93,545.91,0.00015694754464285715 +gfx950,256,704,512,5120,triton,0,0,12.769906,triton,289.04,544.0,0.0001553622159090909 +gfx950,256,736,512,5120,triton,0,0,12.701343,triton,303.81,562.41,0.00013799252717391303 +gfx950,256,768,512,5120,triton,0,0,12.44915,triton,323.44,589.6,0.00015767415364583334 +gfx950,256,800,512,5120,triton,0,0,13.043337,triton,321.57,577.82,0.000166015625 +gfx950,256,832,512,5120,triton,0,0,13.476663,triton,323.68,573.83,0.00017136793870192306 +gfx950,256,864,512,5120,triton,0,0,13.404798,triton,337.93,591.57,0.00011980975115740741 +gfx950,256,896,512,5120,triton,0,0,13.85629,triton,339.02,586.48,0.00011335100446428572 +gfx950,256,928,512,5120,triton,0,0,13.609232,triton,357.51,611.58,0.0001599542025862069 +gfx950,256,960,512,5120,triton,0,0,14.123247,triton,356.37,603.24,0.00012613932291666668 +gfx950,256,992,512,5120,triton,0,0,13.984268,triton,371.91,623.29,0.00017522996471774194 +gfx950,256,1024,512,5120,triton,0,0,14.222406,triton,377.48,626.68,0.0001201629638671875 +gfx950,256,1088,512,5120,triton,0,0,18.33401,triton,311.13,507.59,0.00014899758731617648 +gfx950,256,1152,512,5120,triton,0,0,17.865286,triton,338.07,542.91,0.00014750162760416666 +gfx950,256,1216,512,5120,triton,0,0,18.97994,triton,335.9,531.75,0.00014616313733552632 +gfx950,256,1280,512,5120,triton,0,0,17.71868,triton,378.75,591.79,0.00014190673828125 +gfx950,256,1344,512,5120,triton,0,0,18.857102,triton,373.68,576.92,0.00013950892857142856 +gfx950,256,1408,512,5120,triton,0,0,18.18392,triton,405.96,619.9,0.0001470392400568182 +gfx950,256,1472,512,5120,triton,0,0,19.082172,triton,404.44,611.32,0.00013533882472826087 +gfx950,256,1536,512,5120,triton,0,0,17.9312,triton,449.11,672.49,0.00013097127278645834 +gfx950,256,1600,512,5120,triton,0,0,19.316023,triton,434.28,644.64,0.000162353515625 +gfx950,256,1664,512,5120,triton,0,0,18.44074,triton,473.09,696.56,0.00015845665564903847 +gfx950,256,1728,512,5120,triton,0,0,20.019173,triton,452.55,661.28,0.00014693648726851852 +gfx950,256,1792,512,5120,triton,0,0,17.873235,triton,525.66,762.68,0.00012969970703125 +gfx950,256,1856,512,5120,triton,0,0,19.945,triton,487.88,703.17,0.00014206458782327585 +gfx950,256,1920,512,5120,triton,0,0,18.328385,triton,549.22,786.64,0.00013834635416666666 +gfx950,256,1984,512,5120,triton,0,0,20.511124,triton,507.13,722.1,0.00014668126260080646 +gfx950,256,2048,512,5120,triton,0,0,18.408175,triton,583.3,825.96,0.00012969970703125 +gfx950,256,4096,512,5120,triton,0,0,29.080859,triton,738.45,955.52,0.00015544891357421875 +gfx950,256,6144,512,5120,triton,0,0,36.399731,triton,884.96,1109.08,0.0001452763875325521 +gfx950,256,8192,512,5120,triton,0,0,39.52373,triton,1086.68,1339.78,0.00014495849609375 +gfx950,256,10240,512,5120,triton,0,0,62.609835,triton,857.49,1046.74,0.00013599395751953124 +gfx950,256,12288,512,5120,triton,0,0,66.349347,triton,970.99,1177.39,0.00013637542724609375 +gfx950,256,14336,512,5120,triton,0,0,70.669889,triton,1063.56,1283.46,0.00014019012451171875 +gfx950,256,16384,512,5120,triton,0,0,71.525263,triton,1200.97,1444.03,0.00016033649444580078 +gfx950,256,32768,512,5120,triton,0,0,138.14818,triton,1243.58,1476.3,0.00013780593872070312 diff --git a/aiter/jit/core.py b/aiter/jit/core.py index bd023fd0a8a..cb37dbb3381 100644 --- a/aiter/jit/core.py +++ b/aiter/jit/core.py @@ -125,6 +125,13 @@ def mp_lock( f"{AITER_ROOT_DIR}/aiter/configs/a8w8_blockscale_tuned_gemm.csv", ) +# Native E8M0 group32 scales have a different operand contract from the +# FP32 128x128 blockscale family, so shape-identical rows must stay separate. +AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_GROUP32 = os.getenv( + "AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_GROUP32", + f"{AITER_ROOT_DIR}/aiter/configs/a8w8_blockscale_group32_tuned_gemm.csv", +) + AITER_CONFIG_FMOE = os.getenv( "AITER_CONFIG_FMOE", f"{AITER_ROOT_DIR}/aiter/configs/tuned_fmoe.csv", @@ -243,6 +250,14 @@ def AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_FILE(self): "a8w8_blockscale_tuned_gemm", ) + @property + def AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_GROUP32_FILE(self): + return self.get_config_file( + "AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_GROUP32", + AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_GROUP32, + "a8w8_blockscale_group32_tuned_gemm", + ) + @property def AITER_CONFIG_FMOE_FILE(self): return self.get_config_file( diff --git a/aiter/ops/gemm_op_a8w8.py b/aiter/ops/gemm_op_a8w8.py index 05b5264d9f7..40c3ad13967 100644 --- a/aiter/ops/gemm_op_a8w8.py +++ b/aiter/ops/gemm_op_a8w8.py @@ -803,8 +803,28 @@ def _blockscale_triton( x_scale: Tensor, w_scale: Tensor, dtype: torch.dtype, + *, + group32: bool = False, + split_k: int | None = None, ) -> Tensor: - """Run the triton kernel on CK's inputs: row-major x_scale, (N, K) weight, JIT per arch.""" + """Run Triton on unshuffled weights with the selected scale format.""" + if group32: + from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale_group32 import ( + gemm_a8w8_blockscale_group32, + ) + + assert WQ.ndim == w_scale.ndim == 2, "Expected matrix weights and scales" + group_n = 32 if w_scale.shape[0] == -(-WQ.shape[0] // 32) else 1 + return gemm_a8w8_blockscale_group32( + XQ, + WQ, + x_scale, + w_scale, + dtype=dtype, + weight_group_rows=group_n, + split_k=split_k, + ) + from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale import ( gemm_a8w8_blockscale as _gemm_a8w8_blockscale_triton, ) @@ -821,6 +841,7 @@ def gemm_a8w8_blockscale_fake( w_scale: Tensor, dtype: torch.dtype = dtypes.bf16, isBpreshuffled=False, + split_k: int | None = None, ) -> torch.Tensor: m = XQ.shape[0] n = WQ.shape[0] @@ -828,7 +849,7 @@ def gemm_a8w8_blockscale_fake( return Y -@torch_compile_guard(gen_fake=gemm_a8w8_blockscale_fake) +@torch_compile_guard(mutates_args=[], gen_fake=gemm_a8w8_blockscale_fake) def gemm_a8w8_blockscale( XQ: Tensor, WQ: Tensor, @@ -836,62 +857,107 @@ def gemm_a8w8_blockscale( w_scale: Tensor, dtype: torch.dtype = dtypes.bf16, isBpreshuffled: bool = False, + split_k: int | None = None, ) -> torch.Tensor: - assert dtype in [ - dtypes.bf16, - dtypes.fp16, - ], f"Output {dtype=} is currently not supported in gemm_a8w8" + """Blockscaled A8W8 GEMM with configuration-first backend dispatch. + + Native E8M0 group32 scales (typed tensors or uint8 views) use + AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_GROUP32; + FP32 128x128 scales use AITER_CONFIG_GEMM_A8W8_BLOCKSCALE. Both tables are + queried through get_CKGEMM_config before choosing a fallback. Group32 + currently supports libtype="triton", also its default on a config miss. + Triton tile and split-K parameters come from its own tuning tables. + For native group32 operands, split_k optionally overrides the positive + partition count without changing configured backend selection. + """ + is_group32 = ( + x_scale.dtype in (dtypes.fp8_e8m0, torch.uint8) + and w_scale.dtype in (dtypes.fp8_e8m0, torch.uint8) + and XQ.ndim == WQ.ndim == 2 + and x_scale.shape == (XQ.shape[0], XQ.shape[1] // 32) + and w_scale.shape + in ( + (WQ.shape[0], WQ.shape[1] // 32), + (-(-WQ.shape[0] // 32), WQ.shape[1] // 32), + ) + ) + # A malformed byte-scale layout must not fall through to FP32 CK dispatch. + assert ( + is_group32 or x_scale.dtype == w_scale.dtype == dtypes.fp32 + ), "Expected E8M0 group32 scale shapes (typed or uint8), or FP32 128x128 scales" + assert ( + split_k is None or is_group32 + ), "split_k override requires native group32 operands" + assert dtype in (dtypes.bf16, dtypes.fp16) or ( + is_group32 and dtype == dtypes.fp32 + ), f"Output {dtype=} is currently not supported in gemm_a8w8" m = XQ.shape[0] n = WQ.shape[0] k = XQ.shape[1] - Y = torch.empty(m, n, dtype=dtype, device=XQ.device) if isBpreshuffled: + assert not is_group32, "Group32 FP8 GEMM requires native weights" if get_gfx() in ["gfx950"] and m >= 16 and k >= 512 and dtype == dtypes.bf16: + Y = torch.empty(m, n, dtype=dtype, device=XQ.device) return gfx950_a8w8_blockscale_ASM(XQ, WQ, x_scale, w_scale, Y) else: assert 0, "asm kernel only support B preshuffle and m >= 16" - else: - if not _hip_blockscale_supported(): - return _blockscale_triton(XQ, WQ, x_scale, w_scale, dtype) - config = get_CKGEMM_config( - m, n, k, AITER_CONFIGS.AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_FILE + + tuned_file = ( + AITER_CONFIGS.AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_GROUP32_FILE + if is_group32 + else AITER_CONFIGS.AITER_CONFIG_GEMM_A8W8_BLOCKSCALE_FILE + ) + config = get_CKGEMM_config(m, n, k, tuned_file) + if config is not None: + libtype = config["libtype"] + if libtype == "triton": + return _blockscale_triton( + XQ, WQ, x_scale, w_scale, dtype, group32=is_group32, split_k=split_k + ) + # CK/CKTile currently consume FP32 128x128 scales. A misconfigured + # group32 row must fail instead of reinterpreting its E8M0 bytes. + assert not is_group32, f"Unsupported libtype {libtype} for group32 GEMM" + splitK = int(config.get("splitK", 0)) + kernelName = str(config.get("kernelName", "")) + Y = torch.empty(m, n, dtype=dtype, device=XQ.device) + if libtype == "ck": + return gemm_a8w8_blockscale_ck( + XQ, + WQ, + x_scale, + w_scale, + Y, + splitK=splitK, + kernelName=kernelName, + ) + elif libtype == "cktile": + return gemm_a8w8_blockscale_cktile( + XQ, + WQ, + x_scale, + w_scale, + Y, + splitK=splitK, + kernelName=kernelName, + ) + else: + assert 0, f"Unsupported libtype {libtype} for gemm_a8w8_blockscale" + + if is_group32 or not _hip_blockscale_supported(): + return _blockscale_triton( + XQ, WQ, x_scale, w_scale, dtype, group32=is_group32, split_k=split_k ) - if config is not None: - libtype = config["libtype"] - splitK = int(config.get("splitK", 0)) - kernelName = str(config.get("kernelName", "")) - if libtype == "ck": - return gemm_a8w8_blockscale_ck( - XQ, - WQ, - x_scale, - w_scale, - Y, - splitK=splitK, - kernelName=kernelName, - ) - elif libtype == "cktile": - return gemm_a8w8_blockscale_cktile( - XQ, - WQ, - x_scale, - w_scale, - Y, - splitK=splitK, - kernelName=kernelName, - ) - else: - assert 0, f"Unsupported libtype {libtype} for gemm_a8w8_blockscale" - min_m = _BLOCKSCALE_TRITON_FALLBACK_MIN_M.get(get_gfx()) - if min_m is not None and m >= min_m: - return _blockscale_triton(XQ, WQ, x_scale, w_scale, dtype) - try: - return gemm_a8w8_blockscale_ck(XQ, WQ, x_scale, w_scale, Y) - except RuntimeError as e: - raise RuntimeError( - f"gemm_a8w8_blockscale failed for shape M={m}, N={n}, K={k}, " - f"{dtype=}, config={config}: {e}" - ) from e + min_m = _BLOCKSCALE_TRITON_FALLBACK_MIN_M.get(get_gfx()) + if min_m is not None and m >= min_m: + return _blockscale_triton(XQ, WQ, x_scale, w_scale, dtype) + Y = torch.empty(m, n, dtype=dtype, device=XQ.device) + try: + return gemm_a8w8_blockscale_ck(XQ, WQ, x_scale, w_scale, Y) + except RuntimeError as e: + raise RuntimeError( + f"gemm_a8w8_blockscale failed for shape M={m}, N={n}, K={k}, " + f"{dtype=}, config={config}: {e}" + ) from e def flatmm_a8w8_blockscale_ASM( diff --git a/aiter/ops/triton/_triton_kernels/gemm/basic/gemm_a8w8_blockscale_group32.py b/aiter/ops/triton/_triton_kernels/gemm/basic/gemm_a8w8_blockscale_group32.py new file mode 100644 index 00000000000..ecff2cbc0cb --- /dev/null +++ b/aiter/ops/triton/_triton_kernels/gemm/basic/gemm_a8w8_blockscale_group32.py @@ -0,0 +1,295 @@ +# SPDX-License-Identifier: MIT +# Copyright (C) 2026, Advanced Micro Devices, Inc. All rights reserved. +"""Native E4M3/E8M0 group32 GEMMs, including packed-K small-M reductions.""" + +import triton +import triton.language as tl + +from aiter.ops.triton.utils._triton.kernel_repr import make_kernel_repr +from aiter.ops.triton.utils.gemm_config_utils import get_gemm_config + +_gemm_group32_repr = make_kernel_repr( + "_gemm_a8w8_blockscale_group32_kernel", + [ + "BLOCK_SIZE_M", + "BLOCK_SIZE_N", + "BLOCK_SIZE_K", + "GROUP_N", + "SPLITK_BLOCK_SIZE", + "FUSED_SPLITS", + "N_FIRST", + "B_CACHE_MODIFIER", + "N", + "K", + "LAUNCH_OPTIONS", + ], +) +_gemm_group32_packed_repr = make_kernel_repr( + "_gemm_a8w8_blockscale_group32_packed_kernel", + [ + "BLOCK_SIZE_M", + "BLOCK_SIZE_N", + "BLOCK_SIZE_K", + "K_PACK", + "SPLITK_BLOCK_SIZE", + "FUSED_SPLITS", + "B_CACHE_MODIFIER", + "N", + "K", + "LAUNCH_OPTIONS", + ], +) + + +@triton.jit +def _tile_on_xcd(N: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, SPLITS: tl.constexpr): + """(pid_m, pid_n, tile, split) of a 1D grid, all splits of a tile on one XCD. + + CTAs go to XCDs by pid % 8, a multiple of every gfx950 XCD count. + """ + pid = tl.program_id(0) + tile = (pid // 8 // SPLITS) * 8 + pid % 8 + grid_n: tl.constexpr = (N + BLOCK_SIZE_N - 1) // BLOCK_SIZE_N + return tile // grid_n, tile % grid_n, tile, pid // 8 % SPLITS + + +@triton.jit +def _sum_splits_on_xcd( + partial, c_ptrs, mask, slot, counter, split, SPLITS: tl.constexpr +): + """The tile's last CTA to arrive sums its splits in split order. + + The splits share one XCD's L2, so a completed store is visible without an + agent-scope release and its L2 writeback. + """ + BM: tl.constexpr = partial.shape[0] + BN: tl.constexpr = partial.shape[1] + local = tl.arange(0, BM)[:, None] * BN + tl.arange(0, BN)[None, :] + tl.store(slot + split * BM * BN + local, partial) + # debug_barrier is a bare s_barrier: drain this wave's stores first. + tl.inline_asm_elementwise( + "s_waitcnt vmcnt(0)", "=v,v", [split], dtype=tl.int32, is_pure=False, pack=1 + ) + tl.debug_barrier() + if tl.atomic_add(counter, 1, sem="acq_rel", scope="cta") == SPLITS - 1: + total = tl.zeros((BM, BN), tl.float32) + for s in tl.static_range(SPLITS): + # .cv skips this CU's L1, which never saw the other splits. + total += tl.load(slot + s * BM * BN + local, cache_modifier=".cv") + tl.store(c_ptrs, total.to(c_ptrs.dtype.element_ty), mask) + tl.store(counter, 0) + + +@triton.jit(repr=_gemm_group32_repr, do_not_specialize=["M"]) +def _gemm_a8w8_blockscale_group32_kernel( + a_ptr, + b_ptr, + a_scale_ptr, + b_scale_ptr, + c_ptr, + ws_ptr, + cnt_ptr, + # Runtime, not constexpr: M is the token count, and specializing on it + # recompiles the kernel for every prefill chunk length. + M, + N: tl.constexpr, + K: tl.constexpr, + GROUP_N: tl.constexpr, + SPLITK_BLOCK_SIZE: tl.constexpr, + BLOCK_SIZE_M: tl.constexpr, + BLOCK_SIZE_N: tl.constexpr, + BLOCK_SIZE_K: tl.constexpr, + LAUNCH_OPTIONS: tl.constexpr, + N_FIRST: tl.constexpr = False, + FUSED_SPLITS: tl.constexpr = 1, + B_CACHE_MODIFIER: tl.constexpr = None, +): + """E4M3 x E4M3 with E8M0 group scales, on the CDNA4 microscaling MFMA. + + A tile spans BLOCK_SIZE_K/32 scale groups. Below a K tile of 64, + dot_scaled lowers to slower BF16 emulation instead of microscaling MFMA. + FUSED_SPLITS > 1 sums split-K partials in this launch on a 1D grid; + otherwise c_ptr takes one FP32 partial per program_id(2). + B_CACHE_MODIFIER applies to the weight and weight-scale loads. + """ + tl.static_assert(BLOCK_SIZE_K >= 64 and BLOCK_SIZE_K % 32 == 0) + if FUSED_SPLITS > 1: + pid_m, pid_n, tile, split = _tile_on_xcd(N, BLOCK_SIZE_N, FUSED_SPLITS) + if pid_m * BLOCK_SIZE_M >= M: + return + else: + # The first grid dimension advances fastest. N-first traversal reuses + # A; M-first traversal reuses B. The wrapper swaps the launch dimensions. + pid_m = tl.program_id(1 if N_FIRST else 0) + pid_n = tl.program_id(0 if N_FIRST else 1) + split = tl.program_id(2) + row = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + col = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) + groups: tl.constexpr = K // 32 + ks = split * SPLITK_BLOCK_SIZE + tl.arange(0, BLOCK_SIZE_K) + gs = split * SPLITK_BLOCK_SIZE // 32 + tl.arange(0, BLOCK_SIZE_K // 32) + rows = row[:, None] < M + cols = col[:, None] < N + a_ptrs = a_ptr + row[:, None] * K + ks[None, :] + b_ptrs = b_ptr + col[:, None] * K + ks[None, :] + as_ptrs = a_scale_ptr + row[:, None] * groups + gs[None, :] + # One scale row per GROUP_N output columns, so the column index is + # divided rather than the grid expanded. + bs_ptrs = b_scale_ptr + (col[:, None] // GROUP_N) * groups + gs[None, :] + + accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), tl.float32) + # A constexpr trip count lets the pipeliner schedule statically. Only a + # last split with fewer K tiles than the others stops early, at K. + LAST: tl.constexpr = K - (K - 1) // SPLITK_BLOCK_SIZE * SPLITK_BLOCK_SIZE + if (LAST + BLOCK_SIZE_K - 1) // BLOCK_SIZE_K == SPLITK_BLOCK_SIZE // BLOCK_SIZE_K: + span = SPLITK_BLOCK_SIZE + else: + span = tl.minimum(SPLITK_BLOCK_SIZE, K - split * SPLITK_BLOCK_SIZE) + for i in range(0, span, BLOCK_SIZE_K): + # K masks only when K has a partial tile. + a_mask, b_mask, as_mask, bs_mask = rows, cols, rows, cols + if K % BLOCK_SIZE_K != 0: + live = (ks + i < K)[None, :] + held = (gs + i // 32 < groups)[None, :] + a_mask, b_mask = rows & live, cols & live + as_mask, bs_mask = rows & held, cols & held + a = tl.load(a_ptrs, a_mask, other=0.0) + b = tl.load(b_ptrs, b_mask, other=0.0, cache_modifier=B_CACHE_MODIFIER) + a_code = tl.load(as_ptrs, as_mask, other=127) + b_code = tl.load(bs_ptrs, bs_mask, other=127, cache_modifier=B_CACHE_MODIFIER) + # acc= leaves the sum in the matrix core's registers: one rounding per + # tile instead of two, and no separate vector add. + accumulator = tl.dot_scaled( + a, a_code, "e4m3", b.T, b_code, "e4m3", acc=accumulator + ) + a_ptrs += BLOCK_SIZE_K + b_ptrs += BLOCK_SIZE_K + as_ptrs += BLOCK_SIZE_K // 32 + bs_ptrs += BLOCK_SIZE_K // 32 + c_ptrs = c_ptr + row[:, None] * N + col[None, :] + if FUSED_SPLITS > 1: + _sum_splits_on_xcd( + accumulator, + c_ptrs, + rows & cols.T, + ws_ptr + tile * (FUSED_SPLITS * BLOCK_SIZE_M * BLOCK_SIZE_N), + cnt_ptr + tile, + split, + FUSED_SPLITS, + ) + else: + tl.store(c_ptrs + split * M * N, accumulator, rows & cols.T) + + +@triton.jit(repr=_gemm_group32_packed_repr, do_not_specialize=["M"]) +def _gemm_a8w8_blockscale_group32_packed_kernel( + a_ptr, + b_ptr, + a_scale_ptr, + b_scale_ptr, + c_ptr, + ws_ptr, + cnt_ptr, + M, + N: tl.constexpr, + K: tl.constexpr, + BLOCK_SIZE_M: tl.constexpr, + BLOCK_SIZE_N: tl.constexpr, + BLOCK_SIZE_K: tl.constexpr, + K_PACK: tl.constexpr, + LAUNCH_OPTIONS: tl.constexpr, + SPLITK_BLOCK_SIZE: tl.constexpr = 0, + FUSED_SPLITS: tl.constexpr = 1, + B_CACHE_MODIFIER: tl.constexpr = None, +): + """Small-M group32 GEMM with K panels packed into MFMA rows/columns. + + Each packed row/column retains its own E8M0 scales. Only matching panel + pairs contribute to the output; cross-panel products are discarded. One + CTA owns the K reduction, or with FUSED_SPLITS > 1 one SPLITK_BLOCK_SIZE + of it, summed in this launch. BLOCK_SIZE_M is an unpacked token tile, + independent of runtime M. B_CACHE_MODIFIER applies to the weight and + weight-scale loads. + """ + tl.static_assert(K_PACK == 1 or K_PACK == 2 or K_PACK == 4) + tl.static_assert(BLOCK_SIZE_M * K_PACK >= 16) + tl.static_assert(BLOCK_SIZE_K >= 128 and BLOCK_SIZE_K % 32 == 0) + STEP: tl.constexpr = BLOCK_SIZE_K * K_PACK + if FUSED_SPLITS > 1: + tl.static_assert(SPLITK_BLOCK_SIZE % STEP == 0) + pid_m, pid_n, tile, split = _tile_on_xcd(N, BLOCK_SIZE_N, FUSED_SPLITS) + if pid_m * BLOCK_SIZE_M >= M: + return + start = split * SPLITK_BLOCK_SIZE + stop = tl.minimum(start + SPLITK_BLOCK_SIZE, K) + else: + pid_m = tl.program_id(0) + pid_n = tl.program_id(1) + start = 0 + stop = K + rows = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M * K_PACK) // K_PACK + a_panel = tl.arange(0, BLOCK_SIZE_M * K_PACK) % K_PACK + cols = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N * K_PACK) // K_PACK + b_panel = tl.arange(0, BLOCK_SIZE_N * K_PACK) % K_PACK + ks = tl.arange(0, BLOCK_SIZE_K) + gs = tl.arange(0, BLOCK_SIZE_K // 32) + groups: tl.constexpr = K // 32 + accumulator = tl.zeros((BLOCK_SIZE_M * K_PACK, BLOCK_SIZE_N * K_PACK), tl.float32) + for base in range(start, stop, STEP): + ak = base + a_panel[:, None] * BLOCK_SIZE_K + ks[None, :] + bk = base + b_panel[:, None] * BLOCK_SIZE_K + ks[None, :] + a = tl.load( + a_ptr + rows[:, None] * K + ak, + (rows[:, None] < M) & (ak < K), + other=0.0, + ) + b = tl.load( + b_ptr + cols[:, None] * K + bk, + (cols[:, None] < N) & (bk < K), + other=0.0, + cache_modifier=B_CACHE_MODIFIER, + ) + ag = base // 32 + a_panel[:, None] * (BLOCK_SIZE_K // 32) + gs[None, :] + bg = base // 32 + b_panel[:, None] * (BLOCK_SIZE_K // 32) + gs[None, :] + a_code = tl.load( + a_scale_ptr + rows[:, None] * groups + ag, + (rows[:, None] < M) & (ag < groups), + other=127, + ) + b_code = tl.load( + b_scale_ptr + (cols[:, None] // 32) * groups + bg, + (cols[:, None] < N) & (bg < groups), + other=127, + cache_modifier=B_CACHE_MODIFIER, + ) + accumulator = tl.dot_scaled( + a, a_code, "e4m3", b.T, b_code, "e4m3", acc=accumulator + ) + panels = accumulator.reshape(BLOCK_SIZE_M, K_PACK, BLOCK_SIZE_N, K_PACK).trans( + 0, 2, 1, 3 + ) + pair = tl.arange(0, K_PACK) + diagonal = tl.where( + pair[None, None, :, None] == pair[None, None, None, :], panels, 0.0 + ) + output = tl.sum(tl.sum(diagonal, 3), 2) + row = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + col = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) + c_ptrs = c_ptr + row[:, None] * N + col[None, :] + mask = (row[:, None] < M) & (col[None, :] < N) + if FUSED_SPLITS > 1: + _sum_splits_on_xcd( + output, + c_ptrs, + mask, + ws_ptr + tile * (FUSED_SPLITS * BLOCK_SIZE_M * BLOCK_SIZE_N), + cnt_ptr + tile, + split, + FUSED_SPLITS, + ) + else: + tl.store(c_ptrs, output, mask) + + +def _get_config(M: int, N: int, K: int): + return get_gemm_config("GEMM-A8W8_BLOCKSCALE_GROUP32", M, N, K) diff --git a/aiter/ops/triton/configs/CLAUDE.md b/aiter/ops/triton/configs/CLAUDE.md index 985e8185238..119aa5813c1 100644 --- a/aiter/ops/triton/configs/CLAUDE.md +++ b/aiter/ops/triton/configs/CLAUDE.md @@ -191,7 +191,9 @@ The returned config is a fresh deep copy, safe to mutate. - `M_LEQ_x` is searched over `STANDARD_M_BOUNDS = (1, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096, 8192)`. A caller may override with `bounds=(...)`, which must be strictly increasing positive - ints. + ints. A file may instead declare a top-level `M_BOUNDS` list to keep + nonstandard tuning boundaries in JSON. Explicit `bounds=` wins; files + without `M_BOUNDS` retain the standard bounds. - `any` must exist unless every reachable `M` is covered by an explicit bound. A `KeyError` at lookup time usually means it is missing. - The deprecated `{"large": …, "small": …}` shape must not be introduced. diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/DEFAULT.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/DEFAULT.json new file mode 100644 index 00000000000..cb5a173e362 --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/DEFAULT.json @@ -0,0 +1,16 @@ +{ + "any": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1152-K=5120.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1152-K=5120.json new file mode 100644 index 00000000000..7fa212b7fea --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1152-K=5120.json @@ -0,0 +1,1183 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "NUM_KSPLIT": 3, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 8, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1536-K=5120.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1536-K=5120.json new file mode 100644 index 00000000000..4d6b29064fd --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1536-K=5120.json @@ -0,0 +1,1179 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=16384-K=1280.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=16384-K=1280.json new file mode 100644 index 00000000000..a6432a621d8 --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=16384-K=1280.json @@ -0,0 +1,1122 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "K_PACK": 4, + "NUM_KSPLIT": 1, + "num_warps": 4, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 128, + "K_PACK": 4, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 128, + "K_PACK": 4, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "NUM_KSPLIT": 1, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1792-K=5120.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1792-K=5120.json new file mode 100644 index 00000000000..e5c77e684f8 --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=1792-K=5120.json @@ -0,0 +1,1131 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg" + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=2304-K=5120.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=2304-K=5120.json new file mode 100644 index 00000000000..9d7f68c75ef --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=2304-K=5120.json @@ -0,0 +1,1137 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 8, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 8, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 8, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=25600-K=6144.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=25600-K=6144.json new file mode 100644 index 00000000000..43c1029423c --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=25600-K=6144.json @@ -0,0 +1,1102 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 3, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 3, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 3, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=32768-K=1280.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=32768-K=1280.json new file mode 100644 index 00000000000..06be301e94e --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=32768-K=1280.json @@ -0,0 +1,1106 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=4096-K=1280.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=4096-K=1280.json new file mode 100644 index 00000000000..2f1e7293c8d --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=4096-K=1280.json @@ -0,0 +1,1167 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "num_warps": 4, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "num_warps": 4, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 1, + "num_warps": 4, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 2, + "NUM_KSPLIT": 1, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=4608-K=5120.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=4608-K=5120.json new file mode 100644 index 00000000000..ad927a532d3 --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=4608-K=5120.json @@ -0,0 +1,1102 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=512-K=5120.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=512-K=5120.json new file mode 100644 index 00000000000..484be8ce21d --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=512-K=5120.json @@ -0,0 +1,1261 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "NUM_KSPLIT": 5, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 10, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 5, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 6, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=1152.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=1152.json new file mode 100644 index 00000000000..4120b736633 --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=1152.json @@ -0,0 +1,1135 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "K_PACK": 4, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "K_PACK": 4, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "K_PACK": 4, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=15360.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=15360.json new file mode 100644 index 00000000000..ffcd2c6bb6b --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=15360.json @@ -0,0 +1,1117 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 15, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 1, + "REDUCE_BLOCK_SIZE_N": 128, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 8, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 15, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 1, + "REDUCE_BLOCK_SIZE_N": 128, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 8, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 15, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 1, + "REDUCE_BLOCK_SIZE_N": 128, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 8, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "NUM_KSPLIT": 3, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 4, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 128, + "FUSED_SPLITK": false + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 128, + "FUSED_SPLITK": false + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 128, + "FUSED_SPLITK": false + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 128, + "FUSED_SPLITK": false + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 128, + "FUSED_SPLITK": false + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 128, + "FUSED_SPLITK": false + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 128, + "FUSED_SPLITK": false + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=2048.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=2048.json new file mode 100644 index 00000000000..048b23d1e12 --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=2048.json @@ -0,0 +1,1097 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 4, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "K_PACK": 4, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg" + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=2304.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=2304.json new file mode 100644 index 00000000000..e2718e3e3cf --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=2304.json @@ -0,0 +1,1106 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=4096.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=4096.json new file mode 100644 index 00000000000..beb4290ff1e --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=4096.json @@ -0,0 +1,1108 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 1, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 1, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 1, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "NUM_KSPLIT": 4, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "NUM_KSPLIT": 4, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=576.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=576.json new file mode 100644 index 00000000000..26094c9bbfe --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=576.json @@ -0,0 +1,1347 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "NUM_KSPLIT": 1, + "num_warps": 2, + "num_stages": 2, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 3, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 2, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=8192.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=8192.json new file mode 100644 index 00000000000..a7e0d854386 --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=5120-K=8192.json @@ -0,0 +1,1101 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 4, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 4, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 4, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 512, + "K_PACK": 4, + "NUM_KSPLIT": 4, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 4, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 3, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": true + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64, + "FUSED_SPLITK": false + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 2, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 2, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 256, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=8192-K=1280.json b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=8192-K=1280.json new file mode 100644 index 00000000000..57620212ffd --- /dev/null +++ b/aiter/ops/triton/configs/gfx950/triton/gemm/gemm_a8w8_blockscale_group32/GEMM-A8W8_BLOCKSCALE_GROUP32-N=8192-K=1280.json @@ -0,0 +1,1167 @@ +{ + "M_BOUNDS": [ + 1, + 2, + 4, + 8, + 16, + 32, + 48, + 64, + 80, + 96, + 112, + 128, + 144, + 160, + 176, + 192, + 208, + 224, + 240, + 256, + 288, + 320, + 352, + 384, + 416, + 448, + 480, + 512, + 544, + 576, + 608, + 640, + 672, + 704, + 736, + 768, + 800, + 832, + 864, + 896, + 928, + 960, + 992, + 1024, + 1088, + 1152, + 1216, + 1280, + 1344, + 1408, + 1472, + 1536, + 1600, + 1664, + 1728, + 1792, + 1856, + 1920, + 1984, + 2048, + 4096, + 6144, + 8192, + 10240, + 12288, + 14336, + 16384, + 32768 + ], + "M_LEQ_1": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_2": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_4": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_8": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 0 + } + }, + "M_LEQ_16": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 1024, + "K_PACK": 1, + "NUM_KSPLIT": 1, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1, + "cache_modifier": ".cg" + } + }, + "M_LEQ_32": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 2, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_48": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_64": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 32, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 4, + "num_stages": 3, + "matrix_instr_nonkdim": 16, + "waves_per_eu": 1 + } + }, + "M_LEQ_80": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_96": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_112": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_128": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 512, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_144": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_160": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_176": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_208": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_224": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_256": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 3, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_288": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_320": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_352": { + "BLOCK_SIZE_M": 64, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "packed": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "K_PACK": 1, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16 + } + }, + "M_LEQ_384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 1, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 64 + }, + "M_LEQ_416": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_448": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_480": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_512": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_544": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_576": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_608": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_640": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_672": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_704": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_736": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_768": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_800": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_832": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_864": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_896": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_928": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_960": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_992": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": ".cg", + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1024": { + "BLOCK_SIZE_M": 256, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1088": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1152": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 64, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1216": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1280": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1344": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1408": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1472": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1536": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1600": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1664": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1728": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1792": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_1856": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1920": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_1984": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_2048": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_4096": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_6144": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_8192": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": false + }, + "M_LEQ_10240": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_12288": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_14336": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_16384": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 4, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "M_LEQ_32768": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 128, + "num_warps": 8, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 32, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": true, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + }, + "any": { + "BLOCK_SIZE_M": 128, + "BLOCK_SIZE_N": 128, + "BLOCK_SIZE_K": 256, + "num_warps": 4, + "num_stages": 1, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 0, + "cache_modifier": null, + "NUM_KSPLIT": 1, + "N_FIRST": false, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32 + } +} diff --git a/aiter/ops/triton/gemm/basic/gemm_a8w8_blockscale_group32.py b/aiter/ops/triton/gemm/basic/gemm_a8w8_blockscale_group32.py new file mode 100644 index 00000000000..89049b42169 --- /dev/null +++ b/aiter/ops/triton/gemm/basic/gemm_a8w8_blockscale_group32.py @@ -0,0 +1,253 @@ +# SPDX-License-Identifier: MIT +# Copyright (C) 2026, Advanced Micro Devices, Inc. All rights reserved. + +import torch +import triton + +from aiter.jit.utils.chip_info import get_gfx_runtime +from aiter.ops.triton._triton_kernels.common.splitk_reduce import ( + _gemm_splitk_reduce_kernel, +) +from aiter.ops.triton._triton_kernels.gemm.basic.gemm_a8w8_blockscale_group32 import ( + _gemm_a8w8_blockscale_group32_kernel, + _gemm_a8w8_blockscale_group32_packed_kernel, + _get_config, +) +from aiter.utility.graph_alloc import ROUTES_INSIDE_CAPTURE, persistent_alloc + +# One arrival counter per output tile of a fused split-K launch. +_SPLIT_COUNTERS = 1 << 16 +_split_counters_by_stream: dict[tuple[torch.device, int], torch.Tensor] = {} + + +def _split_counters(device: torch.device) -> torch.Tensor | None: + """Zeroed arrival counters per (device, stream), as for the a16w16 split-K + semaphores; kernels re-zero what they use. None when a stream's first use + falls inside a capture that cannot allocate outside the graph pool.""" + key = (device, torch.cuda.current_stream(device).cuda_stream) + counters = _split_counters_by_stream.get(key) + if counters is None: + if not ROUTES_INSIDE_CAPTURE and torch.cuda.is_current_stream_capturing(): + return None + with persistent_alloc(device): + counters = torch.zeros(_SPLIT_COUNTERS, dtype=torch.int32, device=device) + _split_counters_by_stream[key] = counters + return counters + + +def _split_k_partition(k: int, splits: int, step: int) -> tuple[int, int]: + """Partition size rounded up to whole steps, and the non-empty count.""" + part = -(-k // splits) + size = -(-part // step) * step + return size, -(-k // size) + + +def _fused_workspace(y, tiles, splits, block_m, block_n): + """(grid, partials, counters) for fused split-K, or None to run unfused.""" + assert tiles <= _SPLIT_COUNTERS, "Too many output tiles for fused split-K" + counters = _split_counters(y.device) + if counters is None: + return None + partials = torch.empty( + tiles * splits * block_m * block_n, dtype=torch.float32, device=y.device + ) + return (-(-tiles // 8) * 8 * splits,), partials, counters + + +def gemm_a8w8_blockscale_group32( + x: torch.Tensor, + w: torch.Tensor, + x_scale: torch.Tensor, + w_scale: torch.Tensor, + dtype: torch.dtype = torch.bfloat16, + y: torch.Tensor | None = None, + weight_group_rows: int = 32, + split_k: int | None = None, + config: dict | None = None, +) -> torch.Tensor: + """Compute X @ W.T using native E4M3 operands and E8M0 group32 scales. + + x and w are contiguous (M, K) and (N, K) matrices. Activation scales are + row-major (M, K // 32); compact weight scales are + (ceil(N / weight_group_rows), K // 32), with weight_group_rows 1 or 32. + E8M0 and E4M3 tensors may also be passed as their uint8 views. Weights and + scales are neither preshuffled nor expanded. + + Accumulation is FP32, with BF16, FP16 or FP32 output. y optionally supplies + a contiguous (M, N) output buffer. split_k is a positive partition count; + None uses NUM_KSPLIT and the optional packed-K variant from the tuning table. + config optionally overrides the table, following the Triton GEMM interface. + Explicit partitions are rounded to whole K tiles, then empty partitions + are removed. Split-K uses the shared AITER reduction without atomics; + a table FUSED_SPLITK (or packed NUM_KSPLIT) instead sums each tile's + partitions within the launch, on one XCD. A table cache_modifier applies + to the weight loads. + + This backend requires gfx950 microscaling MFMA. Small M is limited by + weight bandwidth and occupancy; packing K panels fills MFMA rows without + a global partial buffer. Larger M uses output tiling and activation reuse. + """ + assert x.ndim == w.ndim == 2, "Expected two matrix operands" + M, K = x.shape + N, weight_k = w.shape + assert ( + K > 0 and K % 32 == 0 and weight_k == K and N > 0 + ), "Expected matching positive K divisible by 32 and positive N" + assert weight_group_rows in (1, 32), "Weight scales must be 1x32 or 32x32" + assert x.dtype in (torch.float8_e4m3fn, torch.uint8), "x must be E4M3" + assert w.dtype in (torch.float8_e4m3fn, torch.uint8), "w must be E4M3" + assert x_scale.dtype in (torch.float8_e8m0fnu, torch.uint8), "x_scale must be E8M0" + assert w_scale.dtype in (torch.float8_e8m0fnu, torch.uint8), "w_scale must be E8M0" + assert x_scale.shape == (M, K // 32), "Invalid activation scale shape" + assert w_scale.shape == ( + -(-N // weight_group_rows), + K // 32, + ), "Invalid weight scale shape" + assert all( + t.is_cuda and t.device == x.device and t.is_contiguous() + for t in (x, w, x_scale, w_scale) + ), "Operands must be contiguous on the same GPU" + assert dtype in ( + torch.bfloat16, + torch.float16, + torch.float32, + ), "Output must be BF16, FP16 or FP32" + assert split_k is None or ( + type(split_k) is int and split_k > 0 + ), "split_k must be a positive integer or None" + if y is None: + y = torch.empty((M, N), dtype=dtype, device=x.device) + else: + assert ( + y.shape == (M, N) and y.dtype == dtype and y.device == x.device + ), "Invalid output shape, dtype or device" + assert y.is_contiguous(), "Output must be contiguous" + if M == 0: + return y + assert get_gfx_runtime() == "gfx950", "Group32 FP8 GEMM requires gfx950" + + x = x.view(torch.float8_e4m3fn) if x.dtype == torch.uint8 else x + w = w.view(torch.float8_e4m3fn) if w.dtype == torch.uint8 else w + x_scale = x_scale.view(torch.uint8) + w_scale = w_scale.view(torch.uint8) + if config is None: + config, _ = _get_config(M, N, K) + packed = ( + config.get("packed") if weight_group_rows == 32 and split_k is None else None + ) + launch_config = config if packed is None else packed + launch_options = { + key: launch_config[key] + for key in ("num_warps", "num_stages", "waves_per_eu", "matrix_instr_nonkdim") + } + # Triton's repr callback sees constexpr arguments, not compiler options. + # Carry the same values into the name without a second source of tuning. + launch_repr = tuple(launch_options.items()) + if packed is not None: + block_m, block_n = packed["BLOCK_SIZE_M"], packed["BLOCK_SIZE_N"] + block_k, k_pack = packed["BLOCK_SIZE_K"], packed["K_PACK"] + grid_m, grid_n = -(-M // block_m), -(-N // block_n) + split_k_size, num_splits = _split_k_partition( + K, packed.get("NUM_KSPLIT", 1), block_k * k_pack + ) + workspace = None + if num_splits > 1: + workspace = _fused_workspace( + y, grid_m * grid_n, num_splits, block_m, block_n + ) + if workspace is None: + workspace = (grid_m, grid_n), y, y + split_k_size, num_splits = 0, 1 + grid, partials, counters = workspace + _gemm_a8w8_blockscale_group32_packed_kernel[grid]( + x, + w, + x_scale, + w_scale, + y, + partials, + counters, + M, + N, + K, + block_m, + block_n, + block_k, + k_pack, + LAUNCH_OPTIONS=launch_repr, + SPLITK_BLOCK_SIZE=split_k_size, + FUSED_SPLITS=num_splits, + B_CACHE_MODIFIER=packed.get("cache_modifier"), + **launch_options, + ) + return y + + block_m = config["BLOCK_SIZE_M"] + block_n = config["BLOCK_SIZE_N"] + block_k = config["BLOCK_SIZE_K"] + n_first = config["N_FIRST"] + grid_m, grid_n = -(-M // block_m), -(-N // block_n) + fused = split_k is None and config.get("FUSED_SPLITK", False) + split_k = config["NUM_KSPLIT"] if split_k is None else split_k + split_k_size, num_splits = _split_k_partition(K, split_k, block_k) + workspace = None + if fused and num_splits > 1: + workspace = _fused_workspace(y, grid_m * grid_n, num_splits, block_m, block_n) + if workspace is not None: + grid, partials, counters = workspace + out, fused_splits = y, num_splits + else: + out = ( + y + if num_splits == 1 + else torch.empty((num_splits, M, N), dtype=torch.float32, device=x.device) + ) + grid = (grid_n, grid_m, num_splits) if n_first else (grid_m, grid_n, num_splits) + partials = counters = out + fused_splits = 1 + _gemm_a8w8_blockscale_group32_kernel[grid]( + x, + w, + x_scale, + w_scale, + out, + partials, + counters, + M, + N, + K, + weight_group_rows, + split_k_size, + block_m, + block_n, + block_k, + N_FIRST=n_first, + LAUNCH_OPTIONS=launch_repr, + FUSED_SPLITS=fused_splits, + B_CACHE_MODIFIER=config.get("cache_modifier"), + **launch_options, + ) + if fused_splits == 1 and num_splits > 1: + reduce_m = config["REDUCE_BLOCK_SIZE_M"] + reduce_n = config["REDUCE_BLOCK_SIZE_N"] + _gemm_splitk_reduce_kernel[(-(-M // reduce_m), -(-N // reduce_n))]( + out, + y, + None, + M, + N, + M * N, + N, + 1, + N, + 1, + BLOCK_SIZE_M=reduce_m, + BLOCK_SIZE_N=reduce_n, + ACTUAL_KSPLIT=num_splits, + MAX_KSPLIT=triton.next_power_of_2(num_splits), + ADD_BIAS=False, + activation=None, + use_activation=False, + KERNEL_NAME="_gemm_a8w8_blockscale_group32_reduce_kernel", + ) + return y diff --git a/aiter/ops/triton/utils/gemm_config_utils.py b/aiter/ops/triton/utils/gemm_config_utils.py index 14d75d14af9..17c00f1ed83 100644 --- a/aiter/ops/triton/utils/gemm_config_utils.py +++ b/aiter/ops/triton/utils/gemm_config_utils.py @@ -86,8 +86,15 @@ def _get_gemm_config_cached( config_dict, is_tuned = specialized_config, True break - # use standard bounds unless custom bounds are passed - search_bounds = bounds if bounds is not None else STANDARD_M_BOUNDS + # Explicit bounds override file-defined bounds; legacy files use the standard bounds. + search_bounds = ( + bounds if bounds is not None else config_dict.get("M_BOUNDS", STANDARD_M_BOUNDS) + ) + assert ( + len(search_bounds) > 0 + and all(type(x) is int and x > 0 for x in search_bounds) + and all(x < y for x, y in itertools.pairwise(search_bounds)) + ), "M_BOUNDS must be a non-empty sequence of strictly increasing positive integers" # Search for M_LEQ_x keys for bound in search_bounds: @@ -110,6 +117,18 @@ def _get_gemm_config_cached( ) +def _copy_gemm_config(config: dict) -> dict: + # JSON leaves are immutable. Avoid deepcopy's per-scalar overhead on the + # GEMM launch path while keeping nested variant configs safe to mutate. + result = config.copy() + for key, value in config.items(): + if isinstance(value, dict): + result[key] = _copy_gemm_config(value) + elif isinstance(value, list): + result[key] = copy.deepcopy(value) + return result + + def get_gemm_config( config_name: str, M: int, @@ -129,7 +148,7 @@ def get_gemm_config( 2. If B, N and K are provided, try B-specialized config: {config_name}-B={B}-N={N}-K={K}.json 3. If N and K are provided, try to load specialized config: {config_name}-N={N}-K={K}.json Or if specialized_filename is provided, use: {config_name}-{specialized_filename}.json - 4. Search for M_LEQ_x keys in order of bounds (default: STANDARD_M_BOUNDS) + 4. Search M_LEQ_x using explicit bounds, file M_BOUNDS or STANDARD_M_BOUNDS 5. If no M_LEQ_x matches, search for M_GEQ_x keys in reverse order 6. Fall back to "any" if no bounds match @@ -138,7 +157,7 @@ def get_gemm_config( M: M dimension of the GEMM N: N dimension of the GEMM (optional) K: K dimension of the GEMM (optional) - bounds: Custom bounds to use instead of STANDARD_M_BOUNDS (optional) + bounds: Override the file M_BOUNDS or STANDARD_M_BOUNDS (optional) specialized_filename: Custom specialized filename suffix (optional) backend: Backend whose config directory to read, "triton" (default) or "gluon". Declared by the caller; there is no fallback to @@ -152,7 +171,7 @@ def get_gemm_config( config, is_tuned = _get_gemm_config_cached( config_name, M, N, K, bounds, specialized_filename, backend, B ) - return copy.deepcopy(config), is_tuned + return _copy_gemm_config(config), is_tuned def add_default_gemm_config_params(config: dict) -> dict: diff --git a/op_tests/op_benchmarks/triton/bench_gemm_a8w8_blockscale_group32.py b/op_tests/op_benchmarks/triton/bench_gemm_a8w8_blockscale_group32.py new file mode 100644 index 00000000000..cbecc2baf1a --- /dev/null +++ b/op_tests/op_benchmarks/triton/bench_gemm_a8w8_blockscale_group32.py @@ -0,0 +1,131 @@ +# SPDX-License-Identifier: MIT +# Copyright (C) 2026, Advanced Micro Devices, Inc. All rights reserved. +"""Benchmark native group32 FP8 GEMM through its production config dispatch.""" + +import functools +import math + +import torch +import triton + +from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale_group32 import ( + gemm_a8w8_blockscale_group32, +) +from op_tests.op_benchmarks.triton.utils.argparse import ( + add_argparse_ff, + get_ff_args, + get_parser, +) +from op_tests.op_benchmarks.triton.utils.benchmark_utils import ( + get_caller_name_no_ext, + get_model_benchmark_object, + get_shape_benchmark_object, + print_vgpr, +) +from op_tests.triton_tests.gemm.basic.test_gemm_a8w8_blockscale_group32 import ( + generate_inputs, + run_torch, +) + +DTYPES = {"bf16": torch.bfloat16, "fp16": torch.float16, "fp32": torch.float32} + + +def bench_gemm_fn(m, n, k, metric, args): + dtype = DTYPES[args.dtype] + x, w, xs, ws = generate_inputs(m, n, k, args.weight_group_rows) + y = torch.empty((m, n), dtype=dtype, device=x.device) + run = functools.partial( + gemm_a8w8_blockscale_group32, + x, + w, + xs, + ws, + dtype=dtype, + y=y, + weight_group_rows=args.weight_group_rows, + ) + if args.test: + expected = run_torch(x, w, xs, ws, args.weight_group_rows) + torch.testing.assert_close( + run(), + expected.to(dtype), + rtol={"bf16": 0.016, "fp16": 0.002, "fp32": 3e-5}[args.dtype], + atol=5e-5 * expected.abs().max().item(), + ) + ms = ( + triton.testing.do_bench_cudagraph(run, rep=100) + if args.cudagraph + else triton.testing.do_bench(run, warmup=25, rep=100) + ) + if metric == "time": + return ms + if metric == "throughput": + return 2.0 * m * n * k / ms * 1e-9 + if metric == "bandwidth": + num_bytes = sum(t.numel() * t.element_size() for t in (x, w, xs, ws, y)) + return num_bytes / ms * 1e-6 + raise ValueError(f"Unknown metric: {metric}") + + +def run_benchmark(args): + name = get_caller_name_no_ext() + benchmark = ( + get_model_benchmark_object(name, args) + if args.model + else get_shape_benchmark_object(name, args) + ) + + @triton.testing.perf_report([benchmark]) + def bench( + M, + metric, + N=None, + K=None, + hidden_dim=None, + intermediate_dim=None, + layer=None, + **kwargs, + ): + if args.model: + if layer == "fc1": + N = intermediate_dim * (1 if args.no_glu else 2) + K = hidden_dim + else: + N, K = hidden_dim, math.ceil(intermediate_dim / args.tp) + if not args.model or layer == "fc1": + N = math.ceil(N / args.tp) + return bench_gemm_fn(M, N, K, metric, args) + + bench.run(save_path="." if args.o else None, print_data=True) + + +def parse_args(argv=None): + parser = add_argparse_ff(get_parser(kernel_name="Native group32 A8W8 GEMM")) + parser.add_argument("--weight-group-rows", type=int, choices=(1, 32), default=32) + parser.add_argument("--dtype", choices=tuple(DTYPES), default="bf16") + parser.add_argument( + "--cudagraph", action="store_true", help="Time CUDA graph replay" + ) + parser.add_argument( + "-test", action="store_true", help="Check against FP64 reference" + ) + args, _ = get_ff_args(parser, args=argv) + if args.shape is not None and (len(args.shape) != 3 or args.model): + parser.error("Use --shape M N K or --model with an optional -M") + if args.layout != "TN": + parser.error("Native group32 weights require --layout TN") + if args.tp < 1: + parser.error("-tp must be positive") + return args + + +def main(argv=None): + args = parse_args(argv) + if args.print_vgpr: + print_vgpr(lambda: run_benchmark(args), get_caller_name_no_ext()) + else: + run_benchmark(args) + + +if __name__ == "__main__": + main() diff --git a/op_tests/test_gemm_group32_interface.py b/op_tests/test_gemm_group32_interface.py new file mode 100644 index 00000000000..bcfe8377e87 --- /dev/null +++ b/op_tests/test_gemm_group32_interface.py @@ -0,0 +1,275 @@ +# SPDX-License-Identifier: MIT +# Copyright (C) 2026, Advanced Micro Devices, Inc. All rights reserved. +"""CPU contracts for A8W8 format dispatch and JSON-defined GEMM bounds.""" + +import json + +import pytest +import torch + +from aiter.jit.core import AITER_CONFIGS +from aiter.ops import gemm_op_a8w8 +from aiter.ops.triton.gemm.basic import gemm_a8w8_blockscale_group32 as group32 +from aiter.ops.triton.utils import gemm_config_utils +from aiter.ops.triton.utils.config_utils import load_config_json + + +@pytest.fixture +def config_dir(tmp_path, monkeypatch): + monkeypatch.setattr( + gemm_config_utils, "resolve_config_dir", lambda *args, **kwargs: str(tmp_path) + ) + gemm_config_utils._get_gemm_config_cached.cache_clear() + load_config_json.cache_clear() + yield tmp_path + gemm_config_utils._get_gemm_config_cached.cache_clear() + load_config_json.cache_clear() + + +def test_file_bounds_and_explicit_override(config_dir): + table = { + "M_BOUNDS": [3, 7], + "M_LEQ_3": {"value": 3}, + "M_LEQ_7": {"value": 7}, + "any": {"value": 99}, + } + (config_dir / "DEFAULT.json").write_text(json.dumps(table)) + assert gemm_config_utils.get_gemm_config("GEMM-TEST", 2)[0]["value"] == 3 + assert gemm_config_utils.get_gemm_config("GEMM-TEST", 4)[0]["value"] == 7 + assert gemm_config_utils.get_gemm_config("GEMM-TEST", 8)[0]["value"] == 99 + assert ( + gemm_config_utils.get_gemm_config("GEMM-TEST", 2, bounds=(7,))[0]["value"] == 7 + ) + + +def test_legacy_bounds_and_specialized_nested_copy(config_dir): + (config_dir / "DEFAULT.json").write_text( + json.dumps( + {"M_LEQ_3": {"value": 3}, "M_LEQ_4": {"value": 4}, "any": {"value": 99}} + ) + ) + assert gemm_config_utils.get_gemm_config("GEMM-TEST", 2)[0]["value"] == 4 + (config_dir / "GEMM-TEST-N=64-K=32.json").write_text( + json.dumps( + { + "M_BOUNDS": [3], + "M_LEQ_3": {"packed": {"K_PACK": 4}, "nested": [{"value": 1}]}, + "any": {"value": 99}, + } + ) + ) + first, tuned = gemm_config_utils.get_gemm_config("GEMM-TEST", 2, 64, 32) + assert tuned + first["packed"]["K_PACK"] = 1 + first["nested"][0]["value"] = 2 + assert gemm_config_utils.get_gemm_config("GEMM-TEST", 2, 64, 32)[0]["nested"] == [ + {"value": 1} + ] + assert ( + gemm_config_utils.get_gemm_config("GEMM-TEST", 2, 64, 32)[0]["packed"]["K_PACK"] + == 4 + ) + + +@pytest.mark.parametrize("bounds", [[], [4, 3], [3, 3], [0, 4], [1.5, 4], [True, 4]]) +def test_invalid_file_bounds(config_dir, bounds): + (config_dir / "DEFAULT.json").write_text( + json.dumps({"M_BOUNDS": bounds, "any": {"value": 1}}) + ) + with pytest.raises(AssertionError, match="M_BOUNDS"): + gemm_config_utils.get_gemm_config("GEMM-TEST", 2) + + +@pytest.mark.parametrize("group_n", [1, 32]) +@pytest.mark.parametrize("libtype", [None, "triton", "ck", "cktile", "unknown"]) +@pytest.mark.parametrize("split_k", [None, 3]) +@pytest.mark.parametrize( + "scale_dtypes", + [ + (torch.float8_e8m0fnu, torch.float8_e8m0fnu), + (torch.uint8, torch.uint8), + (torch.float8_e8m0fnu, torch.uint8), + (torch.uint8, torch.float8_e8m0fnu), + ], +) +def test_public_native_group32_route( + monkeypatch, group_n, libtype, split_k, scale_dtypes +): + calls = [] + expected = torch.empty((3, 65), dtype=torch.float32) + + def backend(x, w, xs, ws, **kwargs): + calls.append(kwargs) + return expected + + def lookup(m, n, k, tuned_file): + assert (m, n, k) == (3, 65, 64) + assert tuned_file.endswith("a8w8_blockscale_group32_tuned_gemm.csv") + calls.append("lookup") + return None if libtype is None else {"libtype": libtype} + + monkeypatch.setattr(gemm_op_a8w8, "get_CKGEMM_config", lookup) + monkeypatch.setattr(group32, "gemm_a8w8_blockscale_group32", backend) + x = torch.empty((3, 64), dtype=torch.float8_e4m3fn) + w = torch.empty((65, 64), dtype=torch.float8_e4m3fn) + xs = torch.empty((3, 2), dtype=scale_dtypes[0]) + ws = torch.empty(((65 + group_n - 1) // group_n, 2), dtype=scale_dtypes[1]) + if libtype not in (None, "triton"): + with pytest.raises(AssertionError, match="Unsupported libtype"): + gemm_op_a8w8.gemm_a8w8_blockscale( + x, w, xs, ws, dtype=torch.float32, split_k=split_k + ) + assert calls == ["lookup"] + return + actual = gemm_op_a8w8.gemm_a8w8_blockscale( + x, w, xs, ws, dtype=torch.float32, split_k=split_k + ) + assert actual is expected + assert calls == [ + "lookup", + {"dtype": torch.float32, "weight_group_rows": group_n, "split_k": split_k}, + ] + + +@pytest.mark.parametrize("libtype", ["ck", "cktile"]) +def test_public_legacy_ck_route_and_readonly_schema(monkeypatch, libtype): + calls = [] + + def ck(x, w, xs, ws, out, **kwargs): + calls.append(kwargs) + return out + + monkeypatch.setattr(gemm_op_a8w8, "_hip_blockscale_supported", lambda: True) + monkeypatch.setattr( + gemm_op_a8w8, + "get_CKGEMM_config", + lambda *args: {"libtype": libtype, "splitK": 2, "kernelName": "existing"}, + ) + monkeypatch.setattr(gemm_op_a8w8, f"gemm_a8w8_blockscale_{libtype}", ck) + x = torch.empty((3, 128), dtype=torch.float8_e4m3fn) + w = torch.empty((128, 128), dtype=torch.float8_e4m3fn) + xs = torch.empty((3, 1), dtype=torch.float32) + ws = torch.empty((1, 1), dtype=torch.float32) + actual = gemm_op_a8w8.gemm_a8w8_blockscale(x, w, xs, ws) + assert actual.shape == (3, 128) + assert calls == [{"splitK": 2, "kernelName": "existing"}] + schema = torch.ops.aiter.gemm_a8w8_blockscale.default._schema + assert all(arg.alias_info is None for arg in schema.arguments) + + +@pytest.fixture +def backend_config_files(tmp_path, monkeypatch): + monkeypatch.setattr(gemm_op_a8w8, "get_gfx", lambda: "gfx950") + monkeypatch.setattr(gemm_op_a8w8, "get_cu_num", lambda: 256) + monkeypatch.setattr(gemm_op_a8w8, "_CKGEMM_CONFIG_CACHE", {}) + monkeypatch.setattr(gemm_op_a8w8, "_CKGEMM_HAS_GFX", {}) + gemm_op_a8w8.get_CKGEMM_config.cache_clear() + AITER_CONFIGS.get_config_file.cache_clear() + files = {} + for suffix in ("", "_GROUP32"): + path = tmp_path / f"blockscale{suffix.lower()}.csv" + monkeypatch.setenv(f"AITER_CONFIG_GEMM_A8W8_BLOCKSCALE{suffix}", str(path)) + files[suffix] = path + yield files + gemm_op_a8w8.get_CKGEMM_config.cache_clear() + AITER_CONFIGS.get_config_file.cache_clear() + + +@pytest.mark.parametrize("libtype", [None, "triton", "unknown"]) +def test_scale_formats_have_independent_backend_configs( + backend_config_files, monkeypatch, libtype +): + # Identical M/N/K must not make the FP32 128x128 winner apply to E8M0. + header = "gfx,cu_num,M,N,K,libtype,splitK,kernelName\n" + backend_config_files[""].write_text(header + "gfx950,256,4,128,128,ck,2,legacy\n") + backend_config_files["_GROUP32"].write_text( + header + + ("" if libtype is None else f"gfx950,256,4,128,128,{libtype},0,native\n") + ) + calls = [] + + def ck(x, w, xs, ws, out, **kwargs): + assert xs.dtype == ws.dtype == torch.float32 + assert kwargs == {"splitK": 2, "kernelName": "legacy"} + calls.append("ck") + return out + + def native(x, w, xs, ws, **kwargs): + assert xs.dtype == ws.dtype == torch.float8_e8m0fnu + calls.append("group32") + return torch.empty((4, 128), dtype=kwargs["dtype"]) + + monkeypatch.setattr(gemm_op_a8w8, "gemm_a8w8_blockscale_ck", ck) + monkeypatch.setattr(group32, "gemm_a8w8_blockscale_group32", native) + x = torch.empty((4, 128), dtype=torch.float8_e4m3fn) + w = torch.empty((128, 128), dtype=torch.float8_e4m3fn) + gemm_op_a8w8.gemm_a8w8_blockscale(x, w, torch.empty((4, 1)), torch.empty((1, 1))) + assert calls == ["ck"] + xs = torch.empty((4, 4), dtype=torch.float8_e8m0fnu) + ws = torch.empty((4, 4), dtype=torch.float8_e8m0fnu) + if libtype == "unknown": + with pytest.raises(AssertionError, match="Unsupported libtype unknown"): + gemm_op_a8w8.gemm_a8w8_blockscale(x, w, xs, ws) + assert calls == ["ck"] + else: + gemm_op_a8w8.gemm_a8w8_blockscale(x, w, xs, ws) + assert calls == ["ck", "group32"] + + +@pytest.mark.parametrize("configured", [False, True]) +def test_legacy_triton_config_and_fallback(monkeypatch, configured): + calls = [] + + def lookup(*args): + calls.append("lookup") + return {"libtype": "triton"} if configured else None + + def triton(x, w, xs, ws, dtype, **kwargs): + calls.append("triton") + return torch.empty((3, 128), dtype=dtype) + + monkeypatch.setattr(gemm_op_a8w8, "get_CKGEMM_config", lookup) + monkeypatch.setattr(gemm_op_a8w8, "_hip_blockscale_supported", lambda: False) + monkeypatch.setattr(gemm_op_a8w8, "_blockscale_triton", triton) + gemm_op_a8w8.gemm_a8w8_blockscale( + torch.empty((3, 128), dtype=torch.float8_e4m3fn), + torch.empty((128, 128), dtype=torch.float8_e4m3fn), + torch.empty((3, 1)), + torch.empty((1, 1)), + ) + assert calls == ["lookup", "triton"] + + +def test_legacy_scale_format_rejects_native_split_override(): + with pytest.raises( + AssertionError, match="split_k override requires native group32" + ): + gemm_op_a8w8.gemm_a8w8_blockscale( + torch.empty((3, 128), dtype=torch.float8_e4m3fn), + torch.empty((128, 128), dtype=torch.float8_e4m3fn), + torch.empty((3, 1)), + torch.empty((1, 1)), + split_k=3, + ) + + +@pytest.mark.parametrize( + "invalid", ["activation_groups", "weight_groups", "mixed_formats"] +) +def test_invalid_byte_scales_never_reach_legacy_dispatch(monkeypatch, invalid): + def lookup(*args): + pytest.fail("Invalid byte scales reached backend configuration lookup") + + monkeypatch.setattr(gemm_op_a8w8, "get_CKGEMM_config", lookup) + x = torch.empty((3, 128), dtype=torch.float8_e4m3fn) + w = torch.empty((65, 128), dtype=torch.float8_e4m3fn) + xs = torch.empty((3, 4), dtype=torch.uint8) + ws = torch.empty((3, 4), dtype=torch.uint8) + if invalid == "activation_groups": + xs = torch.empty((3, 1), dtype=torch.uint8) + elif invalid == "weight_groups": + ws = torch.empty((3, 1), dtype=torch.uint8) + else: + xs = xs.float() + with pytest.raises(AssertionError, match="Expected E8M0 group32 scale shapes"): + gemm_op_a8w8.gemm_a8w8_blockscale(x, w, xs, ws) diff --git a/op_tests/triton_tests/gemm/basic/test_gemm_a8w8_blockscale_group32.py b/op_tests/triton_tests/gemm/basic/test_gemm_a8w8_blockscale_group32.py new file mode 100644 index 00000000000..f58c1c68f54 --- /dev/null +++ b/op_tests/triton_tests/gemm/basic/test_gemm_a8w8_blockscale_group32.py @@ -0,0 +1,447 @@ +# SPDX-License-Identifier: MIT +# Copyright (C) 2026, Advanced Micro Devices, Inc. All rights reserved. +"""Numerical and replay contracts for native group32 FP8 projections.""" + +import pytest +import torch + +if not torch.cuda.is_available(): + pytest.skip("ROCm GPU required", allow_module_level=True) + +from aiter.jit.utils.chip_info import get_gfx +from aiter.ops import gemm_op_a8w8 +from aiter.ops.gemm_op_a8w8 import gemm_a8w8_blockscale +from aiter.ops.triton.gemm.basic import gemm_a8w8_blockscale_group32 as group32_op +from aiter.ops.triton.gemm.basic.gemm_a8w8_blockscale_group32 import ( + gemm_a8w8_blockscale_group32, +) +from aiter.ops.triton.utils.gemm_config_utils import get_gemm_config +from aiter.test_common import checkAllclose + +pytestmark = pytest.mark.skipif(get_gfx() != "gfx950", reason="CDNA4 packed MFMA") + + +def generate_inputs(m, n, k, weight_group_rows=32): + torch.manual_seed(m + n + k) + x = torch.randn(m, k, device="cuda").to(torch.float8_e4m3fn) + weight = torch.randn(n, k, device="cuda").to(torch.float8_e4m3fn) + xs = torch.randint(122, 131, (m, k // 32), device="cuda", dtype=torch.uint8) + ws = torch.randint( + 122, + 131, + (-(-n // weight_group_rows), k // 32), + device="cuda", + dtype=torch.uint8, + ) + return x, weight, xs.view(torch.float8_e8m0fnu), ws.view(torch.float8_e8m0fnu) + + +def run_torch(x, weight, xs, ws, weight_group_rows=32): + # FP64 is independent of the MFMA's internal block accumulation and of + # either implementation's K reduction order. + a = x.double() * xs.double().repeat_interleave(32, -1) + b = weight.double() * ws.double().repeat_interleave(weight_group_rows, 0)[ + : weight.shape[0] + ].repeat_interleave(32, -1) + return a @ b.T + + +@pytest.mark.parametrize( + "m,n,k", + [ + (1, 5120, 576), + (3, 2053, 1280), + (4, 8192, 1280), + (8, 5120, 1152), + (16, 4096, 1280), + (31, 4096, 1280), + (3, 5120, 2048), + (8, 5120, 2304), + (4, 5120, 4096), + (3, 2304, 5120), + (1, 5120, 8192), + # Untuned geometries exercise the default configuration with tails. + (63, 8193, 1280), + (129, 4097, 576), + (255, 16385, 1152), + (1023, 4097, 1280), + ], +) +@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16, torch.float32]) +def test_group32_projection_panel_scales_and_tails(m, n, k, dtype): + x, weight, xs, ws = generate_inputs(m, n, k) + actual = gemm_a8w8_blockscale_group32(x, weight, xs, ws, dtype=dtype) + expected = run_torch(x, weight, xs, ws) + peak = expected.abs().max().item() + # Keep the established native-MFMA FP32 bound. BF16 additionally rounds + # output; near cancellation still uses the same peak-relative floor. + torch.testing.assert_close( + actual.float(), + expected.to(dtype).float(), + rtol=( + 0.016 + if dtype == torch.bfloat16 + else 0.002 if dtype == torch.float16 else 3e-5 + ), + atol=5e-5 * peak, + ) + + +@pytest.mark.parametrize("m,packed", [(3, True), (129, False)]) +@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16, torch.float32]) +def test_tuned_variants_use_production_config(m, packed, dtype): + n, k = 8192, 1280 + config, tuned = get_gemm_config("GEMM-A8W8_BLOCKSCALE_GROUP32", m, n, k) + assert tuned + assert ("packed" in config) == packed + if not packed: + assert config["N_FIRST"] + x, w, xs, ws = generate_inputs(m, n, k) + actual = gemm_a8w8_blockscale_group32(x, w, xs, ws, dtype=dtype) + expected = run_torch(x, w, xs, ws) + torch.testing.assert_close( + actual.float(), + expected.to(dtype).float(), + rtol=( + 0.016 + if dtype == torch.bfloat16 + else 0.002 if dtype == torch.float16 else 3e-5 + ), + atol=5e-5 * expected.abs().max().item(), + ) + + +@pytest.mark.parametrize( + "a_code,b_code", [(0, 254), (254, 0), (128, 0), (255, 127), (127, 255)] +) +def test_group32_projection_extreme_scale_codes(a_code, b_code): + x = torch.ones(3, 1280, device="cuda").to(torch.float8_e4m3fn) + weight = torch.ones(2048, 1280, device="cuda").to(torch.float8_e4m3fn) + xs = torch.full((3, 40), a_code, device="cuda", dtype=torch.uint8).view( + torch.float8_e8m0fnu + ) + ws = torch.full((64, 40), b_code, device="cuda", dtype=torch.uint8).view( + torch.float8_e8m0fnu + ) + actual = gemm_a8w8_blockscale_group32(x, weight, xs, ws, dtype=torch.float32) + if 255 in (a_code, b_code): + assert actual.isnan().all() + else: + expected = torch.full_like(actual, 1280 * 2.0 ** (a_code + b_code - 254)) + torch.testing.assert_close(actual, expected, rtol=0, atol=0) + + +@pytest.mark.parametrize( + "rows,n,k,split_k", [(3, 4096, 1280, None), (33, 8193, 576, 3)] +) +def test_group32_projection_graph_reads_live_inputs_and_scales(rows, n, k, split_k): + x, weight, xs, ws = generate_inputs(2 * rows, n, k) + gemm_a8w8_blockscale_group32( + x, weight, xs, ws, dtype=torch.float32, split_k=split_k + ) + torch.cuda.synchronize() + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + actual = gemm_a8w8_blockscale_group32( + x, weight, xs, ws, dtype=torch.float32, split_k=split_k + ) + for factor in (2, 0.5): + x.copy_((x.float() * factor).to(x.dtype)) + xs.view(torch.uint8).add_(1) + graph.replay() + expected = run_torch(x, weight, xs, ws) + torch.testing.assert_close( + actual.float(), + expected.float(), + rtol=3e-5, + atol=5e-5 * expected.abs().max().item(), + ) + + +@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16, torch.float32]) +def test_row_scaled_fp8_projection_with_token_and_column_tails(dtype): + x, weight, xs, ws = generate_inputs(97, 8193, 576) + # Per-row scales share the scheduling kernel but not the compact weight grid. + ws = ws.repeat_interleave(32, 0)[: weight.shape[0]].contiguous() + actual = gemm_a8w8_blockscale_group32( + x, weight, xs, ws, weight_group_rows=1, dtype=dtype, split_k=3 + ) + expected = (x.double() * xs.double().repeat_interleave(32, -1)) @ ( + weight.double() * ws.double().repeat_interleave(32, -1) + ).T + torch.testing.assert_close( + actual.float(), + expected.to(dtype).float(), + rtol=( + 0.016 + if dtype == torch.bfloat16 + else 0.002 if dtype == torch.float16 else 3e-5 + ), + atol=5e-5 * expected.abs().max().item(), + ) + + +@pytest.mark.parametrize( + "m,n,k", [(0, 65, 64), (3, 2053, 1280), (66, 8193, 576), (512, 4096, 1280)] +) +@pytest.mark.parametrize("group_n", [1, 32]) +@pytest.mark.parametrize("configured", [False, True]) +@pytest.mark.parametrize("raw_views", [False, True]) +def test_public_group32_dispatch(m, n, k, group_n, configured, raw_views, monkeypatch): + if configured: + monkeypatch.setattr( + gemm_op_a8w8, "get_CKGEMM_config", lambda *args: {"libtype": "triton"} + ) + x, w, xs, ws = generate_inputs(m, n, k) + if group_n == 1: + ws = ws.repeat_interleave(32, 0)[:n].contiguous() + expected = gemm_a8w8_blockscale_group32(x, w, xs, ws, weight_group_rows=group_n) + operands = (x, w, xs, ws) + if raw_views: + operands = tuple(t.view(torch.uint8) for t in operands) + actual = gemm_a8w8_blockscale(*operands) + torch.testing.assert_close(actual, expected, rtol=0, atol=0) + + +@pytest.mark.parametrize("split_k", [None, 1, 3, 100]) +def test_raw_views_and_preallocated_output(split_k): + x, w, xs, ws = generate_inputs(63, 8193, 576) + output = torch.empty((63, 8193), dtype=torch.bfloat16, device="cuda") + actual = gemm_a8w8_blockscale_group32( + x.view(torch.uint8), + w.view(torch.uint8), + xs.view(torch.uint8), + ws.view(torch.uint8), + y=output, + split_k=split_k, + ) + assert actual is output + expected = run_torch(x, w, xs, ws).to(output.dtype) + checkAllclose( + expected, + actual, + rtol=0.016, + atol=5e-5 * expected.abs().max().item(), + catastrophic_check=True, + ) + + +@pytest.mark.parametrize( + "invalid", ["stride", "scale_shape", "scale_dtype", "output", "split"] +) +def test_invalid_group32_contract(invalid): + x, w, xs, ws = generate_inputs(3, 2053, 1280) + kwargs = {} + if invalid == "stride": + x = x.T.contiguous().T + elif invalid == "scale_shape": + ws = ws[:-1] + elif invalid == "scale_dtype": + xs = xs.float() + elif invalid == "output": + kwargs["y"] = torch.empty((3, 2052), dtype=torch.bfloat16, device="cuda") + else: + kwargs["split_k"] = 0 + with pytest.raises(AssertionError): + gemm_a8w8_blockscale_group32(x, w, xs, ws, **kwargs) + + +@pytest.mark.parametrize("split_k", [None, 3]) +def test_public_group32_compile_dynamic_rows(split_k): + def forward(x, w, xs, ws): + return gemm_a8w8_blockscale(x, w, xs, ws, split_k=split_k) + + compiled = torch.compile(forward, fullgraph=True, dynamic=True) + for m in (3, 7, 65): + x, w, xs, ws = generate_inputs(m, 4096, 1280) + torch.testing.assert_close( + compiled(x, w, xs, ws), forward(x, w, xs, ws), rtol=0, atol=0 + ) + + +@pytest.mark.parametrize("configured", [False, True]) +@pytest.mark.parametrize("split_k", [None, 3]) +def test_public_group32_graph_replay(configured, split_k, monkeypatch): + if configured: + monkeypatch.setattr( + gemm_op_a8w8, "get_CKGEMM_config", lambda *args: {"libtype": "triton"} + ) + x, w, xs, ws = generate_inputs(3, 4096, 1280) + gemm_a8w8_blockscale(x, w, xs, ws, split_k=split_k) + torch.cuda.synchronize() + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + actual = gemm_a8w8_blockscale(x, w, xs, ws, split_k=split_k) + x.copy_((x.float() * 0.5).to(x.dtype)) + xs.view(torch.uint8).add_(1) + graph.replay() + expected = gemm_a8w8_blockscale_group32(x, w, xs, ws, split_k=split_k) + torch.testing.assert_close(actual, expected, rtol=0, atol=0) + + +@pytest.mark.parametrize("n", [1, 2, 16, 31, 32, 33]) +@pytest.mark.parametrize("group_n", [1, 32]) +@pytest.mark.parametrize("raw_views", [False, True]) +def test_public_small_n_scale_layouts(n, group_n, raw_views): + x, w, xs, ws = generate_inputs(3, n, 128, weight_group_rows=group_n) + expected = run_torch(x, w, xs, ws, weight_group_rows=group_n) + operands = (x, w, xs, ws) + if raw_views: + operands = tuple(t.view(torch.uint8) for t in operands) + actual = gemm_a8w8_blockscale(*operands, dtype=torch.float32) + torch.testing.assert_close( + actual.double(), expected, rtol=3e-5, atol=5e-5 * expected.abs().max().item() + ) + + +_FUSED_TILE = { + "BLOCK_SIZE_M": 16, + "BLOCK_SIZE_N": 32, + "BLOCK_SIZE_K": 256, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + "NUM_KSPLIT": 5, + "N_FIRST": False, + "REDUCE_BLOCK_SIZE_M": 32, + "REDUCE_BLOCK_SIZE_N": 32, + "FUSED_SPLITK": True, +} +_FUSED_PACKED = dict( + _FUSED_TILE, + packed={ + "BLOCK_SIZE_M": 8, + "BLOCK_SIZE_N": 16, + "BLOCK_SIZE_K": 256, + "K_PACK": 2, + "NUM_KSPLIT": 5, + "num_warps": 2, + "num_stages": 2, + "waves_per_eu": 0, + "matrix_instr_nonkdim": 16, + }, +) +_FUSED = {"tile": _FUSED_TILE, "packed": _FUSED_PACKED} + + +def _assert_matches_reference(actual, x, w, xs, ws, dtype, weight_group_rows=32): + expected = run_torch(x, w, xs, ws, weight_group_rows) + torch.testing.assert_close( + actual.float(), + expected.to(dtype).float(), + rtol=( + 0.016 + if dtype == torch.bfloat16 + else 0.002 if dtype == torch.float16 else 3e-5 + ), + atol=5e-5 * expected.abs().max().item(), + ) + + +@pytest.mark.parametrize("variant", ["tile", "packed"]) +@pytest.mark.parametrize( + "m,n,k", [(1, 512, 5120), (3, 2053, 1280), (16, 1152, 5120), (37, 5120, 2304)] +) +@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16, torch.float32]) +def test_fused_split_k_matches_reference(variant, m, n, k, dtype): + x, w, xs, ws = generate_inputs(m, n, k) + actual = gemm_a8w8_blockscale_group32( + x, w, xs, ws, dtype=dtype, config=_FUSED[variant] + ) + _assert_matches_reference(actual, x, w, xs, ws, dtype) + + +@pytest.mark.parametrize("variant", ["tile", "packed"]) +@pytest.mark.parametrize("m,n,k", [(3, 2053, 1280), (37, 5120, 2336)]) +def test_weight_cache_modifier_keeps_results(variant, m, n, k): + x, w, xs, ws = generate_inputs(m, n, k) + config = _FUSED[variant] + cached = dict(config, cache_modifier=".cg") + if variant == "packed": + cached["packed"] = dict(config["packed"], cache_modifier=".cg") + run = gemm_a8w8_blockscale_group32 + expected = run(x, w, xs, ws, dtype=torch.float32, config=config) + actual = run(x, w, xs, ws, dtype=torch.float32, config=cached) + torch.testing.assert_close(actual, expected, rtol=0, atol=0) + + +def test_fused_split_k_row_scales(): + x, w, xs, ws = generate_inputs(5, 2053, 1280, weight_group_rows=1) + actual = gemm_a8w8_blockscale_group32( + x, w, xs, ws, weight_group_rows=1, dtype=torch.float32, config=_FUSED_TILE + ) + _assert_matches_reference(actual, x, w, xs, ws, torch.float32, 1) + + +@pytest.mark.parametrize("variant", ["tile", "packed"]) +def test_fused_split_k_is_deterministic_and_rezeroes_counters(variant): + x, w, xs, ws = generate_inputs(4, 1152, 5120) + run = lambda: gemm_a8w8_blockscale_group32( + x, w, xs, ws, dtype=torch.float32, config=_FUSED[variant] + ) + first = run() + # A stale counter or an early partial read would change a later sum. + assert all(torch.equal(first, run()) for _ in range(50)) + torch.cuda.synchronize() + assert not group32_op._split_counters(x.device).any() + + +@pytest.mark.parametrize("variant", ["tile", "packed"]) +def test_fused_split_k_graph_replay_on_warmed_stream(variant): + x, w, xs, ws = generate_inputs(6, 1152, 5120) + stream = torch.cuda.Stream() + with torch.cuda.stream(stream): + gemm_a8w8_blockscale_group32( + x, w, xs, ws, dtype=torch.float32, config=_FUSED[variant] + ) + stream.synchronize() + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph, stream=stream): + actual = gemm_a8w8_blockscale_group32( + x, w, xs, ws, dtype=torch.float32, config=_FUSED[variant] + ) + for factor in (2, 0.5): + x.copy_((x.float() * factor).to(x.dtype)) + xs.view(torch.uint8).add_(1) + graph.replay() + torch.cuda.synchronize() + _assert_matches_reference(actual, x, w, xs, ws, torch.float32) + + +@pytest.mark.parametrize("variant", ["tile", "packed"]) +def test_fused_split_k_first_use_inside_capture(variant): + # Before torch 2.10 counters cannot be allocated mid-capture: runs unfused. + x, w, xs, ws = generate_inputs(6, 1152, 5120) + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + actual = gemm_a8w8_blockscale_group32( + x, w, xs, ws, dtype=torch.float32, config=_FUSED[variant] + ) + graph.replay() + torch.cuda.synchronize() + _assert_matches_reference(actual, x, w, xs, ws, torch.float32) + + +@pytest.mark.parametrize("variant", ["tile", "packed"]) +def test_fused_split_k_concurrent_streams(variant): + inputs = [generate_inputs(m, 1152, 5120) for m in (3, 5)] + expected = [ + gemm_a8w8_blockscale_group32(*t, dtype=torch.float32, config=_FUSED[variant]) + for t in inputs + ] + torch.cuda.synchronize() + streams = [torch.cuda.Stream() for _ in inputs] + outputs = [[], []] + # Unsynchronized interleaving keeps both streams' reductions in flight. + for _ in range(40): + for i, (stream, operands) in enumerate(zip(streams, inputs)): + with torch.cuda.stream(stream): + outputs[i].append( + gemm_a8w8_blockscale_group32( + *operands, dtype=torch.float32, config=_FUSED[variant] + ) + ) + torch.cuda.synchronize() + for want, got in zip(expected, outputs): + assert all(torch.equal(want, y) for y in got)