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
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch import nn
import torch.nn.functional as F
import torch.utils.data
from torch.nn.modules.loss import _Loss
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor


@triton.jit
def triton_per_fused_abs_cat_mean_sub_0(in_out_ptr0, in_ptr0, in_ptr1,
    in_ptr2, in_ptr3, in_ptr4, in_ptr5, in_ptr6, in_ptr7, in_ptr8, in_ptr9,
    in_ptr10, in_ptr11, xnumel, rnumel):
    XBLOCK: tl.constexpr = 1
    rnumel = 432
    RBLOCK: tl.constexpr = 512
    xoffset = tl.program_id(0) * XBLOCK
    tl.full([1], xoffset, tl.int32)
    tl.full([RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[:]
    rmask = rindex < rnumel
    r1 = rindex // 36 % 3
    r0 = rindex % 36
    r2 = rindex // 108
    tmp0 = r1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 1, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), rmask &
        tmp4, eviction_policy='evict_last', other=0.0)
    tmp6 = tmp5 * tmp5
    tmp7 = tl.load(in_ptr1 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), rmask &
        tmp4, eviction_policy='evict_last', other=0.0)
    tmp8 = tmp7 * tmp7
    tmp9 = tmp6 + tmp8
    tmp10 = 1e-06
    tmp11 = tmp9 + tmp10
    tmp12 = libdevice.sqrt(tmp11)
    tmp13 = tl.full(tmp12.shape, 0.0, tmp12.dtype)
    tmp14 = tl.where(tmp4, tmp12, tmp13)
    tmp15 = tmp0 >= tmp3
    tmp16 = tl.full([1], 2, tl.int64)
    tmp17 = tmp0 < tmp16
    tmp18 = tmp15 & tmp17
    tmp19 = tl.load(in_ptr2 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp18, eviction_policy='evict_last', other=0.0)
    tmp20 = tmp19 * tmp19
    tmp21 = tl.load(in_ptr3 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp18, eviction_policy='evict_last', other=0.0)
    tmp22 = tmp21 * tmp21
    tmp23 = tmp20 + tmp22
    tmp24 = tmp23 + tmp10
    tmp25 = libdevice.sqrt(tmp24)
    tmp26 = tl.full(tmp25.shape, 0.0, tmp25.dtype)
    tmp27 = tl.where(tmp18, tmp25, tmp26)
    tmp28 = tmp0 >= tmp16
    tl.full([1], 3, tl.int64)
    tmp31 = tl.load(in_ptr4 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp28, eviction_policy='evict_last', other=0.0)
    tmp32 = tmp31 * tmp31
    tmp33 = tl.load(in_ptr5 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp28, eviction_policy='evict_last', other=0.0)
    tmp34 = tmp33 * tmp33
    tmp35 = tmp32 + tmp34
    tmp36 = tmp35 + tmp10
    tmp37 = libdevice.sqrt(tmp36)
    tmp38 = tl.full(tmp37.shape, 0.0, tmp37.dtype)
    tmp39 = tl.where(tmp28, tmp37, tmp38)
    tmp40 = tl.where(tmp18, tmp27, tmp39)
    tmp41 = tl.where(tmp4, tmp14, tmp40)
    tmp42 = tl.load(in_ptr6 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp4, eviction_policy='evict_last', other=0.0)
    tmp43 = tmp42 * tmp42
    tmp44 = tl.load(in_ptr7 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp4, eviction_policy='evict_last', other=0.0)
    tmp45 = tmp44 * tmp44
    tmp46 = tmp43 + tmp45
    tmp47 = tmp46 + tmp10
    tmp48 = libdevice.sqrt(tmp47)
    tmp49 = tl.full(tmp48.shape, 0.0, tmp48.dtype)
    tmp50 = tl.where(tmp4, tmp48, tmp49)
    tmp51 = tl.load(in_ptr8 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp18, eviction_policy='evict_last', other=0.0)
    tmp52 = tmp51 * tmp51
    tmp53 = tl.load(in_ptr9 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp18, eviction_policy='evict_last', other=0.0)
    tmp54 = tmp53 * tmp53
    tmp55 = tmp52 + tmp54
    tmp56 = tmp55 + tmp10
    tmp57 = libdevice.sqrt(tmp56)
    tmp58 = tl.full(tmp57.shape, 0.0, tmp57.dtype)
    tmp59 = tl.where(tmp18, tmp57, tmp58)
    tmp60 = tl.load(in_ptr10 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp28, eviction_policy='evict_last', other=0.0)
    tmp61 = tmp60 * tmp60
    tmp62 = tl.load(in_ptr11 + tl.broadcast_to(r0 + 36 * r2, [RBLOCK]), 
        rmask & tmp28, eviction_policy='evict_last', other=0.0)
    tmp63 = tmp62 * tmp62
    tmp64 = tmp61 + tmp63
    tmp65 = tmp64 + tmp10
    tmp66 = libdevice.sqrt(tmp65)
    tmp67 = tl.full(tmp66.shape, 0.0, tmp66.dtype)
    tmp68 = tl.where(tmp28, tmp66, tmp67)
    tmp69 = tl.where(tmp18, tmp59, tmp68)
    tmp70 = tl.where(tmp4, tmp50, tmp69)
    tmp71 = tmp41 - tmp70
    tmp72 = tl_math.abs(tmp71)
    tmp73 = tl.broadcast_to(tmp72, [RBLOCK])
    tmp75 = tl.where(rmask, tmp73, 0)
    tmp76 = triton_helpers.promote_to_tensor(tl.sum(tmp75, 0))
    tmp77 = 432.0
    tmp78 = tmp76 / tmp77
    tl.debug_barrier()
    tl.store(in_out_ptr0 + tl.full([1], 0, tl.int32), tmp78, None)


def call(args):
    arg0_1, arg1_1, arg2_1, arg3_1 = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(arg1_1, (1, 1, 3, 3), (9, 9, 3, 1))
    assert_size_stride(arg2_1, (1, 1, 3, 3), (9, 9, 3, 1))
    assert_size_stride(arg3_1, (4, 4, 4, 4), (64, 16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = extern_kernels.convolution(reinterpret_tensor(arg0_1, (4, 1,
            4, 4), (64, 0, 4, 1), 0), arg1_1, stride=(1, 1), padding=(2, 2),
            dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf0, (4, 1, 6, 6), (36, 36, 6, 1))
        buf1 = extern_kernels.convolution(reinterpret_tensor(arg0_1, (4, 1,
            4, 4), (64, 0, 4, 1), 0), arg2_1, stride=(1, 1), padding=(2, 2),
            dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf1, (4, 1, 6, 6), (36, 36, 6, 1))
        buf2 = extern_kernels.convolution(reinterpret_tensor(arg0_1, (4, 1,
            4, 4), (64, 0, 4, 1), 16), arg1_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf2, (4, 1, 6, 6), (36, 36, 6, 1))
        buf3 = extern_kernels.convolution(reinterpret_tensor(arg0_1, (4, 1,
            4, 4), (64, 0, 4, 1), 16), arg2_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf3, (4, 1, 6, 6), (36, 36, 6, 1))
        buf4 = extern_kernels.convolution(reinterpret_tensor(arg0_1, (4, 1,
            4, 4), (64, 0, 4, 1), 32), arg1_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf4, (4, 1, 6, 6), (36, 36, 6, 1))
        buf5 = extern_kernels.convolution(reinterpret_tensor(arg0_1, (4, 1,
            4, 4), (64, 0, 4, 1), 32), arg2_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf5, (4, 1, 6, 6), (36, 36, 6, 1))
        del arg0_1
        buf10 = extern_kernels.convolution(reinterpret_tensor(arg3_1, (4, 1,
            4, 4), (64, 0, 4, 1), 16), arg2_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf10, (4, 1, 6, 6), (36, 36, 6, 1))
        buf11 = extern_kernels.convolution(reinterpret_tensor(arg3_1, (4, 1,
            4, 4), (64, 0, 4, 1), 32), arg1_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf11, (4, 1, 6, 6), (36, 36, 6, 1))
        buf12 = extern_kernels.convolution(reinterpret_tensor(arg3_1, (4, 1,
            4, 4), (64, 0, 4, 1), 32), arg2_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf12, (4, 1, 6, 6), (36, 36, 6, 1))
        buf7 = extern_kernels.convolution(reinterpret_tensor(arg3_1, (4, 1,
            4, 4), (64, 0, 4, 1), 0), arg1_1, stride=(1, 1), padding=(2, 2),
            dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf7, (4, 1, 6, 6), (36, 36, 6, 1))
        buf8 = extern_kernels.convolution(reinterpret_tensor(arg3_1, (4, 1,
            4, 4), (64, 0, 4, 1), 0), arg2_1, stride=(1, 1), padding=(2, 2),
            dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf8, (4, 1, 6, 6), (36, 36, 6, 1))
        del arg2_1
        buf9 = extern_kernels.convolution(reinterpret_tensor(arg3_1, (4, 1,
            4, 4), (64, 0, 4, 1), 16), arg1_1, stride=(1, 1), padding=(2, 2
            ), dilation=(1, 1), transposed=False, output_padding=(0, 0),
            groups=1, bias=None)
        assert_size_stride(buf9, (4, 1, 6, 6), (36, 36, 6, 1))
        del arg1_1
        del arg3_1
        buf14 = empty_strided_cuda((), (), torch.float32)
        buf15 = buf14
        del buf14
        get_raw_stream(0)
        triton_per_fused_abs_cat_mean_sub_0[grid(1)](buf15, buf0, buf1,
            buf2, buf3, buf4, buf5, buf7, buf8, buf9, buf10, buf11, buf12, 
            1, 432, num_warps=4, num_stages=1)
        del buf0
        del buf1
        del buf10
        del buf11
        del buf12
        del buf2
        del buf3
        del buf4
        del buf5
        del buf7
        del buf8
        del buf9
    return buf15,


class Gradient(nn.Module):

    def __init__(self):
        super(Gradient, self).__init__()
        kernel_v = [[0, -1, 0], [0, 0, 0], [0, 1, 0]]
        kernel_h = [[0, 0, 0], [-1, 0, 1], [0, 0, 0]]
        kernel_h = torch.FloatTensor(kernel_h).unsqueeze(0).unsqueeze(0)
        kernel_v = torch.FloatTensor(kernel_v).unsqueeze(0).unsqueeze(0)
        self.weight_h = nn.Parameter(data=kernel_h, requires_grad=False)
        self.weight_v = nn.Parameter(data=kernel_v, requires_grad=False)

    def forward(self, x):
        x0 = x[:, 0]
        x1 = x[:, 1]
        x2 = x[:, 2]
        x0_v = F.conv2d(x0.unsqueeze(1), self.weight_v, padding=2)
        x0_h = F.conv2d(x0.unsqueeze(1), self.weight_h, padding=2)
        x1_v = F.conv2d(x1.unsqueeze(1), self.weight_v, padding=2)
        x1_h = F.conv2d(x1.unsqueeze(1), self.weight_h, padding=2)
        x2_v = F.conv2d(x2.unsqueeze(1), self.weight_v, padding=2)
        x2_h = F.conv2d(x2.unsqueeze(1), self.weight_h, padding=2)
        x0 = torch.sqrt(torch.pow(x0_v, 2) + torch.pow(x0_h, 2) + 1e-06)
        x1 = torch.sqrt(torch.pow(x1_v, 2) + torch.pow(x1_h, 2) + 1e-06)
        x2 = torch.sqrt(torch.pow(x2_v, 2) + torch.pow(x2_h, 2) + 1e-06)
        x = torch.cat([x0, x1, x2], dim=1)
        return x


class L1GradientLossNew(_Loss):
    """
    Gradient loss
    """

    def __init__(self, para):
        super(L1GradientLossNew, self).__init__()
        self.get_grad = Gradient()
        self.L1 = nn.L1Loss()

    def forward(self, input_0, input_1):
        arg1_1 = self.get_grad.weight_h
        arg2_1 = self.get_grad.weight_v
        arg0_1 = input_0
        arg3_1 = input_1
        output = call([arg0_1, arg1_1, arg2_1, arg3_1])
        return output[0]


Measured latency_ms: 0.01522