code: def _backward(
      X, scale, DX, LUT, sizemax, stride_zx, stride_zdx, **meta
  ):
      pidhm = tl.program_id(0)
      pidz = tl.program_id(1)
      TN = meta["TN"]
      BLOCK = meta["BLOCK"]
      # create index ranges
      rxm = pidhm % BLOCK
      rbm = pidhm // BLOCK
      rxn = tl.arange(0, TN) % BLOCK
      rbn = tl.arange(0, TN) // BLOCK
      # extract information from look-up table
      header = LUT + rbm * 2
      size = tl.load(header + 0)
      offset = tl.load(header + 1)
      # bounds checking on lut
      check = rbn < size
      rbmn = tl.where(check, rbn, size - 1)
      # initialize pointers to block-sparse input
      blockid = tl.load(LUT + offset + rbmn * 4)
      X = (
          X
          + pidz * stride_zx
          + blockid * BLOCK * BLOCK
          + rxm * BLOCK
          + rxn
      )
      DX = (
          DX
          + pidz * stride_zdx
          + blockid * BLOCK * BLOCK
          + rxm * BLOCK
          + rxn
      )
      # compute fused softmax backward
      x = tl.load(X, mask=check, other=0)
      dx = tl.load(DX, mask=check, other=0)
      x = x.to(tl.float32)
      dx = dx.to(tl.float32)
      y = x * (dx - tl.sum(x * dx, 0)) * scale
      tl.store(DX, y, mask=check)
  
file_path: H2O/h2o_flexgen/benchmark/third_party/DeepSpeed/deepspeed/ops/sparse_attention/softmax.py
imports: import torch
  import triton
  import triton.language
label_desc: Use image node_id 4 for calling a global function with example usage: _backward(X, scale, DX, LUT, sizemax, stride_zx, stride_zdx) without return types