# AOT ID: ['0_forward']
from ctypes import c_void_p, c_long, c_int
import torch
import math
import random
import os
import tempfile
from math import inf, nan
from torch._inductor.hooks import run_intermediate_hooks
from torch._inductor.utils import maybe_profile
from torch._inductor.codegen.memory_planning import _align as align
from torch import device, empty_strided
from torch._inductor.async_compile import AsyncCompile
from torch._inductor.select_algorithm import extern_kernels
from torch._inductor.codegen.multi_kernel import MultiKernelCall
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid, split_scan_grid, grid_combo_kernels, start_graph, end_graph
from torch._C import _cuda_getCurrentRawStream as get_raw_stream

aten = torch.ops.aten
inductor_ops = torch.ops.inductor
_quantized = torch.ops._quantized
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
alloc_from_pool = torch.ops.inductor._alloc_from_pool
async_compile = AsyncCompile()


# kernel path: runs/run_shard_6/inductor_cache/fy/cfyu3frphog3iitdcz5tgawpqhvmxx3pgmb3qcnek46yuarja7tl.py
# Topologically Sorted Source Nodes: [mu, std, sub, mul, add, sqrt, x_hat, mul_1, out], Original ATen: [aten.mean, aten.std, aten.sub, aten.mul, aten.add, aten.sqrt, aten.div]
# Source node to ATen node mapping:
#   add => add
#   mu => mean
#   mul => mul
#   mul_1 => mul_1
#   out => add_1
#   sqrt => sqrt_1
#   std => sqrt, var
#   sub => sub
#   x_hat => div
# Graph fragment:
#   %mean : [num_users=1] = call_function[target=torch.ops.aten.mean.dim](args = (%primals_1, [0]), kwargs = {})
#   %var : [num_users=1] = call_function[target=torch.ops.aten.var.correction](args = (%primals_1, [0]), kwargs = {correction: 1.0})
#   %sqrt : [num_users=1] = call_function[target=torch.ops.aten.sqrt.default](args = (%var,), kwargs = {})
#   %sub : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%primals_1, %mean), kwargs = {})
#   %mul : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sqrt, %sqrt), kwargs = {})
#   %add : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul, 1e-05), kwargs = {})
#   %sqrt_1 : [num_users=1] = call_function[target=torch.ops.aten.sqrt.default](args = (%add,), kwargs = {})
#   %div : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%sub, %sqrt_1), kwargs = {})
#   %mul_1 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%primals_2, %div), kwargs = {})
#   %add_1 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%mul_1, %primals_3), kwargs = {})
triton_poi_fused_add_div_mean_mul_sqrt_std_sub_0 = async_compile.triton('triton_poi_fused_add_div_mean_mul_sqrt_std_sub_0', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[16], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_div_mean_mul_sqrt_std_sub_0', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 7, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_div_mean_mul_sqrt_std_sub_0(in_out_ptr0, in_ptr0, in_ptr1, in_ptr2, xnumel, XBLOCK : tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (x0), xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + (x2), xmask)
    tmp2 = tl.load(in_ptr1 + (x0), xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (4 + x0), xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr1 + (8 + x0), xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr1 + (12 + x0), xmask, eviction_policy='evict_last')
    tmp32 = tl.load(in_ptr2 + (x0), xmask, eviction_policy='evict_last')
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = 4.0
    tmp10 = tmp8 / tmp9
    tmp11 = tmp1 - tmp10
    tmp12 = tmp2 - tmp10
    tmp13 = tmp12 * tmp12
    tmp14 = tmp3 - tmp10
    tmp15 = tmp14 * tmp14
    tmp16 = tmp13 + tmp15
    tmp17 = tmp5 - tmp10
    tmp18 = tmp17 * tmp17
    tmp19 = tmp16 + tmp18
    tmp20 = tmp7 - tmp10
    tmp21 = tmp20 * tmp20
    tmp22 = tmp19 + tmp21
    tmp23 = 3.0
    tmp24 = tmp22 / tmp23
    tmp25 = libdevice.sqrt(tmp24)
    tmp26 = tmp25 * tmp25
    tmp27 = 1e-05
    tmp28 = tmp26 + tmp27
    tmp29 = libdevice.sqrt(tmp28)
    tmp30 = tmp11 / tmp29
    tmp31 = tmp0 * tmp30
    tmp33 = tmp31 + tmp32
    tl.store(in_out_ptr0 + (x2), tmp33, xmask)
''', device_str='cuda')


async_compile.wait(globals())
del async_compile

def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (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, 1), torch.float32)
        buf1 = buf0; del buf0  # reuse
        # Topologically Sorted Source Nodes: [mu, std, sub, mul, add, sqrt, x_hat, mul_1, out], Original ATen: [aten.mean, aten.std, aten.sub, aten.mul, aten.add, aten.sqrt, aten.div]
        stream0 = get_raw_stream(0)
        triton_poi_fused_add_div_mean_mul_sqrt_std_sub_0.run(buf1, primals_2, primals_1, primals_3, 16, grid=grid(16), stream=stream0)
        del primals_2
        del primals_3
    return (buf1, primals_1, )


def benchmark_compiled_module(times=10, repeat=10):
    from torch._dynamo.testing import rand_strided
    from torch._inductor.utils import print_performance
    primals_1 = rand_strided((4, 4), (4, 1), device='cuda:0', dtype=torch.float32)
    primals_2 = rand_strided((4, ), (1, ), device='cuda:0', dtype=torch.float32)
    primals_3 = rand_strided((4, ), (1, ), device='cuda:0', dtype=torch.float32)
    fn = lambda: call([primals_1, primals_2, primals_3])
    return print_performance(fn, times=times, repeat=repeat)


if __name__ == "__main__":
    from torch._inductor.wrapper_benchmark import compiled_module_main
    compiled_module_main('None', benchmark_compiled_module)

import torch
import torch.nn as nn


class CustomBatchNormAutograd(nn.Module):
    """
    This nn.module implements a custom version of the batch norm operation for MLPs.
    The operations called in self.forward track the history if the input tensors have the
    flag requires_grad set to True. The backward pass does not need to be implemented, it
    is dealt with by the automatic differentiation provided by PyTorch.
    """

    def __init__(self, n_neurons, eps=1e-05):
        """
        Initializes CustomBatchNormAutograd object.
        
        Args:
          n_neurons: int specifying the number of neurons
          eps: small float to be added to the variance for stability
        
        TODO:
          Save parameters for the number of neurons and eps.
          Initialize parameters gamma and beta via nn.Parameter
        """
        super(CustomBatchNormAutograd, self).__init__()
        self.n_neurons = n_neurons
        self.eps = eps
        self.gammas = torch.nn.Parameter(torch.ones(n_neurons))
        self.betas = torch.nn.Parameter(torch.zeros(n_neurons))

    def forward(self, input):
        """
        Compute the batch normalization
        
        Args:
          input: input tensor of shape (n_batch, n_neurons)
        Returns:
          out: batch-normalized tensor
        
        TODO:
          Check for the correctness of the shape of the input tensor.
          Implement batch normalization forward pass as given in the assignment.
          For the case that you make use of torch.var be aware that the flag unbiased=False should be set.
        """
        shape = input.size()
        if len(shape) == 1:
            input.view(1, -1)
            shape = input.size()
        if len(shape) > 2:
            raise ValueError('Expected 1-D or 2-D tensor (got {})'.format(
                str(shape)))
        elif input.shape[1] != self.n_neurons:
            raise ValueError('Expected _ x {} tensor (got {} x {})'.format(
                str(self.n_neurons), str(shape[0]), str(shape[1])))
        x = input
        mu = x.mean(dim=0)
        std = x.std(dim=0)
        x_hat = (x - mu) / torch.sqrt(std * std + self.eps)
        out = self.gammas * x_hat + self.betas
        return out


def get_inputs():
    return [torch.rand([4, 4])]


def get_init_inputs():
    return [[], {'n_neurons': 4}]

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_mean_mul_sqrt_std_sub_0(in_out_ptr0, in_ptr0,
    in_ptr1, in_ptr2, xnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + x0, xmask, eviction_policy='evict_last')
    tmp1 = tl.load(in_ptr1 + x2, xmask)
    tmp2 = tl.load(in_ptr1 + x0, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr1 + (4 + x0), xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr1 + (8 + x0), xmask, eviction_policy='evict_last')
    tmp7 = tl.load(in_ptr1 + (12 + x0), xmask, eviction_policy='evict_last')
    tmp32 = tl.load(in_ptr2 + x0, xmask, eviction_policy='evict_last')
    tmp4 = tmp2 + tmp3
    tmp6 = tmp4 + tmp5
    tmp8 = tmp6 + tmp7
    tmp9 = 4.0
    tmp10 = tmp8 / tmp9
    tmp11 = tmp1 - tmp10
    tmp12 = tmp2 - tmp10
    tmp13 = tmp12 * tmp12
    tmp14 = tmp3 - tmp10
    tmp15 = tmp14 * tmp14
    tmp16 = tmp13 + tmp15
    tmp17 = tmp5 - tmp10
    tmp18 = tmp17 * tmp17
    tmp19 = tmp16 + tmp18
    tmp20 = tmp7 - tmp10
    tmp21 = tmp20 * tmp20
    tmp22 = tmp19 + tmp21
    tmp23 = 3.0
    tmp24 = tmp22 / tmp23
    tmp25 = libdevice.sqrt(tmp24)
    tmp26 = tmp25 * tmp25
    tmp27 = 1e-05
    tmp28 = tmp26 + tmp27
    tmp29 = libdevice.sqrt(tmp28)
    tmp30 = tmp11 / tmp29
    tmp31 = tmp0 * tmp30
    tmp33 = tmp31 + tmp32
    tl.store(in_out_ptr0 + x2, tmp33, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (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, 1), torch.float32)
        buf1 = buf0
        del buf0
        get_raw_stream(0)
        triton_poi_fused_add_div_mean_mul_sqrt_std_sub_0[grid(16)](buf1,
            primals_2, primals_1, primals_3, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        del primals_2
        del primals_3
    return buf1, primals_1


class CustomBatchNormAutogradNew(nn.Module):
    """
    This nn.module implements a custom version of the batch norm operation for MLPs.
    The operations called in self.forward track the history if the input tensors have the
    flag requires_grad set to True. The backward pass does not need to be implemented, it
    is dealt with by the automatic differentiation provided by PyTorch.
    """

    def __init__(self, n_neurons, eps=1e-05):
        """
        Initializes CustomBatchNormAutograd object.
        
        Args:
          n_neurons: int specifying the number of neurons
          eps: small float to be added to the variance for stability
        
        TODO:
          Save parameters for the number of neurons and eps.
          Initialize parameters gamma and beta via nn.Parameter
        """
        super(CustomBatchNormAutogradNew, self).__init__()
        self.n_neurons = n_neurons
        self.eps = eps
        self.gammas = torch.nn.Parameter(torch.ones(n_neurons))
        self.betas = torch.nn.Parameter(torch.zeros(n_neurons))

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