import torch
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.utils.data
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_poi_fused_add_argmax_gather_index_put_lift_fresh_0(in_ptr0,
    out_ptr2, out_ptr4, out_ptr6, out_ptr8, out_ptr10, out_ptr12, out_ptr14,
    out_ptr16, out_ptr18, out_ptr19, out_ptr20, out_ptr21, out_ptr22,
    out_ptr23, out_ptr24, out_ptr25, out_ptr26, out_ptr27, xnumel, XBLOCK:
    tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 16
    x1 = xindex // 16
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 64 * x1), xmask)
    tmp1 = tl.load(in_ptr0 + (16 + x0 + 64 * x1), xmask)
    tmp17 = tl.load(in_ptr0 + (32 + x0 + 64 * x1), xmask)
    tmp32 = tl.load(in_ptr0 + (48 + x0 + 64 * x1), xmask)
    tmp2 = tmp0 > tmp1
    tmp3 = tmp0 == tmp1
    tmp4 = tmp0 != tmp0
    tmp5 = tmp1 != tmp1
    tmp6 = tmp4 > tmp5
    tmp7 = tmp2 | tmp6
    tmp8 = tmp4 & tmp5
    tmp9 = tmp3 | tmp8
    tmp10 = tl.full([1], 0, tl.int64)
    tmp11 = tl.full([1], 1, tl.int64)
    tmp12 = tmp10 < tmp11
    tmp13 = tmp9 & tmp12
    tmp14 = tmp7 | tmp13
    tmp15 = tl.where(tmp14, tmp0, tmp1)
    tmp16 = tl.where(tmp14, tmp10, tmp11)
    tmp18 = tmp15 > tmp17
    tmp19 = tmp15 == tmp17
    tmp20 = tmp15 != tmp15
    tmp21 = tmp17 != tmp17
    tmp22 = tmp20 > tmp21
    tmp23 = tmp18 | tmp22
    tmp24 = tmp20 & tmp21
    tmp25 = tmp19 | tmp24
    tmp26 = tl.full([1], 2, tl.int64)
    tmp27 = tmp16 < tmp26
    tmp28 = tmp25 & tmp27
    tmp29 = tmp23 | tmp28
    tmp30 = tl.where(tmp29, tmp15, tmp17)
    tmp31 = tl.where(tmp29, tmp16, tmp26)
    tmp33 = tmp30 > tmp32
    tmp34 = tmp30 == tmp32
    tmp35 = tmp30 != tmp30
    tmp36 = tmp32 != tmp32
    tmp37 = tmp35 > tmp36
    tmp38 = tmp33 | tmp37
    tmp39 = tmp35 & tmp36
    tmp40 = tmp34 | tmp39
    tmp41 = tl.full([1], 3, tl.int64)
    tmp42 = tmp31 < tmp41
    tmp43 = tmp40 & tmp42
    tmp44 = tmp38 | tmp43
    tl.where(tmp44, tmp30, tmp32)
    tmp46 = tl.where(tmp44, tmp31, tmp41)
    tmp47 = tl.full([1], -4, tl.int64)
    tmp48 = tmp46 + tmp47
    tmp49 = tmp48 < tmp10
    tmp50 = tl.where(tmp49, tmp10, tmp48)
    tmp51 = tmp50 > tmp41
    tmp52 = tl.where(tmp51, tmp41, tmp50)
    tmp53 = tl.full([1], -3, tl.int64)
    tmp54 = tmp46 + tmp53
    tmp55 = tmp54 < tmp10
    tmp56 = tl.where(tmp55, tmp10, tmp54)
    tmp57 = tmp56 > tmp41
    tmp58 = tl.where(tmp57, tmp41, tmp56)
    tmp59 = tl.full([1], -2, tl.int64)
    tmp60 = tmp46 + tmp59
    tmp61 = tmp60 < tmp10
    tmp62 = tl.where(tmp61, tmp10, tmp60)
    tmp63 = tmp62 > tmp41
    tmp64 = tl.where(tmp63, tmp41, tmp62)
    tmp65 = tl.full([1], -1, tl.int64)
    tmp66 = tmp46 + tmp65
    tmp67 = tmp66 < tmp10
    tmp68 = tl.where(tmp67, tmp10, tmp66)
    tmp69 = tmp68 > tmp41
    tmp70 = tl.where(tmp69, tmp41, tmp68)
    tmp71 = tmp46 + tmp10
    tmp72 = tmp71 < tmp10
    tmp73 = tl.where(tmp72, tmp10, tmp71)
    tmp74 = tmp73 > tmp41
    tmp75 = tl.where(tmp74, tmp41, tmp73)
    tmp76 = tmp46 + tmp11
    tmp77 = tmp76 < tmp10
    tmp78 = tl.where(tmp77, tmp10, tmp76)
    tmp79 = tmp78 > tmp41
    tmp80 = tl.where(tmp79, tmp41, tmp78)
    tmp81 = tmp46 + tmp26
    tmp82 = tmp81 < tmp10
    tmp83 = tl.where(tmp82, tmp10, tmp81)
    tmp84 = tmp83 > tmp41
    tmp85 = tl.where(tmp84, tmp41, tmp83)
    tmp86 = tmp46 + tmp41
    tmp87 = tmp86 < tmp10
    tmp88 = tl.where(tmp87, tmp10, tmp86)
    tmp89 = tmp88 > tmp41
    tmp90 = tl.where(tmp89, tmp41, tmp88)
    tmp91 = tl.full([1], 4, tl.int64)
    tmp92 = tmp46 + tmp91
    tmp93 = tmp92 < tmp10
    tmp94 = tl.where(tmp93, tmp10, tmp92)
    tmp95 = tmp94 > tmp41
    tmp96 = tl.where(tmp95, tmp41, tmp94)
    tmp97 = tl.full([XBLOCK], 4, tl.int32)
    tmp98 = tmp52 + tmp97
    tmp99 = tmp52 < 0
    tmp100 = tl.where(tmp99, tmp98, tmp52)
    tl.device_assert((0 <= tmp100) & (tmp100 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp100 < 4')
    tmp102 = tl.load(in_ptr0 + (x0 + 16 * tmp100 + 64 * x1), xmask)
    tmp103 = tmp58 + tmp97
    tmp104 = tmp58 < 0
    tmp105 = tl.where(tmp104, tmp103, tmp58)
    tl.device_assert((0 <= tmp105) & (tmp105 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp105 < 4')
    tmp107 = tl.load(in_ptr0 + (x0 + 16 * tmp105 + 64 * x1), xmask)
    tmp108 = tmp64 + tmp97
    tmp109 = tmp64 < 0
    tmp110 = tl.where(tmp109, tmp108, tmp64)
    tl.device_assert((0 <= tmp110) & (tmp110 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp110 < 4')
    tmp112 = tl.load(in_ptr0 + (x0 + 16 * tmp110 + 64 * x1), xmask)
    tmp113 = tmp70 + tmp97
    tmp114 = tmp70 < 0
    tmp115 = tl.where(tmp114, tmp113, tmp70)
    tl.device_assert((0 <= tmp115) & (tmp115 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp115 < 4')
    tmp117 = tl.load(in_ptr0 + (x0 + 16 * tmp115 + 64 * x1), xmask)
    tmp118 = tmp75 + tmp97
    tmp119 = tmp75 < 0
    tmp120 = tl.where(tmp119, tmp118, tmp75)
    tl.device_assert((0 <= tmp120) & (tmp120 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp120 < 4')
    tmp122 = tl.load(in_ptr0 + (x0 + 16 * tmp120 + 64 * x1), xmask)
    tmp123 = tmp80 + tmp97
    tmp124 = tmp80 < 0
    tmp125 = tl.where(tmp124, tmp123, tmp80)
    tl.device_assert((0 <= tmp125) & (tmp125 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp125 < 4')
    tmp127 = tl.load(in_ptr0 + (x0 + 16 * tmp125 + 64 * x1), xmask)
    tmp128 = tmp85 + tmp97
    tmp129 = tmp85 < 0
    tmp130 = tl.where(tmp129, tmp128, tmp85)
    tl.device_assert((0 <= tmp130) & (tmp130 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp130 < 4')
    tmp132 = tl.load(in_ptr0 + (x0 + 16 * tmp130 + 64 * x1), xmask)
    tmp133 = tmp90 + tmp97
    tmp134 = tmp90 < 0
    tmp135 = tl.where(tmp134, tmp133, tmp90)
    tl.device_assert((0 <= tmp135) & (tmp135 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp135 < 4')
    tmp137 = tl.load(in_ptr0 + (x0 + 16 * tmp135 + 64 * x1), xmask)
    tmp138 = tmp96 + tmp97
    tmp139 = tmp96 < 0
    tmp140 = tl.where(tmp139, tmp138, tmp96)
    tl.device_assert((0 <= tmp140) & (tmp140 < 4) | ~xmask,
        'index out of bounds: 0 <= tmp140 < 4')
    tmp142 = tl.load(in_ptr0 + (x0 + 16 * tmp140 + 64 * x1), xmask)
    tl.store(out_ptr2 + (x0 + 144 * x1), tmp52, xmask)
    tl.store(out_ptr4 + (x0 + 144 * x1), tmp58, xmask)
    tl.store(out_ptr6 + (x0 + 144 * x1), tmp64, xmask)
    tl.store(out_ptr8 + (x0 + 144 * x1), tmp70, xmask)
    tl.store(out_ptr10 + (x0 + 144 * x1), tmp75, xmask)
    tl.store(out_ptr12 + (x0 + 144 * x1), tmp80, xmask)
    tl.store(out_ptr14 + (x0 + 144 * x1), tmp85, xmask)
    tl.store(out_ptr16 + (x0 + 144 * x1), tmp90, xmask)
    tl.store(out_ptr18 + (x0 + 144 * x1), tmp96, xmask)
    tl.store(out_ptr19 + 9 * x2, tmp102, xmask)
    tl.store(out_ptr20 + 9 * x2, tmp107, xmask)
    tl.store(out_ptr21 + 9 * x2, tmp112, xmask)
    tl.store(out_ptr22 + 9 * x2, tmp117, xmask)
    tl.store(out_ptr23 + 9 * x2, tmp122, xmask)
    tl.store(out_ptr24 + 9 * x2, tmp127, xmask)
    tl.store(out_ptr25 + 9 * x2, tmp132, xmask)
    tl.store(out_ptr26 + 9 * x2, tmp137, xmask)
    tl.store(out_ptr27 + 9 * x2, tmp142, xmask)


@triton.jit
def triton_per_fused__to_copy_add_div_mul_sum_1(in_out_ptr0, in_ptr0,
    in_ptr1, xnumel, rnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    rnumel = 9
    RBLOCK: tl.constexpr = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    rmask = rindex < rnumel
    r1 = rindex
    x0 = xindex
    x2 = xindex % 16
    x3 = xindex // 16
    tmp0 = tl.load(in_ptr0 + (r1 + 9 * x0), rmask & xmask, other=0.0)
    tmp8 = tl.load(in_ptr1 + (x2 + 16 * r1 + 144 * x3), rmask & xmask,
        other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.where(rmask & xmask, tmp1, 0)
    tmp4 = tl.sum(tmp3, 1)[:, None]
    tmp5 = 1e-08
    tmp6 = tmp4 + tmp5
    tmp7 = tmp0 / tmp6
    tmp9 = tmp8.to(tl.float32)
    tmp10 = tmp7 * tmp9
    tmp11 = tl.broadcast_to(tmp10, [XBLOCK, RBLOCK])
    tmp13 = tl.where(rmask & xmask, tmp11, 0)
    tmp14 = tl.sum(tmp13, 1)[:, None]
    tl.store(in_out_ptr0 + x0, tmp14, xmask)


def call(args):
    arg0_1, = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4, 4, 4), (64, 16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf30 = empty_strided_cuda((4, 9, 4, 4), (144, 16, 4, 1), torch.int64)
        buf2 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 0)
        buf4 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 16)
        buf6 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 32)
        buf8 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 48)
        buf10 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 64)
        buf12 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 80)
        buf14 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 96)
        buf16 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 112)
        buf18 = reinterpret_tensor(buf30, (4, 1, 4, 4), (144, 16, 4, 1), 128)
        buf28 = empty_strided_cuda((4, 9, 4, 4), (144, 1, 36, 9), torch.float32
            )
        buf19 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 0)
        buf20 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 1)
        buf21 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 2)
        buf22 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 3)
        buf23 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 4)
        buf24 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 5)
        buf25 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 6)
        buf26 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 7)
        buf27 = reinterpret_tensor(buf28, (4, 1, 4, 4), (144, 1, 36, 9), 8)
        get_raw_stream(0)
        triton_poi_fused_add_argmax_gather_index_put_lift_fresh_0[grid(64)](
            arg0_1, buf2, buf4, buf6, buf8, buf10, buf12, buf14, buf16,
            buf18, buf19, buf20, buf21, buf22, buf23, buf24, buf25, buf26,
            buf27, 64, XBLOCK=64, num_warps=1, num_stages=1)
        del arg0_1
        buf29 = empty_strided_cuda((4, 1, 4, 4), (16, 64, 4, 1), torch.float32)
        buf31 = reinterpret_tensor(buf29, (4, 4, 4), (16, 4, 1), 0)
        del buf29
        triton_per_fused__to_copy_add_div_mul_sum_1[grid(64)](buf31, buf28,
            buf30, 64, 9, XBLOCK=8, num_warps=2, num_stages=1)
        del buf10
        del buf12
        del buf14
        del buf16
        del buf18
        del buf19
        del buf2
        del buf20
        del buf21
        del buf22
        del buf23
        del buf24
        del buf25
        del buf26
        del buf27
        del buf28
        del buf30
        del buf4
        del buf6
        del buf8
    return buf31,


class DisparityRegressionNew(nn.Module):

    def __init__(self, maxdisp, win_size):
        super(DisparityRegressionNew, self).__init__()
        self.max_disp = maxdisp
        self.win_size = win_size

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


Measured latency_ms: 0.057877