code: def _layer_norm_forward(
      x_ptr,
      w_ptr,
      b_ptr,
      eps,
      x_invstd_ptr,
      x_mean_ptr,
      y_ptr,
      M: tl.constexpr,
      N: tl.constexpr,
      M_BLOCK: tl.constexpr,
      N_BLOCK: tl.constexpr,
  ):
      m_idx = tl.program_id(0) * M_BLOCK + tl.arange(0, M_BLOCK)
      m_mask = m_idx < M
      n_idx = tl.arange(0, N_BLOCK)
      n_mask = n_idx < N
      mask = m_mask[:, None] & n_mask[None, :]
      x = tl.load(
          x_ptr + N * m_idx[:, None] + n_idx[None, :], mask, other=0
      ).to(tl.float32)
      x_mean = tl.sum(x, 1) / N
      tl.store(x_mean_ptr + m_idx, x_mean, m_mask)
      x_bar = x - x_mean[:, None]
      x_var = tl.sum(x_bar * x_bar, 1) / N
      x_invstd = rsqrt(x_var + eps)
      tl.store(x_invstd_ptr + m_idx, x_invstd, m_mask)
      x_hat = x_bar * x_invstd[:, None]
      w = tl.load(w_ptr + n_idx, n_mask, other=0).to(tl.float32)[
          None, :
      ]
      b = tl.load(b_ptr + n_idx, n_mask, other=0).to(tl.float32)[
          None, :
      ]
      y = w * x_hat + b
      tl.store(y_ptr + N * m_idx[:, None] + n_idx[None, :], y, mask)
  
imports: from packaging.version import Version
  import triton
  import triton.language
  from triton import Config
label_desc: Use image node_id 1 for calling a global function with example usage: _layer_norm_forward(x_ptr, w_ptr, b_ptr, eps, x_invstd_ptr, x_mean_ptr, y_ptr, M, N, M_BLOCK, N_BLOCK) without return types