Basic tiled matmul. The agent improves this.
(
A_ptr, B_ptr, C_ptr,
M, N, K,
stride_am, stride_ak,
stride_bk, stride_bn,
stride_cm, stride_cn,
BLOCK_SIZE_M: tl.constexpr,
BLOCK_SIZE_N: tl.constexpr,
BLOCK_SIZE_K: tl.constexpr,
)
| 23 | |
| 24 | @triton.jit |
| 25 | def matmul_kernel( |
| 26 | A_ptr, B_ptr, C_ptr, |
| 27 | M, N, K, |
| 28 | stride_am, stride_ak, |
| 29 | stride_bk, stride_bn, |
| 30 | stride_cm, stride_cn, |
| 31 | BLOCK_SIZE_M: tl.constexpr, |
| 32 | BLOCK_SIZE_N: tl.constexpr, |
| 33 | BLOCK_SIZE_K: tl.constexpr, |
| 34 | ): |
| 35 | """Basic tiled matmul. The agent improves this.""" |
| 36 | pid_m = tl.program_id(0) |
| 37 | pid_n = tl.program_id(1) |
| 38 | |
| 39 | offs_m = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) |
| 40 | offs_n = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) |
| 41 | offs_k = tl.arange(0, BLOCK_SIZE_K) |
| 42 | |
| 43 | a_ptrs = A_ptr + offs_m[:, None] * stride_am + offs_k[None, :] * stride_ak |
| 44 | b_ptrs = B_ptr + offs_k[:, None] * stride_bk + offs_n[None, :] * stride_bn |
| 45 | |
| 46 | acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) |
| 47 | |
| 48 | for k in range(0, K, BLOCK_SIZE_K): |
| 49 | a = tl.load(a_ptrs, mask=(offs_m[:, None] < M) & (offs_k[None, :] < K), other=0.0) |
| 50 | b = tl.load(b_ptrs, mask=(offs_k[:, None] < K) & (offs_n[None, :] < N), other=0.0) |
| 51 | acc += tl.dot(a, b) |
| 52 | a_ptrs += BLOCK_SIZE_K * stride_ak |
| 53 | b_ptrs += BLOCK_SIZE_K * stride_bk |
| 54 | offs_k += BLOCK_SIZE_K |
| 55 | |
| 56 | c = acc.to(C_ptr.dtype.element_ty) |
| 57 | c_ptrs = C_ptr + offs_m[:, None] * stride_cm + offs_n[None, :] * stride_cn |
| 58 | mask = (offs_m[:, None] < M) & (offs_n[None, :] < N) |
| 59 | tl.store(c_ptrs, c, mask=mask) |
| 60 | |
| 61 | |
| 62 | def kernel_fn(A: torch.Tensor, B: torch.Tensor) -> torch.Tensor: |
nothing calls this directly
no outgoing calls
no test coverage detected