@triton.jit
def _ld_mse_bwd_kernel(grad_ptr, input_ptr, target_ptr, grad_output,
    grad_row_stride, input_row_stride, target_row_stride, n_rows, n_cols,
    BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    col_offsets = tl.arange(0, BLOCK_SIZE)
    mask = col_offsets < n_cols
    grad_ptrs = grad_ptr + pid * grad_row_stride + col_offsets
    input_ptrs = input_ptr + pid * input_row_stride + col_offsets
    target_ptrs = target_ptr + pid * target_row_stride + col_offsets
    input = tl.load(input_ptrs, mask=mask, other=0.0)
    target = tl.load(target_ptrs, mask=mask, other=0.0)
    grad_ = (input - target) * 2 * grad_output / (n_rows * n_cols)
    tl.store(grad_ptrs, grad_, mask=mask)
