@staticmethod
@triton.jit
def backward(grad_input_ptr, grad_output_ptr, output_ptr, size, p,
    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,))
    output_block_ptr = tl.make_block_ptr(output_ptr, shape=(size,), strides
        =(1,), offsets=(offset,), block_shape=(block_size,), order=(0,))
    grad_output = tl.load(grad_output_block_ptr, boundary_check=(0,))
    output = tl.load(output_block_ptr, boundary_check=(0,))
    condition = output > 0.0
    grad_input = tl.where(condition, grad_output * (1 / (1 - p)), 0.0)
    tl.store(grad_input_block_ptr, grad_input, boundary_check=(0,))
