code: def test_layer_norm(M, N, dtype, eps=1e-5, device="cuda"):
      # create data
      x_shape = (M, N)
      w_shape = (x_shape[-1],)
      weight = torch.rand(
          w_shape, dtype=dtype, device="cuda", requires_grad=True
      )
      bias = torch.rand(
          w_shape, dtype=dtype, device="cuda", requires_grad=True
      )
      x = -2.3 + 0.5 * torch.randn(x_shape, dtype=dtype, device="cuda")
      dy = 0.1 * torch.randn_like(x)
      x.requires_grad_(True)
      # forward pass
      y_tri = layer_norm(x, w_shape, weight, bias, eps)
      y_ref = torch.nn.functional.layer_norm(
          x, w_shape, weight, bias, eps
      ).to(dtype)
      # backward pass (triton)
      y_tri.backward(dy, retain_graph=True)
      dx_tri, dw_tri, db_tri = [
          _.grad.clone() for _ in [x, weight, bias]
      ]
      x.grad, weight.grad, bias.grad = None, None, None
      # backward pass (torch)
      y_ref.backward(dy, retain_graph=True)
      dx_ref, dw_ref, db_ref = [
          _.grad.clone() for _ in [x, weight, bias]
      ]
      # compare
      assert torch.allclose(y_tri, y_ref, atol=1e-2, rtol=0)
      assert torch.allclose(dx_tri, dx_ref, atol=1e-2, rtol=0)
      assert torch.allclose(db_tri, db_ref, atol=1e-2, rtol=0)
      assert torch.allclose(dw_tri, dw_ref, atol=1e-2, rtol=0)
  
file_path: triton/python/tutorials/05-layer-norm.py
imports: import torch
  import triton
  import triton.language
label_desc: Use image node_id 4 for calling a global function with example usage: test_layer_norm(M, N, dtype, eps, device) without return types