code: def test_assert(func: str):
      shape = (128,)
      x = torch.arange(0, shape[0], dtype=torch.int32, device="cuda")
      y = torch.zeros(shape, dtype=x.dtype, device="cuda")
      if func == "device_assert":
          kernel_device_assert[(1,)](x, y, BLOCK=shape[0])
      if func == "device_assert_passes":
          # Assert passes; no error.
          kernel_assert_passes[(1,)](x, y, BLOCK=shape[0])
      elif func == "no_debug":
          # TRITON_DEBUG=1 can override the debug flag
          kernel_device_assert_no_debug[(1,)](x, y, BLOCK=shape[0])
      elif func == "assert":
          kernel_assert[(1,)](x, y, BLOCK=shape[0])
      elif func == "static_assert":
          kernel_static_assert[(1,)](x, y, BLOCK=shape[0])
      elif func == "double_assert":
          # Launching a different kernel after the first one asserted used to
          # segfault.  What seems to have happened is:
          #  - The first kernel is enqueued but doesn't run yet.
          #  - We go to launch the second kernel.  Because this is the first time
          #    we're running it, we have to load the kernel into the GPU.
          #  - Loading the kernel takes some time, during which the first launch
          #    completes.
          #  - Now the GPU is in an error state.  We need to detect this inside
          #    the kernel-launch/loading code and bail out properly.  If we don't,
          #    we segfault.
          kernel_device_assert[(1,)](x, y, BLOCK=shape[0])
          kernel_assert_passes[(1,)](x, y, BLOCK=shape[0])
      assert_close(y, x)
  
file_path: triton/python/test/unit/language/assert_helper.py
imports: import sys
  import torch
  from torch.testing import assert_close
  import triton
  import triton.language
label_desc: Use image node_id 6 for calling a global function with example usage: test_assert(func) without return types