nki(matrix_transpose): NKI (Trainium) implementation

results/B200/csv/matrix_transpose_autotune.csv

@@ -1,81 +1,81 @@
-params,dtype,torch_ms,triton_ms,cutile_ms,speedup_triton,speedup_cutile,triton_vs_cutile,tilelang_ms,speedup_tilelang
-n=1024,fp16,0.0206,0.0040,0.0056,5.11,3.70,1.3809,0.0050,4.10
-n=1024,bf16,0.0204,0.0039,0.0055,5.24,3.72,1.4083,0.0052,3.93
-n=1024,fp32,0.0213,0.0076,0.0076,2.81,2.81,1.0011,0.0080,2.65
-n=1024,int8,0.0193,0.0035,0.0035,5.52,5.49,1.0038,0.0056,3.46
-n=2048,fp16,0.0369,0.0079,0.0079,4.64,4.69,0.9897,0.0106,3.47
-n=2048,bf16,0.0371,0.0079,0.0079,4.67,4.71,0.9920,0.0099,3.76
-n=2048,fp32,0.0402,0.0131,0.0131,3.06,3.07,0.9970,0.0140,2.86
-n=2048,int8,0.0346,0.0061,0.0054,5.64,6.41,0.8789,0.0084,4.10
-n=3072,fp16,0.0558,0.0103,0.0105,5.43,5.30,1.0247,0.0129,4.32
-n=3072,bf16,0.0558,0.0104,0.0107,5.38,5.23,1.0302,0.0128,4.34
-n=3072,fp32,0.0608,0.0186,0.0187,3.27,3.24,1.0094,0.0189,3.21
-n=3072,int8,0.0528,0.0076,0.0069,6.91,7.61,0.9076,0.0120,4.39
-n=4096,fp16,0.0700,0.0132,0.0142,5.31,4.92,1.0789,0.0160,4.38
-n=4096,bf16,0.0699,0.0133,0.0138,5.26,5.07,1.0377,0.0161,4.35
-n=4096,fp32,0.0774,0.0238,0.0237,3.25,3.26,0.9964,0.0253,3.06
-n=4096,int8,0.0664,0.0079,0.0080,8.41,8.32,1.0112,0.0150,4.42
-n=5120,fp16,0.0920,0.0168,0.0166,5.48,5.54,0.9906,0.0187,4.92
-n=5120,bf16,0.0919,0.0166,0.0159,5.54,5.77,0.9608,0.0187,4.92
-n=5120,fp32,0.1022,0.0287,0.0283,3.56,3.61,0.9837,0.0291,3.52
-n=5120,int8,0.0840,0.0104,0.0110,8.06,7.64,1.0551,0.0184,4.57
-n=6144,fp16,0.1068,0.0185,0.0188,5.77,5.69,1.0133,0.0214,5.00
-n=6144,bf16,0.1067,0.0183,0.0188,5.82,5.67,1.0259,0.0213,5.01
-n=6144,fp32,0.1183,0.0336,0.0335,3.53,3.53,1.0000,0.0345,3.42
-n=6144,int8,0.1019,0.0122,0.0116,8.34,8.76,0.9520,0.0208,4.90
-n=7168,fp16,0.1230,0.0211,0.0212,5.83,5.80,1.0047,0.0240,5.12
-n=7168,bf16,0.1232,0.0211,0.0213,5.83,5.79,1.0064,0.0239,5.15
-n=7168,fp32,0.1339,0.0392,0.0392,3.42,3.41,1.0007,0.0395,3.39
-n=7168,int8,0.1201,0.0136,0.0132,8.82,9.10,0.9694,0.0238,5.04
-n=8192,fp16,0.1350,0.0239,0.0241,5.64,5.60,1.0073,0.0267,5.05
-n=8192,bf16,0.1352,0.0238,0.0239,5.67,5.65,1.0051,0.0268,5.04
-n=8192,fp32,0.1502,0.0436,0.0436,3.45,3.45,0.9990,0.0445,3.37
-n=8192,int8,0.1291,0.0152,0.0143,8.49,9.02,0.9416,0.0265,4.87
-n=9216,fp16,0.1553,0.0263,0.0265,5.90,5.86,1.0061,0.0296,5.25
-n=9216,bf16,0.1554,0.0267,0.0269,5.82,5.77,1.0083,0.0295,5.27
-n=9216,fp32,0.1712,0.0488,0.0494,3.51,3.47,1.0125,0.0501,3.41
-n=9216,int8,0.1497,0.0161,0.0158,9.29,9.49,0.9794,0.0296,5.05
-n=10240,fp16,0.1815,0.0285,0.0289,6.37,6.28,1.0146,0.0324,5.60
-n=10240,bf16,0.1813,0.0286,0.0290,6.33,6.25,1.0134,0.0324,5.60
-n=10240,fp32,0.1986,0.0537,0.0538,3.70,3.69,1.0010,0.0543,3.66
-n=10240,int8,0.1728,0.0177,0.0171,9.78,10.10,0.9682,0.0325,5.31
-n=11264,fp16,0.1866,0.0318,0.0319,5.86,5.85,1.0029,0.0350,5.33
-n=11264,bf16,0.1865,0.0324,0.0321,5.75,5.80,0.9918,0.0349,5.34
-n=11264,fp32,0.2097,0.0586,0.0596,3.58,3.52,1.0160,0.0604,3.47
-n=11264,int8,0.1795,0.0195,0.0185,9.18,9.72,0.9445,0.0352,5.10
-n=12288,fp16,0.2105,0.0352,0.0346,5.98,6.08,0.9838,0.0376,5.60
-n=12288,bf16,0.2106,0.0345,0.0346,6.10,6.09,1.0010,0.0374,5.63
-n=12288,fp32,0.2314,0.0639,0.0640,3.62,3.62,1.0011,0.0651,3.56
-n=12288,int8,0.2002,0.0206,0.0198,9.72,10.13,0.9592,0.0382,5.25
-n=13312,fp16,0.2213,0.0372,0.0367,5.94,6.04,0.9844,0.0400,5.53
-n=13312,bf16,0.2218,0.0364,0.0372,6.09,5.96,1.0218,0.0403,5.51
-n=13312,fp32,0.2446,0.0689,0.0689,3.55,3.55,0.9993,0.0702,3.48
-n=13312,int8,0.2119,0.0220,0.0211,9.63,10.06,0.9573,0.0408,5.20
-n=14336,fp16,0.2383,0.0390,0.0393,6.11,6.06,1.0076,0.0429,5.55
-n=14336,bf16,0.2383,0.0391,0.0392,6.10,6.08,1.0029,0.0427,5.58
-n=14336,fp32,0.2606,0.0737,0.0743,3.53,3.51,1.0080,0.0754,3.45
-n=14336,int8,0.2320,0.0228,0.0220,10.16,10.52,0.9660,0.0437,5.30
-n=15360,fp16,0.2614,0.0410,0.0413,6.37,6.32,1.0079,0.0457,5.72
-n=15360,bf16,0.2616,0.0414,0.0416,6.32,6.29,1.0044,0.0459,5.70
-n=15360,fp32,0.2898,0.0799,0.0802,3.63,3.61,1.0043,0.0812,3.57
-n=15360,int8,0.2462,0.0246,0.0238,10.03,10.34,0.9697,0.0467,5.27
-n=16384,fp16,0.2652,0.0435,0.0440,6.10,6.03,1.0122,0.0479,5.54
-n=16384,bf16,0.2649,0.0438,0.0441,6.05,6.01,1.0062,0.0477,5.55
-n=16384,fp32,0.2955,0.0850,0.0850,3.48,3.48,0.9996,0.0865,3.42
-n=16384,int8,0.2554,0.0259,0.0245,9.85,10.44,0.9432,0.0494,5.17
-n=17408,fp16,0.2982,0.0467,0.0469,6.38,6.36,1.0035,0.0512,5.82
-n=17408,bf16,0.2983,0.0463,0.0465,6.45,6.41,1.0053,0.0513,5.81
-n=17408,fp32,0.3294,0.0892,0.0892,3.69,3.69,0.9993,0.0907,3.63
-n=17408,int8,0.2827,0.0270,0.0261,10.48,10.84,0.9668,0.0525,5.39
-n=18432,fp16,0.3048,0.0492,0.0494,6.20,6.17,1.0036,0.0533,5.71
-n=18432,bf16,0.3046,0.0494,0.0498,6.17,6.12,1.0081,0.0534,5.70
-n=18432,fp32,0.3369,0.0939,0.0942,3.59,3.58,1.0025,0.0959,3.51
-n=18432,int8,0.2940,0.0282,0.0275,10.44,10.71,0.9749,0.0552,5.33
-n=19456,fp16,0.3245,0.0517,0.0519,6.28,6.25,1.0047,0.0559,5.80
-n=19456,bf16,0.3244,0.0518,0.0518,6.26,6.26,1.0005,0.0559,5.80
-n=19456,fp32,0.3586,0.0992,0.0992,3.62,3.61,1.0003,0.1008,3.56
-n=19456,int8,0.3062,0.0295,0.0288,10.37,10.62,0.9762,0.0580,5.28
-n=20480,fp16,0.3574,0.0541,0.0541,6.61,6.61,1.0001,0.0595,6.01
-n=20480,bf16,0.3569,0.0546,0.0544,6.54,6.56,0.9977,0.0595,6.00
-n=20480,fp32,0.3883,0.1038,0.1040,3.74,3.74,1.0012,0.1053,3.69
-n=20480,int8,0.3460,0.0306,0.0301,11.31,11.51,0.9828,0.0610,5.68
+params,dtype,torch_ms,triton_ms,cutile_ms,speedup_triton,speedup_cutile,triton_vs_cutile,tilelang_ms,speedup_tilelang,torch_nki_ms,nki_ms,speedup_nki
+n=1024,fp16,0.0206,0.0040,0.0056,5.11,3.70,1.3809,0.0050,4.10,0.0634,0.0439,1.44
+n=1024,bf16,0.0204,0.0039,0.0055,5.24,3.72,1.4083,0.0052,3.93,0.0633,0.0487,1.30
+n=1024,fp32,0.0213,0.0076,0.0076,2.81,2.81,1.0011,0.0080,2.65,0.1010,0.1087,0.93
+n=1024,int8,0.0193,0.0035,0.0035,5.52,5.49,1.0038,0.0056,3.46,0.1961,0.1258,1.56
+n=2048,fp16,0.0369,0.0079,0.0079,4.64,4.69,0.9897,0.0106,3.47,0.1283,0.0781,1.64
+n=2048,bf16,0.0371,0.0079,0.0079,4.67,4.71,0.9920,0.0099,3.76,0.1285,0.0875,1.47
+n=2048,fp32,0.0402,0.0131,0.0131,3.06,3.07,0.9970,0.0140,2.86,0.2083,0.1933,1.08
+n=2048,int8,0.0346,0.0061,0.0054,5.64,6.41,0.8789,0.0084,4.10,0.3537,0.2344,1.51
+n=3072,fp16,0.0558,0.0103,0.0105,5.43,5.30,1.0247,0.0129,4.32,0.1588,0.1076,1.48
+n=3072,bf16,0.0558,0.0104,0.0107,5.38,5.23,1.0302,0.0128,4.34,0.1588,0.1192,1.33
+n=3072,fp32,0.0608,0.0186,0.0187,3.27,3.24,1.0094,0.0189,3.21,0.2829,0.2867,0.99
+n=3072,int8,0.0528,0.0076,0.0069,6.91,7.61,0.9076,0.0120,4.39,0.5307,0.3423,1.55
+n=4096,fp16,0.0700,0.0132,0.0142,5.31,4.92,1.0789,0.0160,4.38,0.2208,0.1381,1.60
+n=4096,bf16,0.0699,0.0133,0.0138,5.26,5.07,1.0377,0.0161,4.35,0.2206,0.1568,1.41
+n=4096,fp32,0.0774,0.0238,0.0237,3.25,3.26,0.9964,0.0253,3.06,0.4044,0.3849,1.05
+n=4096,int8,0.0664,0.0079,0.0080,8.41,8.32,1.0112,0.0150,4.42,0.6783,0.4520,1.50
+n=5120,fp16,0.0920,0.0168,0.0166,5.48,5.54,0.9906,0.0187,4.92,0.2528,0.1677,1.51
+n=5120,bf16,0.0919,0.0166,0.0159,5.54,5.77,0.9608,0.0187,4.92,0.2530,0.1684,1.50
+n=5120,fp32,0.1022,0.0287,0.0283,3.56,3.61,0.9837,0.0291,3.52,0.4544,0.4684,0.97
+n=5120,int8,0.0840,0.0104,0.0110,8.06,7.64,1.0551,0.0184,4.57,0.8717,0.5634,1.55
+n=6144,fp16,0.1068,0.0185,0.0188,5.77,5.69,1.0133,0.0214,5.00,0.3197,0.2010,1.59
+n=6144,bf16,0.1067,0.0183,0.0188,5.82,5.67,1.0259,0.0213,5.01,0.3197,0.2008,1.59
+n=6144,fp32,0.1183,0.0336,0.0335,3.53,3.53,1.0000,0.0345,3.42,0.5592,0.5622,0.99
+n=6144,int8,0.1019,0.0122,0.0116,8.34,8.76,0.9520,0.0208,4.90,1.0224,0.6728,1.52
+n=7168,fp16,0.1230,0.0211,0.0212,5.83,5.80,1.0047,0.0240,5.12,0.3596,0.2310,1.56
+n=7168,bf16,0.1232,0.0211,0.0213,5.83,5.79,1.0064,0.0239,5.15,0.3596,0.2450,1.47
+n=7168,fp32,0.1339,0.0392,0.0392,3.42,3.41,1.0007,0.0395,3.39,0.6266,0.6610,0.95
+n=7168,int8,0.1201,0.0136,0.0132,8.82,9.10,0.9694,0.0238,5.04,1.1729,0.7809,1.50
+n=8192,fp16,0.1350,0.0239,0.0241,5.64,5.60,1.0073,0.0267,5.05,0.4115,0.2679,1.54
+n=8192,bf16,0.1352,0.0238,0.0239,5.67,5.65,1.0051,0.0268,5.04,0.4114,0.2874,1.43
+n=8192,fp32,0.1502,0.0436,0.0436,3.45,3.45,0.9990,0.0445,3.37,0.7204,0.7429,0.97
+n=8192,int8,0.1291,0.0152,0.0143,8.49,9.02,0.9416,0.0265,4.87,1.3309,0.8897,1.50
+n=9216,fp16,0.1553,0.0263,0.0265,5.90,5.86,1.0061,0.0296,5.25,0.4479,0.2938,1.52
+n=9216,bf16,0.1554,0.0267,0.0269,5.82,5.77,1.0083,0.0295,5.27,0.4476,0.3229,1.39
+n=9216,fp32,0.1712,0.0488,0.0494,3.51,3.47,1.0125,0.0501,3.41,0.8095,0.8411,0.96
+n=9216,int8,0.1497,0.0161,0.0158,9.29,9.49,0.9794,0.0296,5.05,1.5220,0.9964,1.53
+n=10240,fp16,0.1815,0.0285,0.0289,6.37,6.28,1.0146,0.0324,5.60,0.4906,0.3276,1.50
+n=10240,bf16,0.1813,0.0286,0.0290,6.33,6.25,1.0134,0.0324,5.60,0.4904,0.3256,1.51
+n=10240,fp32,0.1986,0.0537,0.0538,3.70,3.69,1.0010,0.0543,3.66,0.8962,0.9237,0.97
+n=10240,int8,0.1728,0.0177,0.0171,9.78,10.10,0.9682,0.0325,5.31,1.7084,1.1069,1.54
+n=11264,fp16,0.1866,0.0318,0.0319,5.86,5.85,1.0029,0.0350,5.33,0.5632,0.3531,1.60
+n=11264,bf16,0.1865,0.0324,0.0321,5.75,5.80,0.9918,0.0349,5.34,0.5632,0.3515,1.60
+n=11264,fp32,0.2097,0.0586,0.0596,3.58,3.52,1.0160,0.0604,3.47,0.9936,1.0389,0.96
+n=11264,int8,0.1795,0.0195,0.0185,9.18,9.72,0.9445,0.0352,5.10,1.9401,1.2146,1.60
+n=12288,fp16,0.2105,0.0352,0.0346,5.98,6.08,0.9838,0.0376,5.60,0.6072,0.3913,1.55
+n=12288,bf16,0.2106,0.0345,0.0346,6.10,6.09,1.0010,0.0374,5.63,0.6074,0.3890,1.56
+n=12288,fp32,0.2314,0.0639,0.0640,3.62,3.62,1.0011,0.0651,3.56,1.1184,1.1013,1.02
+n=12288,int8,0.2002,0.0206,0.0198,9.72,10.13,0.9592,0.0382,5.25,2.0346,1.3238,1.54
+n=13312,fp16,0.2213,0.0372,0.0367,5.94,6.04,0.9844,0.0400,5.53,0.6649,0.4149,1.60
+n=13312,bf16,0.2218,0.0364,0.0372,6.09,5.96,1.0218,0.0403,5.51,0.6650,0.4506,1.48
+n=13312,fp32,0.2446,0.0689,0.0689,3.55,3.55,0.9993,0.0702,3.48,1.1597,1.2267,0.95
+n=13312,int8,0.2119,0.0220,0.0211,9.63,10.06,0.9573,0.0408,5.20,2.2620,1.4321,1.58
+n=14336,fp16,0.2383,0.0390,0.0393,6.11,6.06,1.0076,0.0429,5.55,0.7059,0.4688,1.51
+n=14336,bf16,0.2383,0.0391,0.0392,6.10,6.08,1.0029,0.0427,5.58,0.7061,0.5133,1.38
+n=14336,fp32,0.2606,0.0737,0.0743,3.53,3.51,1.0080,0.0754,3.45,1.2495,1.3021,0.96
+n=14336,int8,0.2320,0.0228,0.0220,10.16,10.52,0.9660,0.0437,5.30,2.3416,1.5406,1.52
+n=15360,fp16,0.2614,0.0410,0.0413,6.37,6.32,1.0079,0.0457,5.72,0.7430,0.4722,1.57
+n=15360,bf16,0.2616,0.0414,0.0416,6.32,6.29,1.0044,0.0459,5.70,0.7431,0.5119,1.45
+n=15360,fp32,0.2898,0.0799,0.0802,3.63,3.61,1.0043,0.0812,3.57,1.3324,1.4188,0.94
+n=15360,int8,0.2462,0.0246,0.0238,10.03,10.34,0.9697,0.0467,5.27,2.5378,1.6487,1.54
+n=16384,fp16,0.2652,0.0435,0.0440,6.10,6.03,1.0122,0.0479,5.54,0.8053,0.5559,1.45
+n=16384,bf16,0.2649,0.0438,0.0441,6.05,6.01,1.0062,0.0477,5.55,0.8054,0.5580,1.44
+n=16384,fp32,0.2955,0.0850,0.0850,3.48,3.48,0.9996,0.0865,3.42,1.4400,1.4961,0.96
+n=16384,int8,0.2554,0.0259,0.0245,9.85,10.44,0.9432,0.0494,5.17,2.7062,1.7583,1.54
+n=17408,fp16,0.2982,0.0467,0.0469,6.38,6.36,1.0035,0.0512,5.82,0.8744,0.5406,1.62
+n=17408,bf16,0.2983,0.0463,0.0465,6.45,6.41,1.0053,0.0513,5.81,0.8744,0.5788,1.51
+n=17408,fp32,0.3294,0.0892,0.0892,3.69,3.69,0.9993,0.0907,3.63,1.5473,1.6105,0.96
+n=17408,int8,0.2827,0.0270,0.0261,10.48,10.84,0.9668,0.0525,5.39,2.9418,1.8653,1.58
+n=18432,fp16,0.3048,0.0492,0.0494,6.20,6.17,1.0036,0.0533,5.71,0.9112,0.5892,1.55
+n=18432,bf16,0.3046,0.0494,0.0498,6.17,6.12,1.0081,0.0534,5.70,0.9115,0.6047,1.51
+n=18432,fp32,0.3369,0.0939,0.0942,3.59,3.58,1.0025,0.0959,3.51,1.5889,1.6854,0.94
+n=18432,int8,0.2940,0.0282,0.0275,10.44,10.71,0.9749,0.0552,5.33,3.0265,1.9743,1.53
+n=19456,fp16,0.3245,0.0517,0.0519,6.28,6.25,1.0047,0.0559,5.80,0.9437,0.6167,1.53
+n=19456,bf16,0.3244,0.0518,0.0518,6.26,6.26,1.0005,0.0559,5.80,0.9438,0.6424,1.47
+n=19456,fp32,0.3586,0.0992,0.0992,3.62,3.61,1.0003,0.1008,3.56,1.7026,1.8261,0.93
+n=19456,int8,0.3062,0.0295,0.0288,10.37,10.62,0.9762,0.0580,5.28,3.3324,2.0818,1.60
+n=20480,fp16,0.3574,0.0541,0.0541,6.61,6.61,1.0001,0.0595,6.01,0.9852,0.6753,1.46
+n=20480,bf16,0.3569,0.0546,0.0544,6.54,6.56,0.9977,0.0595,6.00,0.9852,0.6699,1.47
+n=20480,fp32,0.3883,0.1038,0.1040,3.74,3.74,1.0012,0.1053,3.69,1.7611,1.8668,0.94
+n=20480,int8,0.3460,0.0306,0.0301,11.31,11.51,0.9828,0.0610,5.68,3.4587,2.1902,1.58