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
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_poi_fused_convolution_relu_0(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)
    x3 = xindex
    x1 = xindex // 4096 % 256
    tmp0 = tl.load(in_out_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr0 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tl.store(in_out_ptr0 + x3, tmp4, None)


@triton.jit
def triton_poi_fused_convolution_relu_1(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)
    x3 = xindex
    x1 = xindex // 4096 % 128
    tmp0 = tl.load(in_out_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr0 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tl.store(in_out_ptr0 + x3, tmp4, None)


@triton.jit
def triton_poi_fused_cat_2(in_ptr0, in_ptr1, in_ptr2, in_ptr3, out_ptr0,
    xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x1 = xindex // 4096 % 256
    x0 = xindex % 4096
    x2 = xindex // 1048576
    x3 = xindex
    tmp0 = x1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 192, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (x0 + 4096 * x1 + 786432 * x2), tmp4, other=0.0)
    tmp6 = tl.load(in_ptr1 + x1, tmp4, eviction_policy='evict_last', other=0.0)
    tmp7 = tmp5 + tmp6
    tmp8 = tl.full([1], 0, tl.int32)
    tmp9 = triton_helpers.maximum(tmp8, tmp7)
    tmp10 = tl.full(tmp9.shape, 0.0, tmp9.dtype)
    tmp11 = tl.where(tmp4, tmp9, tmp10)
    tmp12 = tmp0 >= tmp3
    tl.full([1], 256, tl.int64)
    tmp15 = tl.load(in_ptr2 + (x0 + 4096 * (-192 + x1) + 262144 * x2),
        tmp12, other=0.0)
    tmp16 = tl.load(in_ptr3 + (-192 + x1), tmp12, eviction_policy=
        'evict_last', other=0.0)
    tmp17 = tmp15 + tmp16
    tmp18 = triton_helpers.maximum(tmp8, tmp17)
    tmp19 = tl.full(tmp18.shape, 0.0, tmp18.dtype)
    tmp20 = tl.where(tmp12, tmp18, tmp19)
    tmp21 = tl.where(tmp4, tmp11, tmp20)
    tl.store(out_ptr0 + x3, tmp21, None)


@triton.jit
def triton_poi_fused_cat_3(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)
    x1 = xindex // 4096 % 128
    x0 = xindex % 4096
    x2 = xindex // 524288
    x3 = xindex
    tmp0 = x1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 126, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (x0 + 4096 * x1 + 516096 * x2), tmp4, other=0.0)
    tmp6 = tl.load(in_ptr1 + x1, tmp4, eviction_policy='evict_last', other=0.0)
    tmp7 = tmp5 + tmp6
    tmp8 = tl.full([1], 0, tl.int32)
    tmp9 = triton_helpers.maximum(tmp8, tmp7)
    tmp10 = tl.full(tmp9.shape, 0.0, tmp9.dtype)
    tmp11 = tl.where(tmp4, tmp9, tmp10)
    tmp12 = tmp0 >= tmp3
    tl.full([1], 128, tl.int64)
    tmp15 = tl.load(in_ptr2 + (x0 + 4096 * (-126 + x1) + 8192 * x2), tmp12,
        other=0.0)
    tmp16 = tl.where(tmp4, tmp11, tmp15)
    tl.store(out_ptr0 + x3, tmp16, None)


@triton.jit
def triton_poi_fused_convolution_relu_threshold_backward_4(in_ptr0, in_ptr1,
    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
    x1 = xindex // 4096 % 126
    tmp0 = tl.load(in_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr1 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tmp5 = 0.0
    tmp6 = tmp4 <= tmp5
    tl.store(out_ptr0 + x3, tmp6, None)


@triton.jit
def triton_poi_fused_convolution_relu_threshold_backward_5(in_ptr0, in_ptr1,
    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
    x1 = xindex // 4096 % 64
    tmp0 = tl.load(in_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr1 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tmp5 = 0.0
    tmp6 = tmp4 <= tmp5
    tl.store(out_ptr0 + x3, tmp6, None)


@triton.jit
def triton_poi_fused_convolution_relu_threshold_backward_6(in_ptr0, in_ptr1,
    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
    x1 = xindex // 4096 % 192
    tmp0 = tl.load(in_ptr0 + x3, None)
    tmp1 = tl.load(in_ptr1 + x1, None, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tmp5 = 0.0
    tmp6 = tmp4 <= tmp5
    tl.store(out_ptr0 + x3, tmp6, None)


def call(args):
    (primals_1, primals_2, primals_3, primals_4, primals_5, primals_6,
        primals_7, primals_8, primals_9, primals_10, primals_11, primals_12
        ) = args
    args.clear()
    assert_size_stride(primals_1, (256, 324, 1, 1), (324, 1, 1, 1))
    assert_size_stride(primals_2, (256,), (1,))
    assert_size_stride(primals_3, (4, 324, 64, 64), (1327104, 4096, 64, 1))
    assert_size_stride(primals_4, (192, 256, 3, 3), (2304, 9, 3, 1))
    assert_size_stride(primals_5, (192,), (1,))
    assert_size_stride(primals_6, (128, 2, 7, 7), (98, 49, 7, 1))
    assert_size_stride(primals_7, (128,), (1,))
    assert_size_stride(primals_8, (4, 2, 64, 64), (8192, 4096, 64, 1))
    assert_size_stride(primals_9, (64, 128, 3, 3), (1152, 9, 3, 1))
    assert_size_stride(primals_10, (64,), (1,))
    assert_size_stride(primals_11, (126, 256, 3, 3), (2304, 9, 3, 1))
    assert_size_stride(primals_12, (126,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = extern_kernels.convolution(primals_3, primals_1, stride=(1, 
            1), padding=(0, 0), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf0, (4, 256, 64, 64), (1048576, 4096, 64, 1))
        buf1 = buf0
        del buf0
        get_raw_stream(0)
        triton_poi_fused_convolution_relu_0[grid(4194304)](buf1, primals_2,
            4194304, XBLOCK=1024, num_warps=4, num_stages=1)
        del primals_2
        buf2 = extern_kernels.convolution(buf1, primals_4, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf2, (4, 192, 64, 64), (786432, 4096, 64, 1))
        buf3 = extern_kernels.convolution(primals_8, primals_6, stride=(1, 
            1), padding=(3, 3), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf3, (4, 128, 64, 64), (524288, 4096, 64, 1))
        buf4 = buf3
        del buf3
        triton_poi_fused_convolution_relu_1[grid(2097152)](buf4, primals_7,
            2097152, XBLOCK=1024, num_warps=4, num_stages=1)
        del primals_7
        buf5 = extern_kernels.convolution(buf4, primals_9, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf5, (4, 64, 64, 64), (262144, 4096, 64, 1))
        buf6 = empty_strided_cuda((4, 256, 64, 64), (1048576, 4096, 64, 1),
            torch.float32)
        triton_poi_fused_cat_2[grid(4194304)](buf2, primals_5, buf5,
            primals_10, buf6, 4194304, XBLOCK=1024, num_warps=4, num_stages=1)
        buf7 = extern_kernels.convolution(buf6, primals_11, stride=(1, 1),
            padding=(1, 1), dilation=(1, 1), transposed=False,
            output_padding=(0, 0), groups=1, bias=None)
        assert_size_stride(buf7, (4, 126, 64, 64), (516096, 4096, 64, 1))
        buf8 = empty_strided_cuda((4, 128, 64, 64), (524288, 4096, 64, 1),
            torch.float32)
        triton_poi_fused_cat_3[grid(2097152)](buf7, primals_12, primals_8,
            buf8, 2097152, XBLOCK=1024, num_warps=4, num_stages=1)
        buf9 = empty_strided_cuda((4, 126, 64, 64), (516096, 4096, 64, 1),
            torch.bool)
        triton_poi_fused_convolution_relu_threshold_backward_4[grid(2064384)](
            buf7, primals_12, buf9, 2064384, XBLOCK=512, num_warps=8,
            num_stages=1)
        del buf7
        del primals_12
        buf10 = empty_strided_cuda((4, 64, 64, 64), (262144, 4096, 64, 1),
            torch.bool)
        triton_poi_fused_convolution_relu_threshold_backward_5[grid(1048576)](
            buf5, primals_10, buf10, 1048576, XBLOCK=1024, num_warps=4,
            num_stages=1)
        del buf5
        del primals_10
        buf11 = empty_strided_cuda((4, 192, 64, 64), (786432, 4096, 64, 1),
            torch.bool)
        triton_poi_fused_convolution_relu_threshold_backward_6[grid(3145728)](
            buf2, primals_5, buf11, 3145728, XBLOCK=1024, num_warps=4,
            num_stages=1)
        del buf2
        del primals_5
    return (buf8, primals_1, primals_3, primals_4, primals_6, primals_8,
        primals_9, primals_11, buf1, buf4, buf6, buf9, buf10, buf11)


class MotionEncoderNew(nn.Module):
    """
    Encodes motion features from the correlation levels of the pyramid
    and the input flow estimate using convolution layers.


    Parameters
    ----------
    corr_radius : int
        Correlation radius of the correlation pyramid
    corr_levels : int
        Correlation levels of the correlation pyramid

    """

    def __init__(self, corr_radius, corr_levels):
        super(MotionEncoderNew, self).__init__()
        cor_planes = corr_levels * (2 * corr_radius + 1) ** 2
        self.convc1 = nn.Conv2d(cor_planes, 256, 1, padding=0)
        self.convc2 = nn.Conv2d(256, 192, 3, padding=1)
        self.convf1 = nn.Conv2d(2, 128, 7, padding=3)
        self.convf2 = nn.Conv2d(128, 64, 3, padding=1)
        self.conv = nn.Conv2d(64 + 192, 128 - 2, 3, padding=1)

    def forward(self, input_0, input_1):
        primals_1 = self.convc1.weight
        primals_2 = self.convc1.bias
        primals_4 = self.convc2.weight
        primals_5 = self.convc2.bias
        primals_6 = self.convf1.weight
        primals_7 = self.convf1.bias
        primals_9 = self.convf2.weight
        primals_10 = self.convf2.bias
        primals_11 = self.conv.weight
        primals_12 = self.conv.bias
        primals_8 = input_0
        primals_3 = input_1
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7, primals_8, primals_9,
            primals_10, primals_11, primals_12])
        return output[0]


Measured latency_ms: 0.042885