# AOT ID: ['0_forward']
from ctypes import c_void_p, c_long, c_int
import torch
import math
import random
import os
import tempfile
from math import inf, nan
from torch._inductor.hooks import run_intermediate_hooks
from torch._inductor.utils import maybe_profile
from torch._inductor.codegen.memory_planning import _align as align
from torch import device, empty_strided
from torch._inductor.async_compile import AsyncCompile
from torch._inductor.select_algorithm import extern_kernels
from torch._inductor.codegen.multi_kernel import MultiKernelCall
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid, split_scan_grid, grid_combo_kernels, start_graph, end_graph
from torch._C import _cuda_getCurrentRawStream as get_raw_stream

aten = torch.ops.aten
inductor_ops = torch.ops.inductor
_quantized = torch.ops._quantized
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
alloc_from_pool = torch.ops.inductor._alloc_from_pool
async_compile = AsyncCompile()


# kernel path: runs/run_shard_6/inductor_cache/pg/cpg5vjqssrlsnglazs6sllqz33jc7cdnlonls4f4kcehzuhy7ewb.py
# Topologically Sorted Source Nodes: [feats_norm], Original ATen: [aten.linalg_vector_norm]
# Source node to ATen node mapping:
#   feats_norm => pow_1, pow_2, sum_1
# Graph fragment:
#   %pow_1 : [num_users=1] = call_function[target=torch.ops.aten.pow.Tensor_Scalar](args = (%primals_1, 2), kwargs = {})
#   %sum_1 : [num_users=1] = call_function[target=torch.ops.aten.sum.dim_IntList](args = (%pow_1, [2], True), kwargs = {})
#   %pow_2 : [num_users=2] = call_function[target=torch.ops.aten.pow.Tensor_Scalar](args = (%sum_1, 0.5), kwargs = {})
triton_poi_fused_linalg_vector_norm_0 = async_compile.triton('triton_poi_fused_linalg_vector_norm_0', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[16], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_linalg_vector_norm_0', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 4, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_linalg_vector_norm_0(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (4*x0), xmask, eviction_policy='evict_last')
    tmp2 = tl.load(in_ptr0 + (1 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr0 + (2 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp8 = tl.load(in_ptr0 + (3 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp1 = tmp0 * tmp0
    tmp3 = tmp2 * tmp2
    tmp4 = tmp1 + tmp3
    tmp6 = tmp5 * tmp5
    tmp7 = tmp4 + tmp6
    tmp9 = tmp8 * tmp8
    tmp10 = tmp7 + tmp9
    tmp11 = libdevice.sqrt(tmp10)
    tl.store(out_ptr0 + (x0), tmp11, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_6/inductor_cache/n5/cn5teeeignn4sp3kp3ujrwymtwcb7yaywvx3h7kpqyi4ixhma2ot.py
# Topologically Sorted Source Nodes: [pairwise_sim], Original ATen: [aten.div]
# Source node to ATen node mapping:
#   pairwise_sim => div
# Graph fragment:
#   %div : [num_users=2] = call_function[target=torch.ops.aten.div.Tensor](args = (%bmm_1, %bmm), kwargs = {})
triton_poi_fused_div_1 = async_compile.triton('triton_poi_fused_div_1', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[64], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_div_1', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 2, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_div_1(in_out_ptr0, in_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (x0), xmask)
    tmp1 = tl.load(in_out_ptr0 + (x0), xmask)
    tmp2 = tmp0 / tmp1
    tl.store(in_out_ptr0 + (x0), tmp2, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_6/inductor_cache/zt/cztirch7jrz5e7flqpqcr5acbutu5sz47zvhuaocrfzdpb7umhxt.py
# Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_1, sub, mul_2, sub_1, pairwise_potential_E], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
# Source node to ATen node mapping:
#   W_sym => div_1
#   add => add
#   mul_1 => mul_1
#   mul_2 => mul_2
#   pairwise_potential => mul
#   pairwise_potential_E => sum_2
#   sub => sub
#   sub_1 => sub_1
# Graph fragment:
#   %add : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%primals_2, %permute_2), kwargs = {})
#   %div_1 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%add, 2), kwargs = {})
#   %mul : [num_users=20] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div, %div_1), kwargs = {})
#   %mul_1 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%permute_3, %mul), kwargs = {})
#   %sub : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (1, %permute_3), kwargs = {})
#   %mul_2 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub, %mul), kwargs = {})
#   %sub_1 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%mul_1, %mul_2), kwargs = {})
#   %sum_2 : [num_users=1] = call_function[target=torch.ops.aten.sum.dim_IntList](args = (%sub_1, [2], True), kwargs = {})
triton_poi_fused_add_div_mul_rsub_sub_sum_2 = async_compile.triton('triton_poi_fused_add_div_mul_rsub_sub_sum_2', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[16], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_div_mul_rsub_sub_sum_2', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 16, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_div_mul_rsub_sub_sum_2(in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = (xindex // 4)
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (16*x1)), xmask)
    tmp2 = tl.load(in_ptr1 + (4*x2), xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + (4*x0), xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr2 + (x0), xmask, eviction_policy='evict_last')
    tmp14 = tl.load(in_ptr0 + (4 + x0 + (16*x1)), xmask)
    tmp16 = tl.load(in_ptr1 + (1 + (4*x2)), xmask, eviction_policy='evict_last')
    tmp17 = tl.load(in_ptr2 + (1 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp18 = tl.load(in_ptr2 + (4 + x0), xmask, eviction_policy='evict_last')
    tmp27 = tl.load(in_ptr0 + (8 + x0 + (16*x1)), xmask)
    tmp29 = tl.load(in_ptr1 + (2 + (4*x2)), xmask, eviction_policy='evict_last')
    tmp30 = tl.load(in_ptr2 + (2 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp31 = tl.load(in_ptr2 + (8 + x0), xmask, eviction_policy='evict_last')
    tmp40 = tl.load(in_ptr0 + (12 + x0 + (16*x1)), xmask)
    tmp42 = tl.load(in_ptr1 + (3 + (4*x2)), xmask, eviction_policy='evict_last')
    tmp43 = tl.load(in_ptr2 + (3 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp44 = tl.load(in_ptr2 + (12 + x0), xmask, eviction_policy='evict_last')
    tmp1 = tl.sigmoid(tmp0)
    tmp5 = tmp3 + tmp4
    tmp6 = 0.5
    tmp7 = tmp5 * tmp6
    tmp8 = tmp2 * tmp7
    tmp9 = tmp1 * tmp8
    tmp10 = 1.0
    tmp11 = tmp10 - tmp1
    tmp12 = tmp11 * tmp8
    tmp13 = tmp9 - tmp12
    tmp15 = tl.sigmoid(tmp14)
    tmp19 = tmp17 + tmp18
    tmp20 = tmp19 * tmp6
    tmp21 = tmp16 * tmp20
    tmp22 = tmp15 * tmp21
    tmp23 = tmp10 - tmp15
    tmp24 = tmp23 * tmp21
    tmp25 = tmp22 - tmp24
    tmp26 = tmp13 + tmp25
    tmp28 = tl.sigmoid(tmp27)
    tmp32 = tmp30 + tmp31
    tmp33 = tmp32 * tmp6
    tmp34 = tmp29 * tmp33
    tmp35 = tmp28 * tmp34
    tmp36 = tmp10 - tmp28
    tmp37 = tmp36 * tmp34
    tmp38 = tmp35 - tmp37
    tmp39 = tmp26 + tmp38
    tmp41 = tl.sigmoid(tmp40)
    tmp45 = tmp43 + tmp44
    tmp46 = tmp45 * tmp6
    tmp47 = tmp42 * tmp46
    tmp48 = tmp41 * tmp47
    tmp49 = tmp10 - tmp41
    tmp50 = tmp49 * tmp47
    tmp51 = tmp48 - tmp50
    tmp52 = tmp39 + tmp51
    tl.store(out_ptr0 + (x2), tmp52, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_6/inductor_cache/4y/c4y7ef5neroioo45nhdxozozllc2bmzg27jz2uy5fyq34ehhexcn.py
# Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_3, sub_2, mul_4, sub_3, pairwise_potential_E_1], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
# Source node to ATen node mapping:
#   W_sym => div_1
#   add => add
#   mul_3 => mul_3
#   mul_4 => mul_4
#   pairwise_potential => mul
#   pairwise_potential_E_1 => sum_3
#   sub_2 => sub_2
#   sub_3 => sub_3
# Graph fragment:
#   %add : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%primals_2, %permute_2), kwargs = {})
#   %div_1 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%add, 2), kwargs = {})
#   %mul : [num_users=20] = call_function[target=torch.ops.aten.mul.Tensor](args = (%div, %div_1), kwargs = {})
#   %mul_3 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%permute_4, %mul), kwargs = {})
#   %sub_2 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (1, %permute_4), kwargs = {})
#   %mul_4 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%sub_2, %mul), kwargs = {})
#   %sub_3 : [num_users=1] = call_function[target=torch.ops.aten.sub.Tensor](args = (%mul_3, %mul_4), kwargs = {})
#   %sum_3 : [num_users=1] = call_function[target=torch.ops.aten.sum.dim_IntList](args = (%sub_3, [2], True), kwargs = {})
triton_poi_fused_add_div_mul_rsub_sub_sum_3 = async_compile.triton('triton_poi_fused_add_div_mul_rsub_sub_sum_3', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[16], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: '*fp32', 4: '*fp32', 5: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3, 4, 5), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_div_mul_rsub_sub_sum_3', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 20, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_div_mul_rsub_sub_sum_3(in_ptr0, in_ptr1, in_ptr2, in_ptr3, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = (xindex // 4)
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + (16*x1)), xmask)
    tmp1 = tl.load(in_ptr1 + (4*x1), xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr2 + (4*x2), xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + (4*x0), xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr3 + (x0), xmask, eviction_policy='evict_last')
    tmp16 = tl.load(in_ptr0 + (4 + x0 + (16*x1)), xmask)
    tmp17 = tl.load(in_ptr1 + (1 + (4*x1)), xmask, eviction_policy='evict_last')
    tmp20 = tl.load(in_ptr2 + (1 + (4*x2)), xmask, eviction_policy='evict_last')
    tmp21 = tl.load(in_ptr3 + (1 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp22 = tl.load(in_ptr3 + (4 + x0), xmask, eviction_policy='evict_last')
    tmp31 = tl.load(in_ptr0 + (8 + x0 + (16*x1)), xmask)
    tmp32 = tl.load(in_ptr1 + (2 + (4*x1)), xmask, eviction_policy='evict_last')
    tmp35 = tl.load(in_ptr2 + (2 + (4*x2)), xmask, eviction_policy='evict_last')
    tmp36 = tl.load(in_ptr3 + (2 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp37 = tl.load(in_ptr3 + (8 + x0), xmask, eviction_policy='evict_last')
    tmp46 = tl.load(in_ptr0 + (12 + x0 + (16*x1)), xmask)
    tmp47 = tl.load(in_ptr1 + (3 + (4*x1)), xmask, eviction_policy='evict_last')
    tmp50 = tl.load(in_ptr2 + (3 + (4*x2)), xmask, eviction_policy='evict_last')
    tmp51 = tl.load(in_ptr3 + (3 + (4*x0)), xmask, eviction_policy='evict_last')
    tmp52 = tl.load(in_ptr3 + (12 + x0), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.sigmoid(tmp2)
    tmp7 = tmp5 + tmp6
    tmp8 = 0.5
    tmp9 = tmp7 * tmp8
    tmp10 = tmp4 * tmp9
    tmp11 = tmp3 * tmp10
    tmp12 = 1.0
    tmp13 = tmp12 - tmp3
    tmp14 = tmp13 * tmp10
    tmp15 = tmp11 - tmp14
    tmp18 = tmp16 + tmp17
    tmp19 = tl.sigmoid(tmp18)
    tmp23 = tmp21 + tmp22
    tmp24 = tmp23 * tmp8
    tmp25 = tmp20 * tmp24
    tmp26 = tmp19 * tmp25
    tmp27 = tmp12 - tmp19
    tmp28 = tmp27 * tmp25
    tmp29 = tmp26 - tmp28
    tmp30 = tmp15 + tmp29
    tmp33 = tmp31 + tmp32
    tmp34 = tl.sigmoid(tmp33)
    tmp38 = tmp36 + tmp37
    tmp39 = tmp38 * tmp8
    tmp40 = tmp35 * tmp39
    tmp41 = tmp34 * tmp40
    tmp42 = tmp12 - tmp34
    tmp43 = tmp42 * tmp40
    tmp44 = tmp41 - tmp43
    tmp45 = tmp30 + tmp44
    tmp48 = tmp46 + tmp47
    tmp49 = tl.sigmoid(tmp48)
    tmp53 = tmp51 + tmp52
    tmp54 = tmp53 * tmp8
    tmp55 = tmp50 * tmp54
    tmp56 = tmp49 * tmp55
    tmp57 = tmp12 - tmp49
    tmp58 = tmp57 * tmp55
    tmp59 = tmp56 - tmp58
    tmp60 = tmp45 + tmp59
    tl.store(out_ptr0 + (x2), tmp60, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_6/inductor_cache/lu/cluvgpnuehlfe66u3wqcj3n6cyvp6elikx7xp4egoya6c67hvs6p.py
# Topologically Sorted Source Nodes: [logits_9], Original ATen: [aten.add]
# Source node to ATen node mapping:
#   logits_9 => add_10
# Graph fragment:
#   %add_10 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%primals_3, %sum_11), kwargs = {})
triton_poi_fused_add_4 = async_compile.triton('triton_poi_fused_add_4', '''
import triton
import triton.language as tl
from triton.compiler.compiler import AttrsDescriptor

from torch._inductor.runtime import triton_helpers, triton_heuristics
from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math
from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, instance_descriptor, DeviceProperties

@triton_heuristics.pointwise(
    size_hints=[64], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=80, major=8, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=108), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2, 3), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_4', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 2, 'num_reduction': 0, 'backend_hash': 'A9C866B4A14FD3277824029365D703C2427B2E685E54EC9B3EF4ADC8D1EEAC1D', 'are_deterministic_algorithms_enabled': False, 'assert_indirect_indexing': True, 'autotune_local_cache': True, 'autotune_pointwise': True, 'autotune_remote_cache': None, 'force_disable_caches': False, 'dynamic_scale_rblock': True, 'max_autotune': False, 'max_autotune_pointwise': False, 'min_split_scan_rblock': 256, 'spill_threshold': 16, 'store_cubin': False},
    min_elem_per_thread=0
)
@triton.jit
def triton_poi_fused_add_4(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x1 = (xindex // 4)
    tmp0 = tl.load(in_ptr0 + (x2), xmask)
    tmp1 = tl.load(in_ptr1 + (x1), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tl.store(out_ptr0 + (x2), tmp2, xmask)
''', device_str='cuda')


async_compile.wait(globals())
del async_compile

def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(primals_2, (1, 4, 4), (16, 4, 1))
    assert_size_stride(primals_3, (4, 4, 4), (16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        # Topologically Sorted Source Nodes: [feats_norm], Original ATen: [aten.linalg_vector_norm]
        stream0 = get_raw_stream(0)
        triton_poi_fused_linalg_vector_norm_0.run(primals_1, buf0, 16, grid=grid(16), stream=stream0)
        buf1 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        # Topologically Sorted Source Nodes: [pairwise_norm], Original ATen: [aten.bmm]
        extern_kernels.bmm(buf0, reinterpret_tensor(buf0, (4, 1, 4), (4, 16, 1), 0), out=buf1)
        buf2 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        # Topologically Sorted Source Nodes: [pairwise_dot], Original ATen: [aten.bmm]
        extern_kernels.bmm(primals_1, reinterpret_tensor(primals_1, (4, 4, 4), (16, 1, 4), 0), out=buf2)
        del primals_1
        buf3 = buf1; del buf1  # reuse
        # Topologically Sorted Source Nodes: [pairwise_sim], Original ATen: [aten.div]
        triton_poi_fused_div_1.run(buf3, buf2, 64, grid=grid(64), stream=stream0)
        buf4 = buf0; del buf0  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_1, sub, mul_2, sub_1, pairwise_potential_E], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_2.run(primals_3, buf3, primals_2, buf4, 16, grid=grid(16), stream=stream0)
        buf5 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_3, sub_2, mul_4, sub_3, pairwise_potential_E_1], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf4, buf3, primals_2, buf5, 16, grid=grid(16), stream=stream0)
        buf6 = buf4; del buf4  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_5, sub_4, mul_6, sub_5, pairwise_potential_E_2], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf5, buf3, primals_2, buf6, 16, grid=grid(16), stream=stream0)
        buf7 = buf5; del buf5  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_7, sub_6, mul_8, sub_7, pairwise_potential_E_3], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf6, buf3, primals_2, buf7, 16, grid=grid(16), stream=stream0)
        buf8 = buf6; del buf6  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_9, sub_8, mul_10, sub_9, pairwise_potential_E_4], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf7, buf3, primals_2, buf8, 16, grid=grid(16), stream=stream0)
        buf9 = buf7; del buf7  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_11, sub_10, mul_12, sub_11, pairwise_potential_E_5], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf8, buf3, primals_2, buf9, 16, grid=grid(16), stream=stream0)
        buf10 = buf8; del buf8  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_13, sub_12, mul_14, sub_13, pairwise_potential_E_6], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf9, buf3, primals_2, buf10, 16, grid=grid(16), stream=stream0)
        buf11 = buf9; del buf9  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_15, sub_14, mul_16, sub_15, pairwise_potential_E_7], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf10, buf3, primals_2, buf11, 16, grid=grid(16), stream=stream0)
        buf12 = buf10; del buf10  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_17, sub_16, mul_18, sub_17, pairwise_potential_E_8], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf11, buf3, primals_2, buf12, 16, grid=grid(16), stream=stream0)
        buf13 = buf11; del buf11  # reuse
        # Topologically Sorted Source Nodes: [add, W_sym, pairwise_potential, mul_19, sub_18, mul_20, sub_19, pairwise_potential_E_9], Original ATen: [aten.add, aten.div, aten.mul, aten.rsub, aten.sub, aten.sum]
        triton_poi_fused_add_div_mul_rsub_sub_sum_3.run(primals_3, buf12, buf3, primals_2, buf13, 16, grid=grid(16), stream=stream0)
        del buf12
        buf14 = buf2; del buf2  # reuse
        # Topologically Sorted Source Nodes: [logits_9], Original ATen: [aten.add]
        triton_poi_fused_add_4.run(primals_3, buf13, buf14, 64, grid=grid(64), stream=stream0)
        del buf13
    return (buf14, primals_2, primals_3, buf3, )


def benchmark_compiled_module(times=10, repeat=10):
    from torch._dynamo.testing import rand_strided
    from torch._inductor.utils import print_performance
    primals_1 = rand_strided((4, 4, 4), (16, 4, 1), device='cuda:0', dtype=torch.float32)
    primals_2 = rand_strided((1, 4, 4), (16, 4, 1), device='cuda:0', dtype=torch.float32)
    primals_3 = rand_strided((4, 4, 4), (16, 4, 1), device='cuda:0', dtype=torch.float32)
    fn = lambda: call([primals_1, primals_2, primals_3])
    return print_performance(fn, times=times, repeat=repeat)


if __name__ == "__main__":
    from torch._inductor.wrapper_benchmark import compiled_module_main
    compiled_module_main('None', benchmark_compiled_module)

import torch
from torch import nn


class CRF(nn.Module):

    def __init__(self, num_nodes, iteration=10):
        """Initialize the CRF module

        Args:
            num_nodes: int, number of nodes/patches within the fully CRF
            iteration: int, number of mean field iterations, e.g. 10
        """
        super(CRF, self).__init__()
        self.num_nodes = num_nodes
        self.iteration = iteration
        self.W = nn.Parameter(torch.zeros(1, num_nodes, num_nodes))

    def forward(self, feats, logits):
        """Performing the CRF. Algorithm details is explained below:

        Within the paper, I formulate the CRF distribution using negative
        energy and cost, e.g. cosine distance, to derive pairwise potentials
        following the convention in energy based models. But for implementation
        simplicity, I use reward, e.g. cosine similarity to derive pairwise
        potentials. So now, pairwise potentials would encourage high reward for
        assigning (y_i, y_j) with the same label if (x_i, x_j) are similar, as
        measured by cosine similarity, pairwise_sim. For

        pairwise_potential_E = torch.sum(
            probs * pairwise_potential - (1 - probs) * pairwise_potential,
            dim=2, keepdim=True
        )

        This is taking the expectation of pairwise potentials using the current
        marginal distribution of each patch being tumor, i.e. probs. There are
        four cases to consider when taking the expectation between (i, j):

        1. i=T,j=T; 2. i=N,j=T; 3. i=T,j=N; 4. i=N,j=N

        probs is the marginal distribution of each i being tumor, therefore
        logits > 0 means tumor and logits < 0 means normal. Given this, the
        full expectation equation should be:

        [probs * +pairwise_potential] + [(1 - probs) * +pairwise_potential] +
                    case 1                            case 2
        [probs * -pairwise_potential] + [(1 - probs) * -pairwise_potential]
                    case 3                            case 4

        positive sign rewards logits to be more tumor and negative sign rewards
        logits to be more normal. But because of label compatibility, i.e. the
        indicator function within equation 3 in the paper, case 2 and case 3
        are dropped, which ends up being:

        probs * pairwise_potential - (1 - probs) * pairwise_potential

        In high level speaking, if (i, j) embedding are different, then
        pairwise_potential, as computed as cosine similarity, would approach 0,
        which then as no affect anyway. if (i, j) embedding are similar, then
        pairwise_potential would be a positive reward. In this case,
        if probs -> 1, then pairwise_potential promotes tumor probability;
        if probs -> 0, then -pairwise_potential promotes normal probability.

        Args:
            feats: 3D tensor with the shape of
            [batch_size, num_nodes, embedding_size], where num_nodes is the
            number of patches within a grid, e.g. 9 for a 3x3 grid;
            embedding_size is the size of extracted feature representation for
            each patch from ResNet, e.g. 512
            logits: 3D tensor with shape of [batch_size, num_nodes, 1], the
            logit of each patch within the grid being tumor before CRF

        Returns:
            logits: 3D tensor with shape of [batch_size, num_nodes, 1], the
            logit of each patch within the grid being tumor after CRF
        """
        feats_norm = torch.norm(feats, p=2, dim=2, keepdim=True)
        pairwise_norm = torch.bmm(feats_norm, torch.transpose(feats_norm, 1, 2)
            )
        pairwise_dot = torch.bmm(feats, torch.transpose(feats, 1, 2))
        pairwise_sim = pairwise_dot / pairwise_norm
        W_sym = (self.W + torch.transpose(self.W, 1, 2)) / 2
        pairwise_potential = pairwise_sim * W_sym
        unary_potential = logits.clone()
        for i in range(self.iteration):
            probs = torch.transpose(logits.sigmoid(), 1, 2)
            pairwise_potential_E = torch.sum(probs * pairwise_potential - (
                1 - probs) * pairwise_potential, dim=2, keepdim=True)
            logits = unary_potential + pairwise_potential_E
        return logits

    def __repr__(self):
        return 'CRF(num_nodes={}, iteration={})'.format(self.num_nodes,
            self.iteration)


def get_inputs():
    return [torch.rand([4, 4, 4]), torch.rand([4, 4, 4])]


def get_init_inputs():
    return [[], {'num_nodes': 4}]

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.triton_helpers import libdevice
from torch import nn
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_linalg_vector_norm_0(in_ptr0, out_ptr0, xnumel, XBLOCK:
    tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + 4 * x0, xmask, eviction_policy='evict_last')
    tmp2 = tl.load(in_ptr0 + (1 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr0 + (2 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp8 = tl.load(in_ptr0 + (3 + 4 * x0), xmask, eviction_policy='evict_last')
    tmp1 = tmp0 * tmp0
    tmp3 = tmp2 * tmp2
    tmp4 = tmp1 + tmp3
    tmp6 = tmp5 * tmp5
    tmp7 = tmp4 + tmp6
    tmp9 = tmp8 * tmp8
    tmp10 = tmp7 + tmp9
    tmp11 = libdevice.sqrt(tmp10)
    tl.store(out_ptr0 + x0, tmp11, xmask)


@triton.jit
def triton_poi_fused_div_1(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + x0, xmask)
    tmp1 = tl.load(in_out_ptr0 + x0, xmask)
    tmp2 = tmp0 / tmp1
    tl.store(in_out_ptr0 + x0, tmp2, xmask)


@triton.jit
def triton_poi_fused_add_div_mul_rsub_sub_sum_2(in_ptr0, in_ptr1, in_ptr2,
    out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = xindex // 4
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 16 * x1), xmask)
    tmp2 = tl.load(in_ptr1 + 4 * x2, xmask, eviction_policy='evict_last')
    tmp3 = tl.load(in_ptr2 + 4 * x0, xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr2 + x0, xmask, eviction_policy='evict_last')
    tmp14 = tl.load(in_ptr0 + (4 + x0 + 16 * x1), xmask)
    tmp16 = tl.load(in_ptr1 + (1 + 4 * x2), xmask, eviction_policy='evict_last'
        )
    tmp17 = tl.load(in_ptr2 + (1 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp18 = tl.load(in_ptr2 + (4 + x0), xmask, eviction_policy='evict_last')
    tmp27 = tl.load(in_ptr0 + (8 + x0 + 16 * x1), xmask)
    tmp29 = tl.load(in_ptr1 + (2 + 4 * x2), xmask, eviction_policy='evict_last'
        )
    tmp30 = tl.load(in_ptr2 + (2 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp31 = tl.load(in_ptr2 + (8 + x0), xmask, eviction_policy='evict_last')
    tmp40 = tl.load(in_ptr0 + (12 + x0 + 16 * x1), xmask)
    tmp42 = tl.load(in_ptr1 + (3 + 4 * x2), xmask, eviction_policy='evict_last'
        )
    tmp43 = tl.load(in_ptr2 + (3 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp44 = tl.load(in_ptr2 + (12 + x0), xmask, eviction_policy='evict_last')
    tmp1 = tl.sigmoid(tmp0)
    tmp5 = tmp3 + tmp4
    tmp6 = 0.5
    tmp7 = tmp5 * tmp6
    tmp8 = tmp2 * tmp7
    tmp9 = tmp1 * tmp8
    tmp10 = 1.0
    tmp11 = tmp10 - tmp1
    tmp12 = tmp11 * tmp8
    tmp13 = tmp9 - tmp12
    tmp15 = tl.sigmoid(tmp14)
    tmp19 = tmp17 + tmp18
    tmp20 = tmp19 * tmp6
    tmp21 = tmp16 * tmp20
    tmp22 = tmp15 * tmp21
    tmp23 = tmp10 - tmp15
    tmp24 = tmp23 * tmp21
    tmp25 = tmp22 - tmp24
    tmp26 = tmp13 + tmp25
    tmp28 = tl.sigmoid(tmp27)
    tmp32 = tmp30 + tmp31
    tmp33 = tmp32 * tmp6
    tmp34 = tmp29 * tmp33
    tmp35 = tmp28 * tmp34
    tmp36 = tmp10 - tmp28
    tmp37 = tmp36 * tmp34
    tmp38 = tmp35 - tmp37
    tmp39 = tmp26 + tmp38
    tmp41 = tl.sigmoid(tmp40)
    tmp45 = tmp43 + tmp44
    tmp46 = tmp45 * tmp6
    tmp47 = tmp42 * tmp46
    tmp48 = tmp41 * tmp47
    tmp49 = tmp10 - tmp41
    tmp50 = tmp49 * tmp47
    tmp51 = tmp48 - tmp50
    tmp52 = tmp39 + tmp51
    tl.store(out_ptr0 + x2, tmp52, xmask)


@triton.jit
def triton_poi_fused_add_div_mul_rsub_sub_sum_3(in_ptr0, in_ptr1, in_ptr2,
    in_ptr3, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = xindex // 4
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + (x0 + 16 * x1), xmask)
    tmp1 = tl.load(in_ptr1 + 4 * x1, xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr2 + 4 * x2, xmask, eviction_policy='evict_last')
    tmp5 = tl.load(in_ptr3 + 4 * x0, xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr3 + x0, xmask, eviction_policy='evict_last')
    tmp16 = tl.load(in_ptr0 + (4 + x0 + 16 * x1), xmask)
    tmp17 = tl.load(in_ptr1 + (1 + 4 * x1), xmask, eviction_policy='evict_last'
        )
    tmp20 = tl.load(in_ptr2 + (1 + 4 * x2), xmask, eviction_policy='evict_last'
        )
    tmp21 = tl.load(in_ptr3 + (1 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp22 = tl.load(in_ptr3 + (4 + x0), xmask, eviction_policy='evict_last')
    tmp31 = tl.load(in_ptr0 + (8 + x0 + 16 * x1), xmask)
    tmp32 = tl.load(in_ptr1 + (2 + 4 * x1), xmask, eviction_policy='evict_last'
        )
    tmp35 = tl.load(in_ptr2 + (2 + 4 * x2), xmask, eviction_policy='evict_last'
        )
    tmp36 = tl.load(in_ptr3 + (2 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp37 = tl.load(in_ptr3 + (8 + x0), xmask, eviction_policy='evict_last')
    tmp46 = tl.load(in_ptr0 + (12 + x0 + 16 * x1), xmask)
    tmp47 = tl.load(in_ptr1 + (3 + 4 * x1), xmask, eviction_policy='evict_last'
        )
    tmp50 = tl.load(in_ptr2 + (3 + 4 * x2), xmask, eviction_policy='evict_last'
        )
    tmp51 = tl.load(in_ptr3 + (3 + 4 * x0), xmask, eviction_policy='evict_last'
        )
    tmp52 = tl.load(in_ptr3 + (12 + x0), xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.sigmoid(tmp2)
    tmp7 = tmp5 + tmp6
    tmp8 = 0.5
    tmp9 = tmp7 * tmp8
    tmp10 = tmp4 * tmp9
    tmp11 = tmp3 * tmp10
    tmp12 = 1.0
    tmp13 = tmp12 - tmp3
    tmp14 = tmp13 * tmp10
    tmp15 = tmp11 - tmp14
    tmp18 = tmp16 + tmp17
    tmp19 = tl.sigmoid(tmp18)
    tmp23 = tmp21 + tmp22
    tmp24 = tmp23 * tmp8
    tmp25 = tmp20 * tmp24
    tmp26 = tmp19 * tmp25
    tmp27 = tmp12 - tmp19
    tmp28 = tmp27 * tmp25
    tmp29 = tmp26 - tmp28
    tmp30 = tmp15 + tmp29
    tmp33 = tmp31 + tmp32
    tmp34 = tl.sigmoid(tmp33)
    tmp38 = tmp36 + tmp37
    tmp39 = tmp38 * tmp8
    tmp40 = tmp35 * tmp39
    tmp41 = tmp34 * tmp40
    tmp42 = tmp12 - tmp34
    tmp43 = tmp42 * tmp40
    tmp44 = tmp41 - tmp43
    tmp45 = tmp30 + tmp44
    tmp48 = tmp46 + tmp47
    tmp49 = tl.sigmoid(tmp48)
    tmp53 = tmp51 + tmp52
    tmp54 = tmp53 * tmp8
    tmp55 = tmp50 * tmp54
    tmp56 = tmp49 * tmp55
    tmp57 = tmp12 - tmp49
    tmp58 = tmp57 * tmp55
    tmp59 = tmp56 - tmp58
    tmp60 = tmp45 + tmp59
    tl.store(out_ptr0 + x2, tmp60, xmask)


@triton.jit
def triton_poi_fused_add_4(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK: tl.
    constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x1 = xindex // 4
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tl.store(out_ptr0 + x2, tmp2, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(primals_2, (1, 4, 4), (16, 4, 1))
    assert_size_stride(primals_3, (4, 4, 4), (16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_linalg_vector_norm_0[grid(16)](primals_1, buf0, 16,
            XBLOCK=16, num_warps=1, num_stages=1)
        buf1 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        extern_kernels.bmm(buf0, reinterpret_tensor(buf0, (4, 1, 4), (4, 16,
            1), 0), out=buf1)
        buf2 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.float32)
        extern_kernels.bmm(primals_1, reinterpret_tensor(primals_1, (4, 4, 
            4), (16, 1, 4), 0), out=buf2)
        del primals_1
        buf3 = buf1
        del buf1
        triton_poi_fused_div_1[grid(64)](buf3, buf2, 64, XBLOCK=64,
            num_warps=1, num_stages=1)
        buf4 = buf0
        del buf0
        triton_poi_fused_add_div_mul_rsub_sub_sum_2[grid(16)](primals_3,
            buf3, primals_2, buf4, 16, XBLOCK=16, num_warps=1, num_stages=1)
        buf5 = empty_strided_cuda((4, 4, 1), (4, 1, 16), torch.float32)
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf4, buf3, primals_2, buf5, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf6 = buf4
        del buf4
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf5, buf3, primals_2, buf6, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf7 = buf5
        del buf5
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf6, buf3, primals_2, buf7, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf8 = buf6
        del buf6
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf7, buf3, primals_2, buf8, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf9 = buf7
        del buf7
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf8, buf3, primals_2, buf9, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf10 = buf8
        del buf8
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf9, buf3, primals_2, buf10, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf11 = buf9
        del buf9
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf10, buf3, primals_2, buf11, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf12 = buf10
        del buf10
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf11, buf3, primals_2, buf12, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        buf13 = buf11
        del buf11
        triton_poi_fused_add_div_mul_rsub_sub_sum_3[grid(16)](primals_3,
            buf12, buf3, primals_2, buf13, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
        del buf12
        buf14 = buf2
        del buf2
        triton_poi_fused_add_4[grid(64)](primals_3, buf13, buf14, 64,
            XBLOCK=64, num_warps=1, num_stages=1)
        del buf13
    return buf14, primals_2, primals_3, buf3


class CRFNew(nn.Module):

    def __init__(self, num_nodes, iteration=10):
        """Initialize the CRF module

        Args:
            num_nodes: int, number of nodes/patches within the fully CRF
            iteration: int, number of mean field iterations, e.g. 10
        """
        super(CRFNew, self).__init__()
        self.num_nodes = num_nodes
        self.iteration = iteration
        self.W = nn.Parameter(torch.zeros(1, num_nodes, num_nodes))

    def __repr__(self):
        return 'CRF(num_nodes={}, iteration={})'.format(self.num_nodes,
            self.iteration)

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