import torch
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid
from torch._C import _cuda_getCurrentRawStream as get_raw_stream
from torch._inductor.runtime.triton_helpers import libdevice
import torch.nn as nn
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda


@triton.jit
def triton_poi_fused_add_div_mul_sqrt_sub_sum_var_0(in_out_ptr0, in_ptr0,
    in_ptr1, in_ptr2, xnumel, XBLOCK: tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x4 = xindex
    x5 = xindex % 64
    tmp0 = tl.load(in_ptr0 + x0, xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + x4, xmask)
    tmp2 = tl.load(in_ptr1 + x5, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (64 + x5), xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr1 + (128 + x5), xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr1 + (192 + x5), xmask, eviction_policy='evict_last')
    tmp31 = tl.load(in_ptr2 + x0, xmask, eviction_policy='evict_last')
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = 0.25
    tmp10 = tmp8 * tmp9
    tmp11 = tmp1 - tmp10
    tmp12 = 4.0
    tmp13 = tmp8 / tmp12
    tmp14 = tmp2 - tmp13
    tmp15 = tmp14 * tmp14
    tmp16 = tmp3 - tmp13
    tmp17 = tmp16 * tmp16
    tmp18 = tmp15 + tmp17
    tmp19 = tmp5 - tmp13
    tmp20 = tmp19 * tmp19
    tmp21 = tmp18 + tmp20
    tmp22 = tmp7 - tmp13
    tmp23 = tmp22 * tmp22
    tmp24 = tmp21 + tmp23
    tmp25 = tmp24 / tmp12
    tmp26 = 1e-05
    tmp27 = tmp25 + tmp26
    tmp28 = libdevice.sqrt(tmp27)
    tmp29 = tmp11 / tmp28
    tmp30 = tmp0 * tmp29
    tmp32 = tmp30 + tmp31
    tl.store(in_out_ptr0 + x4, tmp32, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_2, (4,), (1,))
    assert_size_stride(primals_3, (4,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        buf1 = buf0
        del buf0
        get_raw_stream(0)
        triton_poi_fused_add_div_mul_sqrt_sub_sum_var_0[grid(256)](buf1,
            primals_2, primals_1, primals_3, 256, XBLOCK=128, num_warps=4,
            num_stages=1)
        del primals_2
        del primals_3
    return buf1, primals_1


class CustomBatchNormManualFunction(torch.autograd.Function):
    """
    This torch.autograd.Function implements a functional custom version of the batch norm operation for MLPs.
    Using torch.autograd.Function allows you to write a custom backward function.
    The function will be called from the nn.Module CustomBatchNormManualModule
    Inside forward the tensors are (automatically) not recorded for automatic differentiation since the backward
    pass is done via the backward method.
    The forward pass is not called directly but via the apply() method. This makes sure that the context objects
    are dealt with correctly. Example:
    my_bn_fct = CustomBatchNormManualFunction()
    normalized = fct.apply(input, gamma, beta, eps)
    """

    @staticmethod
    def forward(ctx, input, gamma, beta, eps=1e-05):
        """
        Compute the batch normalization

        Args:
        ctx: context object handling storing and retrival of tensors and constants and specifying
             whether tensors need gradients in backward pass
        input: input tensor of shape (n_batch, n_neurons)
        gamma: variance scaling tensor, applied per neuron, shpae (n_neurons)
        beta: mean bias tensor, applied per neuron, shpae (n_neurons)
        eps: small float added to the variance for stability
        Returns:
        out: batch-normalized tensor
        """
        batch_size = input.shape[0]
        mean = 1 / batch_size * torch.sum(input, dim=0)
        var = input.var(dim=0, unbiased=False)
        norm = (input - mean) / torch.sqrt(var + eps)
        out = gamma * norm + beta
        ctx.save_for_backward(norm, gamma, var)
        ctx.eps = eps
        return out

    @staticmethod
    def backward(ctx, grad_output):
        """
    Compute backward pass of the batch normalization.
    
    Args:
      ctx: context object handling storing and retrival of tensors and constants and specifying
           whether tensors need gradients in backward pass
    Returns:
      out: tuple containing gradients for all input arguments

    """
        normalized, gamma, var = ctx.saved_tensors
        eps = ctx.eps
        B = grad_output.shape[0]
        grad_gamma = (grad_output * normalized).sum(0)
        grad_beta = torch.sum(grad_output, dim=0)
        grad_input = torch.div(gamma, B * torch.sqrt(var + eps)) * (B *
            grad_output - grad_beta - grad_gamma * normalized)
        return grad_input, grad_gamma, grad_beta, None


class CustomBatchNormManualModuleNew(nn.Module):
    """
    This nn.module implements a custom version of the batch norm operation for MLPs.
    In self.forward the functional version CustomBatchNormManualFunction.forward is called.
    The automatic differentiation of PyTorch calls the backward method of this function in the backward pass.
    """

    def __init__(self, n_neurons, eps=1e-05):
        """
        Initializes CustomBatchNormManualModule object.

        Args:
          n_neurons: int specifying the number of neurons
          eps: small float to be added to the variance for stability
        """
        super(CustomBatchNormManualModuleNew, self).__init__()
        self.n_neurons = n_neurons
        self.eps = eps
        self.beta = nn.Parameter(torch.zeros(self.n_neurons))
        self.gamma = nn.Parameter(torch.ones(self.n_neurons))

    def forward(self, input_0):
        primals_2 = self.beta
        primals_3 = self.gamma
        primals_1 = input_0
        output = call([primals_1, primals_2, primals_3])
        return output[0]


Measured latency_ms: 0.016773