|
| 1 | +import torch |
| 2 | +import triton |
| 3 | +import triton.language as tl |
| 4 | +import triton.language.extra.cann.extension as al |
| 5 | + |
| 6 | + |
| 7 | +@triton.jit |
| 8 | +def conv2d_kernel( |
| 9 | + input_ptr, |
| 10 | + weight_ptr, |
| 11 | + bias_ptr, |
| 12 | + output_ptr, |
| 13 | + N: tl.constexpr, |
| 14 | + C_in: tl.constexpr, |
| 15 | + H_in: tl.constexpr, |
| 16 | + W_in: tl.constexpr, |
| 17 | + C_out: tl.constexpr, |
| 18 | + H_out: tl.constexpr, |
| 19 | + W_out: tl.constexpr, |
| 20 | + K_h: tl.constexpr, |
| 21 | + K_w: tl.constexpr, |
| 22 | + stride_h: tl.constexpr, |
| 23 | + stride_w: tl.constexpr, |
| 24 | + padding_h: tl.constexpr, |
| 25 | + padding_w: tl.constexpr, |
| 26 | + groups: tl.constexpr, |
| 27 | +): |
| 28 | + # Load input: (N, C_in, H_in, W_in) |
| 29 | + n_offs = tl.arange(0, N)[:, None, None, None] |
| 30 | + c_offs = tl.arange(0, C_in)[None, :, None, None] |
| 31 | + h_offs = tl.arange(0, H_in)[None, None, :, None] |
| 32 | + w_offs = tl.arange(0, W_in)[None, None, None, :] |
| 33 | + input_tile = tl.load(input_ptr + n_offs * (C_in * H_in * W_in) + c_offs * (H_in * W_in) + h_offs * W_in + w_offs) |
| 34 | + |
| 35 | + # Load weight: (C_out, C_in // groups, K_h, K_w) |
| 36 | + co_offs = tl.arange(0, C_out)[:, None, None, None] |
| 37 | + ci_offs = tl.arange(0, C_in // groups)[None, :, None, None] |
| 38 | + kh_offs = tl.arange(0, K_h)[None, None, :, None] |
| 39 | + kw_offs = tl.arange(0, K_w)[None, None, None, :] |
| 40 | + weight_tile = tl.load(weight_ptr + co_offs * ((C_in // groups) * K_h * K_w) + ci_offs * (K_h * K_w) + |
| 41 | + kh_offs * K_w + kw_offs) |
| 42 | + |
| 43 | + # Load bias: (C_out,) |
| 44 | + bias_tile = tl.load(bias_ptr + tl.arange(0, C_out)) |
| 45 | + |
| 46 | + output = al.conv2d( |
| 47 | + input_tile, |
| 48 | + weight_tile, |
| 49 | + bias_tile, |
| 50 | + groups=groups, |
| 51 | + padding=(padding_h, padding_w), |
| 52 | + stride=(stride_h, stride_w), |
| 53 | + dilation=1, |
| 54 | + ) |
| 55 | + |
| 56 | + # Store output: (N, C_out, H_out, W_out) |
| 57 | + no_offs = tl.arange(0, N)[:, None, None, None] |
| 58 | + co_offs = tl.arange(0, C_out)[None, :, None, None] |
| 59 | + ho_offs = tl.arange(0, H_out)[None, None, :, None] |
| 60 | + wo_offs = tl.arange(0, W_out)[None, None, None, :] |
| 61 | + tl.store(output_ptr + no_offs * (C_out * H_out * W_out) + co_offs * (H_out * W_out) + ho_offs * W_out + wo_offs, |
| 62 | + output) |
| 63 | + |
| 64 | + |
| 65 | +def test_conv2d(): |
| 66 | + N, C_in, H_in, W_in = 2, 16, 32, 32 |
| 67 | + C_out, K_h, K_w = 32, 3, 3 |
| 68 | + stride = (1, 1) |
| 69 | + padding = (1, 1) |
| 70 | + groups = 1 |
| 71 | + H_out = (H_in + 2 * padding[0] - (K_h - 1) - 1) // stride[0] + 1 |
| 72 | + W_out = (W_in + 2 * padding[1] - (K_w - 1) - 1) // stride[1] + 1 |
| 73 | + |
| 74 | + x = torch.randn(N, C_in, H_in, W_in, dtype=torch.float16) |
| 75 | + w = torch.randn(C_out, C_in // groups, K_h, K_w, dtype=torch.float16) |
| 76 | + b = torch.randn(C_out, dtype=torch.float16) |
| 77 | + |
| 78 | + x_npu = x.npu() |
| 79 | + w_npu = w.npu() |
| 80 | + b_npu = b.npu() |
| 81 | + out_npu = torch.empty(N, C_out, H_out, W_out, dtype=torch.float16).npu() |
| 82 | + |
| 83 | + conv2d_kernel[(1, )](x_npu, w_npu, b_npu, out_npu, N=N, C_in=C_in, H_in=H_in, W_in=W_in, C_out=C_out, H_out=H_out, |
| 84 | + W_out=W_out, K_h=K_h, K_w=K_w, stride_h=stride[0], stride_w=stride[1], padding_h=padding[0], |
| 85 | + padding_w=padding[1], groups=groups) |
| 86 | + |
| 87 | + gold = torch.nn.functional.conv2d(x, w, b, stride=stride, padding=padding, groups=groups) |
| 88 | + torch.testing.assert_close(out_npu.cpu(), gold, rtol=1e-2, atol=1e-2) |
| 89 | + |
| 90 | + |
| 91 | +if __name__ == "__main__": |
| 92 | + test_conv2d() |
0 commit comments