code: def layer_norm_bwd_dwdb(
      DW,
      DB,
      FINAL_DW,
      FINAL_DB,
      M,
      N,
      BLOCK_SIZE_M: tl.constexpr,
      BLOCK_SIZE_N: tl.constexpr,
  ):
      # fmt: on
  
      pid = tl.program_id(0)
  
      cols = pid * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
      mask_cols = cols < 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)
  
      for i in range(0, M, BLOCK_SIZE_M):
          rows = i + tl.arange(0, BLOCK_SIZE_M)
          offs = rows[:, None] * N + cols[None, :]
          mask_rm = rows < M
  
          dw += tl.load(
              DW + offs,
              mask=mask_rm[:, None] & mask_cols[None, :],
              other=0.0,
          )
          db += tl.load(
              DB + offs,
              mask=mask_rm[:, None] & mask_cols[None, :],
              other=0.0,
          )
  
      sum_dw = tl.sum(dw, axis=0)
      sum_db = tl.sum(db, axis=0)
  
      cols = pid * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N)
      mask_cols = cols < N
  
      tl.store(FINAL_DW + cols, sum_dw, mask=mask_cols)
      tl.store(FINAL_DB + cols, sum_db, mask=mask_cols)
  
imports: 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