code: def _layer_norm_fwd_fused(
      X,  # pointer to the input
      Y,  # pointer to the output
      W,  # pointer to the weights
      B,  # pointer to the biases
      Mean,  # pointer to the mean
      Rstd,  # pointer to the 1/std
      stride,  # how much to increase the pointer when moving by 1 row
      N,  # number of columns in X
      eps,  # epsilon to avoid division by zero
      BLOCK_SIZE: tl.constexpr,
  ):
      # Map the program id to the row of X and Y it should compute.
      row = tl.program_id(0)
      Y += row * stride
      X += row * stride
      # Compute mean
      mean = 0
      _mean = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
      for off in range(0, N, BLOCK_SIZE):
          cols = off + tl.arange(0, BLOCK_SIZE)
          a = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
          _mean += a
      mean = tl.sum(_mean, axis=0) / N
      # Compute variance
      _var = tl.zeros([BLOCK_SIZE], dtype=tl.float32)
      for off in range(0, N, BLOCK_SIZE):
          cols = off + tl.arange(0, BLOCK_SIZE)
          x = tl.load(X + cols, mask=cols < N, other=0.0).to(tl.float32)
          x = tl.where(cols < N, x - mean, 0.0)
          _var += x * x
      var = tl.sum(_var, axis=0) / N
      rstd = 1 / tl.sqrt(var + eps)
      # Write mean / rstd
      tl.store(Mean + row, mean)
      tl.store(Rstd + row, rstd)
      # Normalize and apply linear transformation
      for off in range(0, N, BLOCK_SIZE):
          cols = off + tl.arange(0, BLOCK_SIZE)
          mask = cols < N
          w = tl.load(W + cols, mask=mask)
          b = tl.load(B + cols, mask=mask)
          x = tl.load(X + cols, mask=mask, other=0.0).to(tl.float32)
          x_hat = (x - mean) * rstd
          y = x_hat * w + b
          # Write output
          tl.store(Y + cols, y, mask=mask)
  
file_path: triton/python/tutorials/05-layer-norm.py
imports: import torch
  import triton
  import triton.language
label_desc: Use image node_id 1 for calling a global function with example usage: _layer_norm_fwd_fused(X, Y, W, B, Mean, Rstd, stride, N, eps, BLOCK_SIZE) without return types