@triton.jit
def triton_cross_entropy_forward(input_ptr, input_stride, target_ptr,
    target_stride, loss_output_ptr, loss_output_stride, num_classes,
    num_valid_targets, ignore_label, smoothing_factor: tl.constexpr,
    reduction_mode: tl.constexpr, BLOCK_SIZE: tl.constexpr):
    row_id = tl.program_id(0).to(tl.int64)
    target_ptr += row_id * target_stride
    target_label = tl.load(target_ptr)
    input_ptr += row_id * input_stride
    if target_label == ignore_label:
        for i in range(0, num_classes, BLOCK_SIZE):
            input_offsets = i + tl.arange(0, BLOCK_SIZE)
            tl.store(input_ptr + input_offsets, 0.0, mask=input_offsets <
                num_classes)
        return
    loss_output_ptr += row_id * loss_output_stride
    max_val = float('-inf')
    normalization_factor = 0.0
    target_input_val = tl.load(input_ptr + target_label)
    smoothing_sum = 0.0
    epsilon = smoothing_factor / num_classes
    for i in range(0, num_classes, BLOCK_SIZE):
        input_offsets = i + tl.arange(0, BLOCK_SIZE)
        input_block = tl.load(input_ptr + input_offsets, mask=input_offsets <
            num_classes, other=float('-inf'))
        block_max = tl.max(input_block)
        if smoothing_factor > 0:
            smoothing_sum += tl.sum(tl.where(input_offsets < num_classes, -
                epsilon * input_block, 0.0))
        new_max = tl.maximum(max_val, block_max)
        normalization_factor = normalization_factor * tl.exp(max_val - new_max
            ) + tl.sum(tl.exp(input_block - new_max))
        max_val = new_max
    for i in range(0, num_classes, BLOCK_SIZE):
        input_offsets = i + tl.arange(0, BLOCK_SIZE)
        input_block = tl.load(input_ptr + input_offsets, mask=input_offsets <
            num_classes, other=float('-inf'))
        if reduction_mode == 'mean':
            input_block = (tl.exp(input_block - max_val) /
                normalization_factor - epsilon) / num_valid_targets
        else:
            input_block = tl.exp(input_block - max_val
                ) / normalization_factor - epsilon
        tl.store(input_ptr + input_offsets, input_block, mask=input_offsets <
            num_classes)
    tl.debug_barrier()
    row_loss = -(target_input_val - max_val - tl.log(normalization_factor))
    if smoothing_factor > 0:
        smooth_loss = smoothing_sum + smoothing_factor * (max_val + tl.log(
            normalization_factor))
        row_loss = row_loss * (1 - smoothing_factor) + smooth_loss
    if reduction_mode == 'mean':
        row_loss /= num_valid_targets
    updated_target_val = tl.load(input_ptr + target_label)
    if reduction_mode == 'mean':
        updated_target_val += -(1 - smoothing_factor) / num_valid_targets
    else:
        updated_target_val += -(1 - smoothing_factor)
    tl.store(loss_output_ptr, row_loss)
    tl.store(input_ptr + target_label, updated_target_val)
