code: def _bwd_preprocess(
      Out,
      DO,
      Delta,
      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
      )
      # compute
      delta = tl.sum(o * do, axis=1)
      # write-back
      tl.store(Delta + off_m, delta)
  
file_path: triton/python/triton/ops/flash_attention.py
imports: import torch
  from ..None import cdiv, jit
  from ..None import language
label_desc: Use image node_id 2 for calling a global function with example usage: _bwd_preprocess(Out, DO, Delta, BLOCK_M, D_HEAD) without return types