@triton.autotune(configs=get_autotune_config(), key=['N', 'C', 'D', 'H',
    'W', 'K', 'D_out', 'H_out', 'W_out', 'T', 'R', 'S', 'stride_d',
    'stride_h', 'stride_w', 'pad_d', 'pad_h', 'pad_w', 'dila_d', 'dila_h',
    'dila_w'])
@triton.jit
def conv3d_kernel(x_ptr, w_ptr, y_ptr, N, C, D, H, W, K, D_out, H_out,
    W_out, T, R, S, stride_d, stride_h, stride_w, pad_d, pad_h, pad_w,
    dila_d, dila_h, dila_w, GEMM_M, GEMM_N, GEMM_K, BLOCK_SIZE_M: tl.
    constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr,
    GROUP_SIZE_M: tl.constexpr):
    pid = tl.program_id(axis=0)
    num_pid_m = tl.cdiv(GEMM_M, BLOCK_SIZE_M)
    num_pid_n = tl.cdiv(GEMM_N, BLOCK_SIZE_N)
    num_pid_in_group = GROUP_SIZE_M * num_pid_n
    group_id = pid // num_pid_in_group
    first_pid_m = group_id * GROUP_SIZE_M
    group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M)
    pid_m = first_pid_m + pid % num_pid_in_group % group_size_m
    pid_n = pid % num_pid_in_group // group_size_m
    gemm_i = (pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M)) % GEMM_M
    gemm_j = (pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)) % GEMM_N
    n = gemm_i // (D_out * H_out * W_out)
    ndhw_residual = gemm_i % (D_out * H_out * W_out)
    d_out = ndhw_residual // (H_out * W_out)
    dhw_residual = ndhw_residual % (H_out * W_out)
    h_out = dhw_residual // W_out
    w_out = dhw_residual % W_out
    k = gemm_j
    accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
    for idx_k in range(0, tl.cdiv(GEMM_K, BLOCK_SIZE_K)):
        gemm_k = idx_k * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K)
        t = gemm_k // (R * S * C)
        trsc_residual = gemm_k % (R * S * C)
        r = trsc_residual // (S * C)
        rsc_residual = gemm_k % (S * C)
        s = rsc_residual // C
        c = rsc_residual % C
        d = d_out[:, None] * stride_d + t[None, :] * dila_d - pad_d
        h = h_out[:, None] * stride_h + r[None, :] * dila_h - pad_h
        w = w_out[:, None] * stride_w + s[None, :] * dila_w - pad_w
        mask_x = (d >= 0) & (d < D) & (h >= 0) & (h < H) & (w >= 0) & (w < W)
        mask_w = (t < T) & (r < R) & (s < S) & (c < C)
        offs_x = n[:, None
            ] * D * H * W * C + d * H * W * C + h * W * C + w * C + c
        offs_w = k[None, :] * T * R * S * C + t[:, None] * R * S * C + r[:,
            None] * S * C + s[:, None] * C + c[:, None]
        x_ptrs = x_ptr + offs_x
        w_ptrs = w_ptr + offs_w
        x_data = tl.load(x_ptrs, mask=mask_x, other=0.0)
        w_data = tl.load(w_ptrs, mask=mask_w[:, None], other=0.0)
        accumulator = tl.dot(x_data, w_data, accumulator)
    c_data = accumulator.to(tl.float16)
    offs_y = gemm_i[:, None] * GEMM_N + gemm_j[None, :]
    mask_y = (gemm_i[:, None] < GEMM_M) & (gemm_j[None, :] < GEMM_N)
    y_ptrs = y_ptr + offs_y
    tl.store(y_ptrs, c_data, mask=mask_y)
