code: def _rms_norm_forward(x, attn_norm_weights, eps):
      if not x.is_contiguous():
          raise ValueError("data must be contiguous")
      if attn_norm_weights is not None:
          if not attn_norm_weights.is_contiguous():
              raise ValueError("weights must be contiguous")
      out = torch.empty_like(x)
      x_arg = x.reshape(-1, x.shape[-1])
      M, N = x_arg.shape
      # Less than 64KB per feature: enqueue fused kernel
      MAX_FUSED_SIZE = 65536 // x.element_size()
      BLOCK_SIZE = min(MAX_FUSED_SIZE, triton.next_power_of_2(N))
      BLOCK_SIZE = max(BLOCK_SIZE, 128)
      BLOCK_SIZE = min(BLOCK_SIZE, 8192)
      # heuristics for number of warps
      num_warps = min(max(BLOCK_SIZE // 256, 1), 8)
      _rms_norm_kernel[(M,)](
          x_arg,
          out,
          attn_norm_weights,
          eps,
          x_arg.stride(0),
          N,
          BLOCK_SIZE=BLOCK_SIZE,
          num_warps=num_warps,
          INCLUDE_WEIGHT=attn_norm_weights is not None,
      )
      return out
  
imports: import torch
  import triton
  import triton.language
label_desc: Use image node_id 3 for calling a global function with example usage: _rms_norm_forward(x, attn_norm_weights, eps) and returns: out