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.functional as F
from torch.nn.modules.utils import _pair
assert_size_stride = torch._C._dynamo.guards.assert_size_stride


@triton.jit
def triton_poi_fused_convolution_0(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl
    .constexpr):
    xnumel = 144
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 9 % 4
    tmp0 = tl.load(in_out_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr0 + x1, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tl.store(in_out_ptr0 + x3, 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,), (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 = extern_kernels.convolution(primals_3, primals_1, stride=(1, 
            1), padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf0, (4, 4, 3, 3), (36, 9, 3, 1))
        buf1 = buf0
        del buf0
        get_raw_stream(0)
        triton_poi_fused_convolution_0[grid(144)](buf1, primals_2, 144,
            XBLOCK=128, num_warps=4, num_stages=1)
        del primals_2
    return buf1, primals_1, primals_3


class SymmetricQuantizeDequantize(torch.autograd.Function):

    @staticmethod
    def forward(ctx, input, precision, clamp_val):
        ctx.save_for_backward(input)
        """
        Compute quantization step size. Mapping (-max_val, max_val) linearly to (-127,127)
        """
        use_max = True
        if use_max:
            max_val = torch.max(torch.abs(input))
        else:
            max_val = clamp_val
        delta = max_val / (2 ** (precision - 1) - 1)
        input_clamped = torch.clamp(input, -max_val, max_val)
        input_q = torch.round(input_clamped / delta)
        if precision == 8:
            input_q = input_q
        elif precision == 16:
            input_q = input_q
        else:
            input_q = input_q
        """
        Dequantize introducing a quantization error in the data
        """
        input_dq = input_q * delta
        input_dq = input_dq
        return input.copy_(input_dq)

    @staticmethod
    def backward(ctx, grad_output):
        _input, = ctx.saved_tensors
        return grad_output, None


class nnConv2dSymQuantNew(nn.Conv2d):
    """ 
    Computes 2d conv output
    Weights are quantized and dequantized introducing a quantization error
    """

    def __init__(self, in_channels, out_channels, kernel_size, stride=1,
        padding=1, dilation=1, groups=1, bias=True, padding_mode='zeros',
        precision=-1, clamp_val=0.5):
        kernel_size = _pair(kernel_size)
        stride = _pair(stride)
        padding = _pair(padding)
        dilation = _pair(dilation)
        super().__init__(in_channels, out_channels, kernel_size, stride,
            padding, dilation, groups, bias, padding_mode)
        self.precision = precision
        self.clamp_val = clamp_val

    def conv2d_forward(self, input, weight):
        return F.conv2d(input, weight, self.bias, self.stride, self.padding,
            self.dilation, self.groups)

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


Measured latency_ms: 0.020997