@staticmethod
@triton.jit
def backward(grad_input_ptr, grad_output_ptr, input_ptr, size, block_size:
    tl.constexpr):
    pid = tl.program_id(0)
    offset = pid * block_size
    grad_input_block_ptr = tl.make_block_ptr(grad_input_ptr, shape=(size,),
        strides=(1,), offsets=(offset,), block_shape=(block_size,), order=(0,))
    grad_output_block_ptr = tl.make_block_ptr(grad_output_ptr, shape=(size,
        ), strides=(1,), offsets=(offset,), block_shape=(block_size,),
        order=(0,))
    input_block_ptr = tl.make_block_ptr(input_ptr, shape=(size,), strides=(
        1,), offsets=(offset,), block_shape=(block_size,), order=(0,))
    grad_output = tl.load(grad_output_block_ptr, boundary_check=(0,))
    input = tl.load(input_block_ptr, boundary_check=(0,))
    condition = input >= 0
    grad_input = tl.where(condition, grad_output, 0)
    tl.store(grad_input_block_ptr, grad_input, boundary_check=(0,))
