code: def _blocksparse_softmax_bwd(
      DA,
      stride_zdx,  #
      DOut,
      stride_zdout,  #
      Out,
      stride_zout,  #
      scale,  #
      LUT,  #
      DR,
      extent,
      stride_zr,
      stride_hr,
      stride_er,  #
      is_causal,  #
      ROW_SIZE: tl.constexpr,  #
      BLOCK_SIZE: tl.constexpr,  #
      IS_DENSE: tl.constexpr,
  ):
      h = tl.program_id(0)
      m = tl.program_id(1)
      z = tl.program_id(2)
      # create index ranges
      hm = h * tl.num_programs(1) + m
      lane_n = tl.arange(0, ROW_SIZE) % BLOCK_SIZE
      block_n = tl.arange(0, ROW_SIZE) // BLOCK_SIZE
      # extract information from LUT
      header = LUT + (hm // BLOCK_SIZE) * 2
      size = tl.load(header + 0)
      offset = tl.load(header + 1)
      # row-col offset
      off_mn = (offset + block_n) * BLOCK_SIZE * BLOCK_SIZE
      off_mn += (m % BLOCK_SIZE) * BLOCK_SIZE
      mask = block_n < size
      # pointers
      As = Out + z * stride_zout + off_mn
      DOuts = DOut + z * stride_zdout + off_mn
      # do not need to read column indices in the dense case
      if IS_DENSE:
          ns = tl.arange(0, ROW_SIZE)
      else:
          off_lut = (
              offset
              + 2
              * tl.num_programs(0)
              * tl.num_programs(1)
              // BLOCK_SIZE
          )
          start_n = tl.load(LUT + off_lut + block_n, mask=mask, other=0)
          ns = start_n * BLOCK_SIZE + lane_n
      # load data
      a = tl.load(As + lane_n, mask=mask, other=0.0)
      a = a.to(tl.float32)
      dout = tl.load(DOuts + lane_n, mask=mask, other=0.0)
      dout = dout.to(tl.float32)
      # compute
      a = tl.where((ns > m) & is_causal & (a == a), 0.0, a)
      da = a * (dout - tl.sum(a * dout, 0))
      # apply relative attention
      if DR is not None:
          DR += z * stride_zr
          DR += h * stride_hr
          off_lo = (extent - m - 1) + ns
          mask_lo = (off_lo >= 0) & (off_lo < extent) & mask
          tl.store(DR + m * extent + off_lo, da, mask=mask_lo)
      da = da * scale
      # convert da
      # write-back
      DAs = DA + z * stride_zdx + off_mn
      tl.store(DAs + lane_n, da, mask=mask)
  
file_path: triton/python/triton/ops/blocksparse/softmax.py
imports: import torch
  from ...None import jit
  from ...None import language
  from ...None import next_power_of_2
label_desc: Use image node_id 3 for calling a global function with example usage: _blocksparse_softmax_bwd(DA, stride_zdx, DOut, stride_zdout, Out, stride_zout, scale, LUT, DR, extent, stride_zr, stride_hr, stride_er, is_causal, ROW_SIZE, BLOCK_SIZE, IS_DENSE) without return types