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 numpy as np
import torch.nn.parallel
import torch.optim
import torch
import torch.nn as nn
import torch.utils.data
import torch.utils.data.distributed
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_mul_0(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
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + x0, xmask)
    tmp1 = tl.load(in_ptr1 + x0, xmask)
    tmp2 = tmp0 * tmp1
    tl.store(out_ptr0 + x0, tmp2, 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, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_3, (4, 4, 4, 4), (64, 16, 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)
        get_raw_stream(0)
        triton_poi_fused_mul_0[grid(256)](primals_1, primals_2, buf0, 256,
            XBLOCK=128, num_warps=4, num_stages=1)
        del primals_1
        buf1 = extern_kernels.convolution(primals_3, buf0, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf1, (4, 4, 3, 3), (36, 9, 3, 1))
    return buf1, primals_2, primals_3, buf0


def get_mask(in_channels, channels, ks):
    in_channels = int(in_channels)
    channels = int(channels)
    if len(ks) == 1:
        mask = np.zeros((int(in_channels), int(channels), int(ks[0])))
    elif len(ks) == 2:
        mask = np.zeros((int(in_channels), int(channels), int(ks[0]), int(
            ks[1])))
    elif len(ks) == 3:
        mask = np.zeros((int(in_channels), int(channels), int(ks[0]), int(
            ks[1]), int(ks[2])))
    else:
        raise Error('not implement yet')
    for _ in range(in_channels):
        mask[_, _ % channels, :, :] = 1.0
    return mask


class DiagonalwiseRefactorizationNew(nn.Module):

    def __init__(self, in_channels, ks, stride=1, groups=1):
        super(DiagonalwiseRefactorizationNew, self).__init__()
        channels = in_channels // groups
        self.in_channels = in_channels
        self.groups = groups
        self.stride = stride
        p = (np.array(ks) - 1) // 2
        self.p = p.tolist()
        self.mask = nn.Parameter(torch.Tensor(get_mask(in_channels,
            channels, ks=ks)), requires_grad=False)
        self.weight = nn.Parameter(torch.Tensor(in_channels, channels, *ks),
            requires_grad=True)
        torch.nn.init.xavier_uniform_(self.weight.data)
        self.weight.data.mul_(self.mask.data)
        if len(ks) == 1:
            self.conv = nn.functional.conv1d
        elif len(ks) == 2:
            self.conv = nn.functional.conv2d
        elif len(ks) == 3:
            self.conv = nn.functional.conv3d
        else:
            raise Error(
                'The kernal size in DiagonalwiseRefactorization is wrong!')

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


Measured latency_ms: 0.023136