import torch
from torch._inductor.select_algorithm import extern_kernels
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 import triton_helpers
from torch._inductor.runtime.triton_helpers import libdevice
import torch.nn as nn
from torch.nn import Conv2d
from torch.nn import ReLU
from torch.nn import InstanceNorm2d
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor


@triton.jit
def triton_poi_fused_0(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    xnumel = 9
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y0 = yindex % 128
    y1 = yindex // 128
    tmp0 = tl.load(in_ptr0 + (x2 + 9 * y3), xmask, eviction_policy='evict_last'
        )
    tl.store(out_ptr0 + (y0 + 128 * x2 + 1152 * y1), tmp0, xmask)


@triton.jit
def triton_poi_fused_1(in_ptr0, out_ptr0, ynumel, xnumel, YBLOCK: tl.
    constexpr, XBLOCK: tl.constexpr):
    ynumel = 512
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    x2 = xindex
    y3 = yindex
    y0 = yindex % 128
    y1 = yindex // 128
    tmp0 = tl.load(in_ptr0 + (x2 + 4096 * y3), ymask, eviction_policy=
        'evict_last')
    tl.store(out_ptr0 + (y0 + 128 * x2 + 524288 * y1), tmp0, ymask)


@triton.jit
def triton_poi_fused_convolution_2(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl
    .constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x0 = xindex % 128
    tmp0 = tl.load(in_out_ptr0 + x2, None)
    tmp1 = tl.load(in_ptr0 + x0, None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tl.store(in_out_ptr0 + x2, tmp2, None)


@triton.jit
def triton_red_fused__native_batch_norm_legit_3(in_ptr0, out_ptr0, out_ptr1,
    out_ptr2, xnumel, rnumel, XBLOCK: tl.constexpr, RBLOCK: tl.constexpr):
    xnumel = 512
    rnumel = 4096
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    tmp2_mean = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp2_m2 = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    tmp2_weight = tl.zeros([XBLOCK, RBLOCK], tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r1 = rindex
        tmp0 = tl.load(in_ptr0 + (128 * r1 + 524288 * (x0 // 128) + x0 % 
            128), rmask & xmask, eviction_policy='evict_first', other=0.0)
        tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
        tmp2_mean_next, tmp2_m2_next, tmp2_weight_next = (triton_helpers.
            welford_reduce(tmp1, tmp2_mean, tmp2_m2, tmp2_weight, roffset == 0)
            )
        tmp2_mean = tl.where(rmask & xmask, tmp2_mean_next, tmp2_mean)
        tmp2_m2 = tl.where(rmask & xmask, tmp2_m2_next, tmp2_m2)
        tmp2_weight = tl.where(rmask & xmask, tmp2_weight_next, tmp2_weight)
    tmp2_tmp, tmp3_tmp, tmp4_tmp = triton_helpers.welford(tmp2_mean,
        tmp2_m2, tmp2_weight, 1)
    tmp2 = tmp2_tmp[:, None]
    tmp3 = tmp3_tmp[:, None]
    tmp4_tmp[:, None]
    tl.store(out_ptr0 + x0, tmp2, xmask)
    tl.store(out_ptr1 + x0, tmp3, xmask)
    tmp5 = 4096.0
    tmp6 = tmp3 / tmp5
    tmp7 = 1e-05
    tmp8 = tmp6 + tmp7
    tmp9 = libdevice.rsqrt(tmp8)
    tl.store(out_ptr2 + x0, tmp9, xmask)


@triton.jit
def triton_poi_fused_relu_4(in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x3 = xindex
    x0 = xindex % 128
    x2 = xindex // 524288
    tmp0 = tl.load(in_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr1 + (x0 + 128 * x2), None, eviction_policy=
        'evict_last')
    tmp3 = tl.load(in_ptr2 + (x0 + 128 * x2), None, eviction_policy=
        'evict_last')
    tmp2 = tmp0 - tmp1
    tmp4 = 4096.0
    tmp5 = tmp3 / tmp4
    tmp6 = 1e-05
    tmp7 = tmp5 + tmp6
    tmp8 = libdevice.rsqrt(tmp7)
    tmp9 = tmp2 * tmp8
    tmp10 = tl.full([1], 0, tl.int32)
    tmp11 = triton_helpers.maximum(tmp10, tmp9)
    tl.store(out_ptr0 + x3, tmp11, None)


@triton.jit
def triton_poi_fused_add_relu_5(in_ptr0, in_ptr1, in_ptr2, in_ptr3,
    out_ptr0, ynumel, xnumel, YBLOCK: tl.constexpr, XBLOCK: tl.constexpr):
    xnumel = 128
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    x2 = xindex
    y3 = yindex
    y1 = yindex // 4096
    y0 = yindex % 4096
    tmp0 = tl.load(in_ptr0 + (x2 + 128 * y3), xmask, eviction_policy=
        'evict_last')
    tmp1 = tl.load(in_ptr1 + (x2 + 128 * y3), xmask, eviction_policy=
        'evict_last')
    tmp2 = tl.load(in_ptr2 + (x2 + 128 * y1), xmask, eviction_policy=
        'evict_last')
    tmp4 = tl.load(in_ptr3 + (x2 + 128 * y1), xmask, eviction_policy=
        'evict_last')
    tmp3 = tmp1 - tmp2
    tmp5 = 4096.0
    tmp6 = tmp4 / tmp5
    tmp7 = 1e-05
    tmp8 = tmp6 + tmp7
    tmp9 = libdevice.rsqrt(tmp8)
    tmp10 = tmp3 * tmp9
    tmp11 = tmp0 + tmp10
    tmp12 = tl.full([1, 1], 0, tl.int32)
    tmp13 = triton_helpers.maximum(tmp12, tmp11)
    tl.store(out_ptr0 + (y0 + 4096 * x2 + 524288 * y1), tmp13, xmask)


@triton.jit
def triton_poi_fused_threshold_backward_6(in_ptr0, out_ptr0, ynumel, xnumel,
    YBLOCK: tl.constexpr, XBLOCK: tl.constexpr):
    ynumel = 512
    yoffset = tl.program_id(1) * YBLOCK
    yindex = yoffset + tl.arange(0, YBLOCK)[None, :]
    ymask = yindex < ynumel
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, YBLOCK], True, tl.int1)
    x2 = xindex
    y3 = yindex
    y0 = yindex % 128
    y1 = yindex // 128
    tmp0 = tl.load(in_ptr0 + (x2 + 4096 * y3), ymask, eviction_policy=
        'evict_last')
    tmp1 = 0.0
    tmp2 = tmp0 <= tmp1
    tl.store(out_ptr0 + (y0 + 128 * x2 + 524288 * y1), tmp2, ymask)


def call(args):
    primals_1, primals_2, primals_3, primals_4, primals_5 = args
    args.clear()
    assert_size_stride(primals_1, (128, 128, 3, 3), (1152, 9, 3, 1))
    assert_size_stride(primals_2, (128,), (1,))
    assert_size_stride(primals_3, (4, 128, 64, 64), (524288, 4096, 64, 1))
    assert_size_stride(primals_4, (128, 128, 3, 3), (1152, 9, 3, 1))
    assert_size_stride(primals_5, (128,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((128, 128, 3, 3), (1152, 1, 384, 128),
            torch.float32)
        get_raw_stream(0)
        triton_poi_fused_0[grid(16384, 9)](primals_1, buf0, 16384, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_1
        buf1 = empty_strided_cuda((4, 128, 64, 64), (524288, 1, 8192, 128),
            torch.float32)
        triton_poi_fused_1[grid(512, 4096)](primals_3, buf1, 512, 4096,
            XBLOCK=32, YBLOCK=32, num_warps=4, num_stages=1)
        del primals_3
        buf2 = empty_strided_cuda((128, 128, 3, 3), (1152, 1, 384, 128),
            torch.float32)
        triton_poi_fused_0[grid(16384, 9)](primals_4, buf2, 16384, 9,
            XBLOCK=16, YBLOCK=64, num_warps=4, num_stages=1)
        del primals_4
        buf3 = extern_kernels.convolution(buf1, buf0, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf3, (4, 128, 64, 64), (524288, 1, 8192, 128))
        buf4 = buf3
        del buf3
        triton_poi_fused_convolution_2[grid(2097152)](buf4, primals_2, 
            2097152, XBLOCK=512, num_warps=8, num_stages=1)
        del primals_2
        buf5 = empty_strided_cuda((1, 512, 1, 1), (512, 1, 512, 512), torch
            .float32)
        buf6 = empty_strided_cuda((1, 512, 1, 1), (512, 1, 512, 512), torch
            .float32)
        buf8 = empty_strided_cuda((1, 512, 1, 1), (512, 1, 512, 512), torch
            .float32)
        triton_red_fused__native_batch_norm_legit_3[grid(512)](buf4, buf5,
            buf6, buf8, 512, 4096, XBLOCK=1, RBLOCK=2048, num_warps=16,
            num_stages=1)
        buf9 = empty_strided_cuda((4, 128, 64, 64), (524288, 1, 8192, 128),
            torch.float32)
        triton_poi_fused_relu_4[grid(2097152)](buf4, buf5, buf6, buf9, 
            2097152, XBLOCK=1024, num_warps=4, num_stages=1)
        buf10 = extern_kernels.convolution(buf9, buf2, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf10, (4, 128, 64, 64), (524288, 1, 8192, 128))
        buf11 = buf10
        del buf10
        triton_poi_fused_convolution_2[grid(2097152)](buf11, primals_5, 
            2097152, XBLOCK=512, num_warps=8, num_stages=1)
        del primals_5
        buf12 = buf6
        del buf6
        buf13 = empty_strided_cuda((1, 512, 1, 1), (512, 1, 512, 512),
            torch.float32)
        buf15 = empty_strided_cuda((1, 512, 1, 1), (512, 1, 512, 512),
            torch.float32)
        triton_red_fused__native_batch_norm_legit_3[grid(512)](buf11, buf12,
            buf13, buf15, 512, 4096, XBLOCK=1, RBLOCK=2048, num_warps=16,
            num_stages=1)
        buf16 = empty_strided_cuda((4, 128, 64, 64), (524288, 4096, 64, 1),
            torch.float32)
        triton_poi_fused_add_relu_5[grid(16384, 128)](buf1, buf11, buf12,
            buf13, buf16, 16384, 128, XBLOCK=32, YBLOCK=32, num_warps=4,
            num_stages=1)
        del buf13
        buf17 = empty_strided_cuda((4, 128, 64, 64), (524288, 1, 8192, 128),
            torch.bool)
        triton_poi_fused_threshold_backward_6[grid(512, 4096)](buf16, buf17,
            512, 4096, XBLOCK=32, YBLOCK=32, num_warps=4, num_stages=1)
    return buf16, buf0, buf1, buf2, buf4, reinterpret_tensor(buf8, (512,),
        (1,), 0), buf9, buf11, reinterpret_tensor(buf15, (512,), (1,), 0
        ), buf17, reinterpret_tensor(buf12, (1, 512, 1, 1), (512, 1, 1, 1), 0
        ), reinterpret_tensor(buf5, (1, 512, 1, 1), (512, 1, 1, 1), 0)


class mySConv(nn.Module):

    def __init__(self, num_filter=128, stride=1, in_channels=128):
        super(mySConv, self).__init__()
        self.conv = Conv2d(out_channels=num_filter, kernel_size=3, stride=
            stride, padding=1, in_channels=in_channels)
        self.bn = InstanceNorm2d(num_features=num_filter)
        self.relu = ReLU()

    def forward(self, x):
        return self.relu(self.bn(self.conv(x)))


class mySBlockNew(nn.Module):

    def __init__(self, num_filter=128):
        super(mySBlockNew, self).__init__()
        self.myconv = mySConv(num_filter=num_filter, stride=1, in_channels=
            num_filter)
        self.conv = Conv2d(out_channels=num_filter, kernel_size=3, padding=
            1, in_channels=num_filter)
        self.bn = InstanceNorm2d(num_features=num_filter)
        self.relu = ReLU()

    def forward(self, input_0):
        primals_1 = self.myconv.conv.weight
        primals_2 = self.myconv.conv.bias
        primals_4 = self.conv.weight
        primals_5 = self.conv.bias
        primals_3 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4, primals_5])
        return output[0]


Measured latency_ms: 0.219836