Same operation as tiled matmul: given A of shape (M,K) and B of shape (K,N), compute C=A⋅B.
C[i,j]=k=0∑K−1A[i,k]⋅B[k,j]The difference is that the kernel is decorated with @triton.autotune over a small space of configs, and Triton picks the fastest config per shape at runtime. Autotune compiles every config on the first call for a given key signature, benchmarks each one, caches the winner, and reuses it for matching subsequent calls. The search space must stay small to fit inside the harness subprocess timeout.
Input: A = [[1, 2], [3, 4]], B = [[5, 6], [7, 8]]
Output: [[19, 22], [43, 50]]
Input: A = [[2, 0], [0, 2]], B = [[1, 2], [3, 4]]
Output: [[2, 4], [6, 8]]
Build each config with triton.Config specifying the meta-parameter dictionary, num_stages, and num_warps. Keep the list short; first-call compile time scales with the size of the space.
The key argument names the runtime values that uniquely determine the best config. For matmul those are M, N, and K; the kernel must accept them as regular int arguments so autotune can hash them.
Inside the kernel body, the autotuner injects the winning config's block sizes via the constexpr meta-parameters. The body is identical to the non-autotuned matmul.
torch.float32 on cuda.atol=1e-2 + rtol=1e-2 * abs(expected) for K≥128.Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.
Same operation as tiled matmul: given A of shape (M,K) and B of shape (K,N), compute C=A⋅B.
C[i,j]=k=0∑K−1A[i,k]⋅B[k,j]The difference is that the kernel is decorated with @triton.autotune over a small space of configs, and Triton picks the fastest config per shape at runtime. Autotune compiles every config on the first call for a given key signature, benchmarks each one, caches the winner, and reuses it for matching subsequent calls. The search space must stay small to fit inside the harness subprocess timeout.
Input: A = [[1, 2], [3, 4]], B = [[5, 6], [7, 8]]
Output: [[19, 22], [43, 50]]
Input: A = [[2, 0], [0, 2]], B = [[1, 2], [3, 4]]
Output: [[2, 4], [6, 8]]
Build each config with triton.Config specifying the meta-parameter dictionary, num_stages, and num_warps. Keep the list short; first-call compile time scales with the size of the space.
The key argument names the runtime values that uniquely determine the best config. For matmul those are M, N, and K; the kernel must accept them as regular int arguments so autotune can hash them.
Inside the kernel body, the autotuner injects the winning config's block sizes via the constexpr meta-parameters. The body is identical to the non-autotuned matmul.
torch.float32 on cuda.atol=1e-2 + rtol=1e-2 * abs(expected) for K≥128.Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.