@triton.autotune(configs=[triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 32}, num_stages=3, num_warps=8),
    triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=3, num_warps=8), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 32}, num_stages=3, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=3, num_warps=4), triton.Config({'BLOCK_SIZE_M': 256,
    'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 32}, num_stages=4, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N':
    32}, num_stages=4, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 32}, num_stages=4, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=4, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 32}, num_stages=4, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 
    32}, num_stages=4, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 32}, num_stages=3, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=3, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 32}, num_stages=3, num_warps=8),
    triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=3, num_warps=8), triton.Config({'BLOCK_SIZE_M': 256,
    'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 32}, num_stages=3, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N':
    32}, num_stages=3, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 32}, num_stages=3, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=3, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 32}, num_stages=3, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 
    32}, num_stages=3, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N': 32}, num_stages=2, num_warps=8),
    triton.Config({'BLOCK_SIZE_M': 256, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=2, num_warps=8), triton.Config({'BLOCK_SIZE_M': 256,
    'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 32}, num_stages=2, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 256, 'BLOCK_SIZE_N':
    32}, num_stages=2, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N': 32}, num_stages=2, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 128, 'BLOCK_SIZE_N':
    32}, num_stages=2, num_warps=4), triton.Config({'BLOCK_SIZE_M': 128,
    'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 32}, num_stages=2, num_warps=4),
    triton.Config({'BLOCK_SIZE_M': 64, 'BLOCK_SIZE_K': 64, 'BLOCK_SIZE_N': 
    32}, num_stages=2, num_warps=4)], key=['M', 'N', 'K'])
@triton.jit
def rz_linear_backward_input_grad_kernel_tf32(a_ptr, b_ptr, c_ptr,
    init_factor, M, N, K, H, stride_am, stride_an, stride_cm, stride_ck, R7:
    int, R6: int, R5: int, R4: int, R3: int, R2: int, R1: int, R0: int,
    BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K:
    tl.constexpr, GROUP_SIZE: tl.constexpr):
    rz_linear_backward_input_grad_core(a_ptr=a_ptr, b_ptr=b_ptr, c_ptr=
        c_ptr, init_factor=init_factor, M=M, N=N, K=K, H=H, stride_am=
        stride_am, stride_an=stride_an, stride_cm=stride_cm, stride_ck=
        stride_ck, R7=R7, R6=R6, R5=R5, R4=R4, R3=R3, R2=R2, R1=R1, R0=R0,
        allow_tf32=True, BLOCK_SIZE_M=BLOCK_SIZE_M, BLOCK_SIZE_N=
        BLOCK_SIZE_N, BLOCK_SIZE_K=BLOCK_SIZE_K, GROUP_SIZE=GROUP_SIZE)
