@triton.jit
def _forward_kernel(c_ptr, s_ptr, u_ptr, col_stride, row_stride, **meta):
    n, n_tilde, dead_index, d_max, tournament_step, BLOCK_SIZE = meta['N'
        ], meta['N_TILDE'], meta['DEAD_INDEX'], meta['D_MAX'], meta['STEP'
        ], meta['BLOCK_SIZE']
    pid_x = tl.program_id(axis=0)
    temp = n_tilde - 1
    i = pid_x + tournament_step
    if pid_x == 0:
        i = 0
    if i >= n_tilde:
        i -= temp
    j = temp - pid_x + tournament_step
    if j >= n_tilde:
        j -= temp
    if i > j:
        i, j = j, i
    if (j == dead_index) | (j > d_max) & (i > d_max):
        return
    theta_offset = i * n - (i + 2) * (i + 1) // 2 + j
    c = tl.load(c_ptr + theta_offset)
    s = tl.load(s_ptr + theta_offset)
    offsets = tl.program_id(axis=1) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    output_offsets_i = i * row_stride + offsets * col_stride
    output_offsets_j = j * row_stride + offsets * col_stride
    maximum = n * row_stride + n * col_stride
    maski = output_offsets_i < maximum
    maskj = output_offsets_j < maximum
    ui = tl.load(u_ptr + output_offsets_i, mask=maski)
    uj = tl.load(u_ptr + output_offsets_j, mask=maskj)
    ioutput = ui * c - uj * s
    joutput = uj * c + ui * s
    ui = tl.store(u_ptr + output_offsets_i, ioutput, mask=maski)
    uj = tl.store(u_ptr + output_offsets_j, joutput, mask=maskj)
