nki(interleave): NKI (Trainium) implementation

results/B200/csv/interleave_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=1000000,fp16,0.0088,0.0026,0.0027,3.38,3.28,1.0385,0.0028,3.11
-n=1000000,bf16,0.0087,0.0025,0.0026,3.42,3.32,1.0400,0.0026,3.34
-n=1000000,fp32,0.0090,0.0035,0.0038,2.57,2.34,1.0857,0.0036,2.49
-n=1000000,int8,0.0083,0.0021,0.0022,3.94,3.77,1.0476,0.0024,3.44
-n=2000000,fp16,0.0127,0.0037,0.0038,3.46,3.31,1.0270,0.0037,3.44
-n=2000000,bf16,0.0124,0.0034,0.0037,3.64,3.36,1.0882,0.0040,3.06
-n=2000000,fp32,0.0146,0.0064,0.0070,2.29,2.07,1.0938,0.0067,2.19
-n=2000000,int8,0.0117,0.0026,0.0027,4.56,4.26,1.0385,0.0033,3.60
-n=3000000,fp16,0.0171,0.0053,0.0057,3.24,3.00,1.0755,0.0054,3.16
-n=3000000,bf16,0.0167,0.0051,0.0054,3.26,3.11,1.0588,0.0056,2.99
-n=3000000,fp32,0.0219,0.0098,0.0098,2.24,2.25,1.0000,0.0096,2.28
-n=3000000,int8,0.0143,0.0029,0.0035,4.99,4.12,1.2069,0.0035,4.08
-n=4000000,fp16,0.0217,0.0069,0.0075,3.15,2.87,1.0870,0.0072,3.03
-n=4000000,bf16,0.0216,0.0066,0.0067,3.29,3.25,1.0152,0.0070,3.09
-n=4000000,fp32,0.0299,0.0117,0.0124,2.55,2.42,1.0598,0.0122,2.46
-n=4000000,int8,0.0181,0.0035,0.0037,5.18,4.91,1.0571,0.0042,4.35
-n=5000000,fp16,0.0266,0.0081,0.0085,3.28,3.15,1.0494,0.0087,3.04
-n=5000000,bf16,0.0266,0.0080,0.0081,3.32,3.27,1.0125,0.0086,3.09
-n=5000000,fp32,0.0373,0.0145,0.0150,2.57,2.49,1.0345,0.0151,2.46
-n=5000000,int8,0.0211,0.0056,0.0041,3.74,5.11,0.7321,0.0046,4.54
-n=6000000,fp16,0.0318,0.0103,0.0096,3.09,3.31,0.9320,0.0105,3.04
-n=6000000,bf16,0.0320,0.0095,0.0106,3.39,3.04,1.1158,0.0105,3.05
-n=6000000,fp32,0.0429,0.0170,0.0172,2.52,2.49,1.0118,0.0171,2.50
-n=6000000,int8,0.0256,0.0049,0.0052,5.24,4.95,1.0612,0.0053,4.79
-n=7000000,fp16,0.0369,0.0108,0.0108,3.41,3.42,1.0000,0.0110,3.35
-n=7000000,bf16,0.0374,0.0109,0.0110,3.44,3.40,1.0092,0.0111,3.38
-n=7000000,fp32,0.0493,0.0194,0.0194,2.54,2.54,1.0000,0.0193,2.56
-n=7000000,int8,0.0303,0.0056,0.0064,5.42,4.77,1.1429,0.0061,4.95
-n=8000000,fp16,0.0421,0.0123,0.0124,3.43,3.40,1.0081,0.0132,3.20
-n=8000000,bf16,0.0425,0.0132,0.0124,3.23,3.43,0.9394,0.0127,3.35
-n=8000000,fp32,0.0557,0.0216,0.0216,2.57,2.58,1.0000,0.0219,2.55
-n=8000000,int8,0.0348,0.0064,0.0067,5.44,5.18,1.0469,0.0074,4.72
-n=9000000,fp16,0.0475,0.0134,0.0137,3.53,3.47,1.0224,0.0146,3.24
-n=9000000,bf16,0.0470,0.0133,0.0137,3.54,3.44,1.0301,0.0137,3.44
-n=9000000,fp32,0.0621,0.0241,0.0242,2.58,2.57,1.0041,0.0246,2.53
-n=9000000,int8,0.0393,0.0071,0.0075,5.50,5.22,1.0563,0.0084,4.68
-n=10000000,fp16,0.0518,0.0145,0.0146,3.56,3.55,1.0069,0.0149,3.47
-n=10000000,bf16,0.0515,0.0145,0.0145,3.55,3.55,1.0000,0.0150,3.44
-n=10000000,fp32,0.0694,0.0263,0.0265,2.65,2.62,1.0076,0.0266,2.61
-n=10000000,int8,0.0442,0.0077,0.0081,5.72,5.43,1.0519,0.0086,5.13
-n=11000000,fp16,0.0560,0.0170,0.0160,3.30,3.51,0.9412,0.0173,3.24
-n=11000000,bf16,0.0561,0.0170,0.0159,3.29,3.53,0.9353,0.0164,3.42
-n=11000000,fp32,0.0796,0.0287,0.0287,2.77,2.78,1.0000,0.0287,2.77
-n=11000000,int8,0.0493,0.0088,0.0090,5.61,5.48,1.0227,0.0095,5.19
-n=12000000,fp16,0.0607,0.0170,0.0173,3.57,3.52,1.0176,0.0170,3.56
-n=12000000,bf16,0.0608,0.0169,0.0171,3.59,3.56,1.0118,0.0171,3.55
-n=12000000,fp32,0.0882,0.0310,0.0310,2.84,2.85,1.0000,0.0314,2.81
-n=12000000,int8,0.0540,0.0091,0.0105,5.96,5.16,1.1538,0.0105,5.16
-n=13000000,fp16,0.0654,0.0182,0.0182,3.59,3.59,1.0000,0.0184,3.56
-n=13000000,bf16,0.0652,0.0183,0.0184,3.57,3.54,1.0055,0.0184,3.54
-n=13000000,fp32,0.0952,0.0330,0.0333,2.88,2.86,1.0091,0.0335,2.84
-n=13000000,int8,0.0582,0.0104,0.0104,5.62,5.58,1.0000,0.0112,5.20
-n=14000000,fp16,0.0697,0.0193,0.0194,3.62,3.60,1.0052,0.0194,3.60
-n=14000000,bf16,0.0699,0.0193,0.0195,3.62,3.59,1.0104,0.0196,3.57
-n=14000000,fp32,0.1015,0.0354,0.0357,2.87,2.84,1.0085,0.0358,2.83
-n=14000000,int8,0.0632,0.0109,0.0113,5.80,5.59,1.0367,0.0113,5.59
-n=15000000,fp16,0.0739,0.0205,0.0207,3.60,3.58,1.0098,0.0204,3.62
-n=15000000,bf16,0.0743,0.0205,0.0206,3.63,3.61,1.0049,0.0206,3.61
-n=15000000,fp32,0.1080,0.0376,0.0380,2.87,2.84,1.0106,0.0378,2.86
-n=15000000,int8,0.0675,0.0125,0.0128,5.41,5.27,1.0240,0.0125,5.42
-n=16000000,fp16,0.0787,0.0218,0.0219,3.61,3.59,1.0046,0.0218,3.61
-n=16000000,bf16,0.0789,0.0218,0.0218,3.62,3.62,1.0000,0.0236,3.34
-n=16000000,fp32,0.1146,0.0399,0.0403,2.87,2.84,1.0100,0.0399,2.87
-n=16000000,int8,0.0716,0.0123,0.0124,5.82,5.76,1.0081,0.0131,5.46
-n=17000000,fp16,0.0835,0.0228,0.0229,3.66,3.65,1.0044,0.0229,3.64
-n=17000000,bf16,0.0835,0.0230,0.0231,3.64,3.62,1.0043,0.0234,3.56
-n=17000000,fp32,0.1217,0.0421,0.0426,2.89,2.86,1.0119,0.0432,2.82
-n=17000000,int8,0.0758,0.0129,0.0130,5.88,5.83,1.0078,0.0131,5.77
-n=18000000,fp16,0.0881,0.0241,0.0242,3.66,3.64,1.0041,0.0241,3.66
-n=18000000,bf16,0.0877,0.0239,0.0240,3.68,3.65,1.0042,0.0240,3.65
-n=18000000,fp32,0.1277,0.0443,0.0448,2.88,2.85,1.0113,0.0456,2.80
-n=18000000,int8,0.0799,0.0143,0.0146,5.58,5.45,1.0210,0.0141,5.68
-n=19000000,fp16,0.0931,0.0251,0.0252,3.70,3.69,1.0040,0.0251,3.71
-n=19000000,bf16,0.0930,0.0250,0.0251,3.73,3.70,1.0040,0.0251,3.71
-n=19000000,fp32,0.1344,0.0469,0.0473,2.87,2.84,1.0085,0.0475,2.83
-n=19000000,int8,0.0842,0.0149,0.0142,5.65,5.92,0.9530,0.0141,5.96
-n=20000000,fp16,0.0985,0.0263,0.0264,3.75,3.73,1.0038,0.0263,3.74
-n=20000000,bf16,0.0983,0.0262,0.0263,3.75,3.74,1.0038,0.0265,3.71
-n=20000000,fp32,0.1404,0.0489,0.0495,2.87,2.84,1.0123,0.0499,2.81
-n=20000000,int8,0.0882,0.0146,0.0148,6.05,5.94,1.0137,0.0147,5.99
+params,dtype,torch_ms,triton_ms,cutile_ms,speedup_triton,speedup_cutile,triton_vs_cutile,torch_nki_ms,nki_ms,speedup_nki
+n=1000000,fp16,0.0088,0.0026,0.0027,3.38,3.28,1.0385,0.0755,0.0277,2.73
+n=1000000,bf16,0.0087,0.0025,0.0026,3.42,3.32,1.0400,0.0754,0.0286,2.64
+n=1000000,fp32,0.0090,0.0035,0.0038,2.57,2.34,1.0857,0.1205,0.0382,3.16
+n=1000000,int8,0.0083,0.0021,0.0022,3.94,3.77,1.0476,0.0525,0.0236,2.22
+n=2000000,fp16,0.0127,0.0037,0.0038,3.46,3.31,1.0270,0.0777,0.0388,2.00
+n=2000000,bf16,0.0124,0.0034,0.0037,3.64,3.36,1.0882,0.0776,0.0389,1.99
+n=2000000,fp32,0.0146,0.0064,0.0070,2.29,2.07,1.0938,0.0922,0.0606,1.52
+n=2000000,int8,0.0117,0.0026,0.0027,4.56,4.26,1.0385,0.0723,0.0290,2.49
+n=3000000,fp16,0.0171,0.0053,0.0057,3.24,3.00,1.0755,0.1936,0.0503,3.85
+n=3000000,bf16,0.0167,0.0051,0.0054,3.26,3.11,1.0588,0.1940,0.0506,3.83
+n=3000000,fp32,0.0219,0.0098,0.0098,2.24,2.25,1.0000,0.3324,0.0844,3.94
+n=3000000,int8,0.0143,0.0029,0.0035,4.99,4.12,1.2069,0.1588,0.0343,4.64
+n=4000000,fp16,0.0217,0.0069,0.0075,3.15,2.87,1.0870,0.1326,0.0605,2.19
+n=4000000,bf16,0.0216,0.0066,0.0067,3.29,3.25,1.0152,0.1327,0.0610,2.17
+n=4000000,fp32,0.0299,0.0117,0.0124,2.55,2.42,1.0598,0.1752,0.1063,1.65
+n=4000000,int8,0.0181,0.0035,0.0037,5.18,4.91,1.0571,0.1303,0.0373,3.49
+n=5000000,fp16,0.0266,0.0081,0.0085,3.28,3.15,1.0494,0.3018,0.0736,4.10
+n=5000000,bf16,0.0266,0.0080,0.0081,3.32,3.27,1.0125,0.3020,0.0766,3.94
+n=5000000,fp32,0.0373,0.0145,0.0150,2.57,2.49,1.0345,0.5474,0.1340,4.08
+n=5000000,int8,0.0211,0.0056,0.0041,3.74,5.11,0.7321,0.2655,0.0461,5.76
+n=6000000,fp16,0.0318,0.0103,0.0096,3.09,3.31,0.9320,0.1850,0.0847,2.18
+n=6000000,bf16,0.0320,0.0095,0.0106,3.39,3.04,1.1158,0.1849,0.0846,2.18
+n=6000000,fp32,0.0429,0.0170,0.0172,2.52,2.49,1.0118,0.2369,0.1581,1.50
+n=6000000,int8,0.0256,0.0049,0.0052,5.24,4.95,1.0612,0.1675,0.0506,3.31
+n=7000000,fp16,0.0369,0.0108,0.0108,3.41,3.42,1.0000,0.4157,0.0965,4.31
+n=7000000,bf16,0.0374,0.0109,0.0110,3.44,3.40,1.0092,0.4161,0.1019,4.08
+n=7000000,fp32,0.0493,0.0194,0.0194,2.54,2.54,1.0000,0.7504,0.1852,4.05
+n=7000000,int8,0.0303,0.0056,0.0064,5.42,4.77,1.1429,0.3461,0.0569,6.08
+n=8000000,fp16,0.0421,0.0123,0.0124,3.43,3.40,1.0081,0.2431,0.1099,2.21
+n=8000000,bf16,0.0425,0.0132,0.0124,3.23,3.43,0.9394,0.2430,0.1127,2.16
+n=8000000,fp32,0.0557,0.0216,0.0216,2.57,2.58,1.0000,0.3154,0.2075,1.52
+n=8000000,int8,0.0348,0.0064,0.0067,5.44,5.18,1.0469,0.2300,0.0605,3.80
+n=9000000,fp16,0.0475,0.0134,0.0137,3.53,3.47,1.0224,0.5377,0.1221,4.40
+n=9000000,bf16,0.0470,0.0133,0.0137,3.54,3.44,1.0301,0.5376,0.1310,4.10
+n=9000000,fp32,0.0621,0.0241,0.0242,2.58,2.57,1.0041,0.9769,0.2307,4.24
+n=9000000,int8,0.0393,0.0071,0.0075,5.50,5.22,1.0563,0.4223,0.0677,6.24
+n=10000000,fp16,0.0518,0.0145,0.0146,3.56,3.55,1.0069,0.2982,0.1343,2.22
+n=10000000,bf16,0.0515,0.0145,0.0145,3.55,3.55,1.0000,0.2983,0.1346,2.22
+n=10000000,fp32,0.0694,0.0263,0.0265,2.65,2.62,1.0076,0.4032,0.2610,1.54
+n=10000000,int8,0.0442,0.0077,0.0081,5.72,5.43,1.0519,0.2807,0.0735,3.82
+n=11000000,fp16,0.0560,0.0170,0.0160,3.30,3.51,0.9412,0.6431,0.1464,4.39
+n=11000000,bf16,0.0561,0.0170,0.0159,3.29,3.53,0.9353,0.6434,0.1464,4.39
+n=11000000,fp32,0.0796,0.0287,0.0287,2.77,2.78,1.0000,1.1971,0.2781,4.30
+n=11000000,int8,0.0493,0.0088,0.0090,5.61,5.48,1.0227,0.5630,0.0791,7.11
+n=12000000,fp16,0.0607,0.0170,0.0173,3.57,3.52,1.0176,0.3536,0.1580,2.24
+n=12000000,bf16,0.0608,0.0169,0.0171,3.59,3.56,1.0118,0.3536,0.1619,2.18
+n=12000000,fp32,0.0882,0.0310,0.0310,2.84,2.85,1.0000,0.4807,0.3038,1.58
+n=12000000,int8,0.0540,0.0091,0.0105,5.96,5.16,1.1538,0.3242,0.0834,3.89
+n=13000000,fp16,0.0654,0.0182,0.0182,3.59,3.59,1.0000,0.7870,0.1713,4.60
+n=13000000,bf16,0.0652,0.0183,0.0184,3.57,3.54,1.0055,0.7880,0.1793,4.39
+n=13000000,fp32,0.0952,0.0330,0.0333,2.88,2.86,1.0091,1.4637,0.3262,4.49
+n=13000000,int8,0.0582,0.0104,0.0104,5.62,5.58,1.0000,0.6391,0.0903,7.08
+n=14000000,fp16,0.0697,0.0193,0.0194,3.62,3.60,1.0052,0.4033,0.1851,2.18
+n=14000000,bf16,0.0699,0.0193,0.0195,3.62,3.59,1.0104,0.4034,0.1873,2.15
+n=14000000,fp32,0.1015,0.0354,0.0357,2.87,2.84,1.0085,0.5679,0.3532,1.61
+n=14000000,int8,0.0632,0.0109,0.0113,5.80,5.59,1.0367,0.3812,0.0966,3.95
+n=15000000,fp16,0.0739,0.0205,0.0207,3.60,3.58,1.0098,0.8741,0.1984,4.41
+n=15000000,bf16,0.0743,0.0205,0.0206,3.63,3.61,1.0049,0.8740,0.2005,4.36
+n=15000000,fp32,0.1080,0.0376,0.0380,2.87,2.84,1.0106,1.6289,0.3748,4.35
+n=15000000,int8,0.0675,0.0125,0.0128,5.41,5.27,1.0240,0.7064,0.1074,6.58
+n=16000000,fp16,0.0787,0.0218,0.0219,3.61,3.59,1.0046,0.4686,0.2074,2.26
+n=16000000,bf16,0.0789,0.0218,0.0218,3.62,3.62,1.0000,0.4686,0.2094,2.24
+n=16000000,fp32,0.1146,0.0399,0.0403,2.87,2.84,1.0100,0.6551,0.3915,1.67
+n=16000000,int8,0.0716,0.0123,0.0124,5.82,5.76,1.0081,0.4268,0.1126,3.79
+n=17000000,fp16,0.0835,0.0228,0.0229,3.66,3.65,1.0044,0.9881,0.2232,4.43
+n=17000000,bf16,0.0835,0.0230,0.0231,3.64,3.62,1.0043,0.9878,0.2261,4.37
+n=17000000,fp32,0.1217,0.0421,0.0426,2.89,2.86,1.0119,1.8419,0.4226,4.36
+n=17000000,int8,0.0758,0.0129,0.0130,5.88,5.83,1.0078,0.8645,0.1212,7.13
+n=18000000,fp16,0.0881,0.0241,0.0242,3.66,3.64,1.0041,0.5239,0.2348,2.23
+n=18000000,bf16,0.0877,0.0239,0.0240,3.68,3.65,1.0042,0.5239,0.2357,2.22
+n=18000000,fp32,0.1277,0.0443,0.0448,2.88,2.85,1.0113,0.6857,0.4461,1.54
+n=18000000,int8,0.0799,0.0143,0.0146,5.58,5.45,1.0210,0.4956,0.1267,3.91
+n=19000000,fp16,0.0931,0.0251,0.0252,3.70,3.69,1.0040,1.0995,0.2429,4.53
+n=19000000,bf16,0.0930,0.0250,0.0251,3.73,3.70,1.0040,1.0996,0.2462,4.47
+n=19000000,fp32,0.1344,0.0469,0.0473,2.87,2.84,1.0085,2.0164,0.4697,4.29
+n=19000000,int8,0.0842,0.0149,0.0142,5.65,5.92,0.9530,0.9226,0.1300,7.10
+n=20000000,fp16,0.0985,0.0263,0.0264,3.75,3.73,1.0038,0.5843,0.2567,2.28
+n=20000000,bf16,0.0983,0.0262,0.0263,3.75,3.74,1.0038,0.5843,0.2608,2.24
+n=20000000,fp32,0.1404,0.0489,0.0495,2.87,2.84,1.0123,0.8081,0.4942,1.64
+n=20000000,int8,0.0882,0.0146,0.0148,6.05,5.94,1.0137,0.5361,0.1315,4.08