code: def _layer_norm_bwd_dwdb(
      DW,  # pointer to the partial sum of weights gradient
      DB,  # pointer to the partial sum of biases gradient
      FINAL_DW,  # pointer to the weights gradient
      FINAL_DB,  # pointer to the biases gradient
      M,  # GROUP_SIZE_M
      N,  # number of columns
      BLOCK_SIZE_M: tl.constexpr,
      BLOCK_SIZE_N: tl.constexpr,
  ):
      # Map the program id to the elements of DW and DB it should compute.
      pid = tl.program_id(0)
      cols = pid * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
      dw = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
      db = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32)
      # Iterate through the rows of DW and DB to sum the partial sums.
      for i in range(0, M, BLOCK_SIZE_M):
          rows = i + tl.arange(0, BLOCK_SIZE_M)
          mask = (rows[:, None] < M) & (cols[None, :] < N)
          offs = rows[:, None] * N + cols[None, :]
          dw += tl.load(DW + offs, mask=mask, other=0.0)
          db += tl.load(DB + offs, mask=mask, other=0.0)
      # Write the final sum to the output.
      sum_dw = tl.sum(dw, axis=0)
      sum_db = tl.sum(db, axis=0)
      tl.store(FINAL_DW + cols, sum_dw, mask=cols < N)
      tl.store(FINAL_DB + cols, sum_db, mask=cols < N)
  
file_path: triton/python/tutorials/05-layer-norm.py
imports: import torch
  import triton
  import triton.language
label_desc: Use image node_id 3 for calling a global function with example usage: _layer_norm_bwd_dwdb(DW, DB, FINAL_DW, FINAL_DB, M, N, BLOCK_SIZE_M, BLOCK_SIZE_N) without return types