code: def add_kernel(
      x_ptr,
      y_ptr,
      output_ptr,
      n_elements,
      BLOCK_SIZE: tl.constexpr,
  ):
      pid = tl.program_id(
          axis=0
      )  # We use a 1D launch grid so axis is 0.
      block_start = pid * BLOCK_SIZE
      offsets = block_start + tl.arange(0, BLOCK_SIZE)
      mask = offsets < n_elements
  
      x_block_ptr = tl.make_block_ptr(
          base=x_ptr,
          shape=(n_elements,),
          strides=(1,),
          offsets=(pid * BLOCK_SIZE,),
          block_shape=(BLOCK_SIZE,),
          order=(0,),
      )
      x = tl.load(
          x_block_ptr, boundary_check=(0,), padding_option="zero"
      )
  
      y = tl.load(y_ptr + offsets, mask=mask)
      output = x + y
      tl.store(output_ptr + offsets, output, mask=mask)
  
file_path: triton/python/test/unit/hopper/test_mixed_io.py
imports: import pytest
  import torch
  from torch.testing import assert_close
  import triton
  import triton.language
label_desc: Use image node_id 1 for calling a global function with example usage: add_kernel(x_ptr, y_ptr, output_ptr, n_elements, BLOCK_SIZE) without return types