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
import math
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_per_fused_mean_0(in_ptr0, out_ptr0, xnumel, rnumel, XBLOCK: tl.
    constexpr):
    xnumel = 16
    RBLOCK: tl.constexpr = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + 16 * x0), xmask, other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.where(xmask, tmp1, 0)
    tmp4 = tl.sum(tmp3, 1)[:, None]
    tl.store(out_ptr0 + x0, tmp4, xmask)


@triton.jit
def triton_poi_fused_copy_1(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 24
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 6
    x1 = xindex // 6
    x2 = xindex
    tmp0 = x0
    tmp1 = tl.full([1], 5, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = -4 + x0
    tmp4 = tl.full([1], 1, tl.int64)
    tmp5 = tmp3 < tmp4
    tmp6 = tmp5 & tmp2
    tmp7 = tmp0 >= tmp4
    tmp8 = tmp0 < tmp1
    tmp9 = tmp7 & tmp8
    tmp10 = tmp9 & tmp6
    tmp11 = tl.load(in_ptr0 + (-1 + x0 + 4 * x1), tmp10 & xmask, other=0.0)
    tmp12 = 16.0
    tmp13 = tmp11 / tmp12
    tmp14 = tl.full(tmp13.shape, 0.0, tmp13.dtype)
    tmp15 = tl.where(tmp10, tmp13, tmp14)
    tmp16 = float('nan')
    tmp17 = tl.where(tmp9, tmp15, tmp16)
    tmp18 = tl.full(tmp17.shape, 0.0, tmp17.dtype)
    tmp19 = tl.where(tmp6, tmp17, tmp18)
    tmp20 = tmp3 >= tmp4
    tmp21 = tmp3 < tmp1
    tmp22 = tmp20 & tmp21
    tmp23 = tmp22 & tmp2
    tmp24 = tl.load(in_ptr0 + (-5 + x0 + 4 * x1), tmp23 & xmask, other=0.0)
    tmp25 = tmp24 / tmp12
    tmp26 = tl.full(tmp25.shape, 0.0, tmp25.dtype)
    tmp27 = tl.where(tmp23, tmp25, tmp26)
    tmp28 = tl.where(tmp22, tmp27, tmp16)
    tmp29 = tl.where(tmp5, tmp19, tmp28)
    tmp30 = tl.full(tmp29.shape, 0.0, tmp29.dtype)
    tmp31 = tl.where(tmp2, tmp29, tmp30)
    tmp32 = tmp0 < tmp4
    tmp33 = 4 + x0
    tmp34 = tmp33 >= tmp4
    tmp35 = tmp33 < tmp1
    tmp36 = tmp34 & tmp35
    tmp37 = tmp36 & tmp32
    tmp38 = tl.load(in_ptr0 + (3 + x0 + 4 * x1), tmp37 & xmask, other=0.0)
    tmp39 = tmp38 / tmp12
    tmp40 = tl.full(tmp39.shape, 0.0, tmp39.dtype)
    tmp41 = tl.where(tmp37, tmp39, tmp40)
    tmp42 = tl.where(tmp36, tmp41, tmp16)
    tmp43 = tl.full(tmp42.shape, 0.0, tmp42.dtype)
    tmp44 = tl.where(tmp32, tmp42, tmp43)
    tmp45 = tl.load(in_ptr0 + (-1 + x0 + 4 * x1), tmp9 & xmask, other=0.0)
    tmp46 = tmp45 / tmp12
    tmp47 = tl.full(tmp46.shape, 0.0, tmp46.dtype)
    tmp48 = tl.where(tmp9, tmp46, tmp47)
    tmp49 = tl.where(tmp9, tmp48, tmp16)
    tmp50 = tl.where(tmp32, tmp44, tmp49)
    tmp51 = tl.where(tmp2, tmp31, tmp50)
    tl.store(out_ptr0 + x2, tmp51, xmask)


@triton.jit
def triton_poi_fused_mul_2(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK: tl.
    constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x1 = xindex // 16
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp2 = tl.sigmoid(tmp1)
    tmp3 = tmp0 * tmp2
    tl.store(out_ptr0 + x2, tmp3, xmask)


def call(args):
    primals_1, primals_2 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_2, (1, 1, 3), (3, 3, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        get_raw_stream(0)
        triton_per_fused_mean_0[grid(16)](primals_1, buf0, 16, 16, XBLOCK=1,
            num_warps=2, num_stages=1)
        buf2 = empty_strided_cuda((4, 1, 6), (6, 6, 1), torch.float32)
        triton_poi_fused_copy_1[grid(24)](buf0, buf2, 24, XBLOCK=32,
            num_warps=1, num_stages=1)
        del buf0
        buf3 = extern_kernels.convolution(buf2, primals_2, stride=(1,),
            padding=(0,), dilation=(1,), transposed=False, output_padding=(
            0,), groups=1, bias=None)
        assert_size_stride(buf3, (4, 1, 4), (4, 4, 1))
        buf4 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused_mul_2[grid(256)](primals_1, buf3, buf4, 256,
            XBLOCK=256, num_warps=4, num_stages=1)
    return buf4, primals_1, primals_2, buf2, buf3


class CecaModuleNew(nn.Module):
    """Constructs a circular ECA module.

    ECA module where the conv uses circular padding rather than zero padding.
    Unlike the spatial dimension, the channels do not have inherent ordering nor
    locality. Although this module in essence, applies such an assumption, it is unnecessary
    to limit the channels on either "edge" from being circularly adapted to each other.
    This will fundamentally increase connectivity and possibly increase performance metrics
    (accuracy, robustness), without significantly impacting resource metrics
    (parameter size, throughput,latency, etc)

    Args:
        channels: Number of channels of the input feature map for use in adaptive kernel sizes
            for actual calculations according to channel.
            gamma, beta: when channel is given parameters of mapping function
            refer to original paper https://arxiv.org/pdf/1910.03151.pdf
            (default=None. if channel size not given, use k_size given for kernel size.)
        kernel_size: Adaptive selection of kernel size (default=3)
    """

    def __init__(self, channels=None, kernel_size=3, gamma=2, beta=1):
        super(CecaModuleNew, self).__init__()
        assert kernel_size % 2 == 1
        if channels is not None:
            t = int(abs(math.log(channels, 2) + beta) / gamma)
            kernel_size = max(t if t % 2 else t + 1, 3)
        self.conv = nn.Conv1d(1, 1, kernel_size=kernel_size, padding=0,
            bias=False)
        self.padding = (kernel_size - 1) // 2

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


Measured latency_ms: 0.033458