Re-implement matrix multiplication using a shared-memory tiled algorithm. Each thread block cooperatively loads tiles of A and B into shared memory, then accumulates partial dot products from those tiles. This reduces global-memory traffic by a factor of TILE_DIM compared to the naive version. The operation is the same GEMM as the naive variant: A has shape (M,K), B has shape (K,N), and the kernel writes C of shape (M,N) where:
C[i,j]=k=0∑K−1A[i,k]⋅B[k,j]All buffers use row-major layout. Each thread block computes a TILE_DIM×TILE_DIM output tile of C by sweeping across the shared dimension K in tile-sized chunks, loading one A tile and one B tile per chunk into shared memory and accumulating C[i,j] in a register. The grader checks correctness, but the naive matmul will be too slow for the larger shapes, so a tiled implementation is required.
Input: M = 2, K = 3, N = 2, A = [[1, 2, 3], [4, 5, 6]], B = [[7, 8], [9, 10], [11, 12]]
Output: C = [[58, 64], [139, 154]]
A single 16×16 block computes the entire output and the inner loop performs one tile iteration along K.
Input: M = 3, K = 2, N = 3, A = [[1, 2], [3, 4], [5, 6]], B = [[1, 2, 3], [4, 5, 6]]
Output: C = [[9, 12, 15], [19, 26, 33], [29, 40, 51]]
Rectangular product with partial tiles in all three dimensions; out-of-range tile loads must write 0 to shared memory.
In each tile step every thread in the block loads one element of the A tile and one element of the B tile from global memory into shared memory. The thread that owns output (i,j) is a natural choice for loading A[i, k0 + threadIdx.x] and B[k0 + threadIdx.y, j].
Two __syncthreads are needed: one after loading the tile so every thread can safely read it, and one after the inner accumulation so no thread overwrites the tile while another is still reading it.
For partial tiles, write 0 to shared memory when the global index is out of range. Skipping the write leaves stale values from the previous iteration in shared memory and corrupts the accumulation.
Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.
Re-implement matrix multiplication using a shared-memory tiled algorithm. Each thread block cooperatively loads tiles of A and B into shared memory, then accumulates partial dot products from those tiles. This reduces global-memory traffic by a factor of TILE_DIM compared to the naive version. The operation is the same GEMM as the naive variant: A has shape (M,K), B has shape (K,N), and the kernel writes C of shape (M,N) where:
C[i,j]=k=0∑K−1A[i,k]⋅B[k,j]All buffers use row-major layout. Each thread block computes a TILE_DIM×TILE_DIM output tile of C by sweeping across the shared dimension K in tile-sized chunks, loading one A tile and one B tile per chunk into shared memory and accumulating C[i,j] in a register. The grader checks correctness, but the naive matmul will be too slow for the larger shapes, so a tiled implementation is required.
Input: M = 2, K = 3, N = 2, A = [[1, 2, 3], [4, 5, 6]], B = [[7, 8], [9, 10], [11, 12]]
Output: C = [[58, 64], [139, 154]]
A single 16×16 block computes the entire output and the inner loop performs one tile iteration along K.
Input: M = 3, K = 2, N = 3, A = [[1, 2], [3, 4], [5, 6]], B = [[1, 2, 3], [4, 5, 6]]
Output: C = [[9, 12, 15], [19, 26, 33], [29, 40, 51]]
Rectangular product with partial tiles in all three dimensions; out-of-range tile loads must write 0 to shared memory.
In each tile step every thread in the block loads one element of the A tile and one element of the B tile from global memory into shared memory. The thread that owns output (i,j) is a natural choice for loading A[i, k0 + threadIdx.x] and B[k0 + threadIdx.y, j].
Two __syncthreads are needed: one after loading the tile so every thread can safely read it, and one after the inner accumulation so no thread overwrites the tile while another is still reading it.
For partial tiles, write 0 to shared memory when the global index is out of range. Skipping the write leaves stale values from the previous iteration in shared memory and corrupts the accumulation.
Accepts: array
Accepts: array
Submit a correct solution to see your kernel performance.