code: def _bwd_preprocess(
      Out,
      DO,
      L,
      NewDO,
      Delta,
      stride_ob,
      stride_oh,
      stride_om,
      stride_ok,
      stride_dob,
      stride_doh,
      stride_dom,
      stride_dok,
      BLOCK_M: tl.constexpr,
      D_HEAD: tl.constexpr,
  ):
      off_m = tl.program_id(0) * BLOCK_M + tl.arange(0, BLOCK_M)
      off_n = tl.arange(0, D_HEAD)
      # load
      o = tl.load(Out + off_m[:, None] * D_HEAD + off_n[None, :]).to(
          tl.float32
      )
      do = tl.load(DO + off_m[:, None] * D_HEAD + off_n[None, :]).to(
          tl.float32
      )
      denom = tl.load(L + off_m).to(tl.float32)
      # compute
      do = do / denom[:, None]
      delta = tl.sum(o * do, axis=1)
      # write-back
      tl.store(NewDO + off_m[:, None] * D_HEAD + off_n[None, :], do)
      tl.store(Delta + off_m, delta)
  
file_path: apex/apex/contrib/openfold_triton/_mha_kernel.py
imports: import triton
  import triton.language
label_desc: Use image node_id 4 for calling a global function with example usage: _bwd_preprocess(Out, DO, L, NewDO, Delta, stride_ob, stride_oh, stride_om, stride_ok, stride_dob, stride_doh, stride_dom, stride_dok, BLOCK_M, D_HEAD) without return types