@triton.jit
def copy_sparse_to_dense_grad_score_kernel(input_ptr, grad_output_ptr,
    grad_scores_ptr, sparse_rows_ptr, num_columns: tl.constexpr,
    num_experts_per_token: tl.constexpr, block_size: tl.constexpr):
    dense_row = tl.program_id(0)
    top_index = tl.program_id(1)
    sparse_row = tl.load(sparse_rows_ptr + dense_row *
        num_experts_per_token + top_index)
    grad_output_ptr += dense_row * num_columns
    input_ptr += sparse_row * num_columns
    offsets = tl.arange(0, block_size)
    if num_columns % block_size == 0:
        grad_scores = tl.load(input_ptr + offsets).to(tl.float32) * tl.load(
            grad_output_ptr + offsets).to(tl.float32)
    else:
        mask = offsets < num_columns
        grad_scores = tl.load(input_ptr + offsets, mask=mask).to(tl.float32
            ) * tl.load(grad_output_ptr + offsets, mask=mask).to(tl.float32)
    for i in range(1, tl.cdiv(num_columns, block_size)):
        offsets += block_size
        if num_columns % block_size == 0:
            grad_scores += tl.load(input_ptr + offsets).to(tl.float32
                ) * tl.load(grad_output_ptr + offsets).to(tl.float32)
        else:
            mask = offsets < num_columns
            grad_scores += tl.load(input_ptr + offsets, mask=mask).to(tl.
                float32) * tl.load(grad_output_ptr + offsets, mask=mask).to(tl
                .float32)
    tl.store(grad_scores_ptr + dense_row * num_experts_per_token +
        top_index, tl.sum(grad_scores))
