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 torch.nn as nn
import torch.nn.init
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_convolution_0(in_ptr0, out_ptr0, ynumel, xnumel,
    YBLOCK: tl.constexpr, XBLOCK: tl.constexpr):
    ynumel = 12
    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 % 3
    y1 = yindex // 3
    tmp0 = tl.load(in_ptr0 + (x2 + 4096 * y3), ymask, eviction_policy=
        'evict_last')
    tl.store(out_ptr0 + (y0 + 3 * x2 + 12288 * y1), tmp0, ymask)


@triton.jit
def triton_poi_fused_convolution_1(in_ptr0, out_ptr0, ynumel, xnumel,
    YBLOCK: tl.constexpr, XBLOCK: tl.constexpr):
    ynumel = 24
    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
    y0 = yindex % 6
    y1 = yindex // 6
    y3 = yindex
    tmp0 = tl.load(in_ptr0 + (y0 + 6 * x2 + 24576 * y1), ymask,
        eviction_policy='evict_last')
    tl.store(out_ptr0 + (x2 + 4096 * y3), tmp0, ymask)


def call(args):
    arg0_1, arg1_1 = args
    args.clear()
    assert_size_stride(arg0_1, (6, 1, 3, 3), (9, 9, 3, 1))
    assert_size_stride(arg1_1, (4, 3, 64, 64), (12288, 4096, 64, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 3, 64, 64), (12288, 1, 192, 3), torch
            .float32)
        get_raw_stream(0)
        triton_poi_fused_convolution_0[grid(12, 4096)](arg1_1, buf0, 12, 
            4096, XBLOCK=64, YBLOCK=16, num_warps=4, num_stages=1)
        del arg1_1
        buf1 = extern_kernels.convolution(buf0, arg0_1, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=3, bias=None)
        assert_size_stride(buf1, (4, 6, 64, 64), (24576, 1, 384, 6))
        del arg0_1
        del buf0
        buf2 = empty_strided_cuda((4, 6, 64, 64), (24576, 4096, 64, 1),
            torch.float32)
        triton_poi_fused_convolution_1[grid(24, 4096)](buf1, buf2, 24, 4096,
            XBLOCK=32, YBLOCK=32, num_warps=4, num_stages=1)
        del buf1
    return buf2,


class SobelNew(nn.Module):

    def __init__(self):
        super(SobelNew, self).__init__()
        kernel = [[[-1, 0, 1], [-2, 0, 2], [-1, 0, 1]], [[-1, -2, -1], [0, 
            0, 0], [1, 2, 1]]]
        kernel = torch.Tensor(kernel).unsqueeze(1).repeat([3, 1, 1, 1])
        self.kernel = nn.Parameter(data=kernel, requires_grad=False)

    def forward(self, input_0):
        arg0_1 = self.kernel
        arg1_1 = input_0
        output = call([arg0_1, arg1_1])
        return output[0]


Measured latency_ms: 0.01488