# AOT ID: ['0_inference']
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_1/inductor_cache/wm/cwmtdji7wew23pxblgrnor37cmtzgq5ehdup65sglxkhagz5loak.py
# Topologically Sorted Source Nodes: [q], Original ATen: [aten.cat]
# Source node to ATen node mapping:
#   q => cat
# Graph fragment:
#   %cat : [num_users=2] = call_function[target=torch.ops.aten.cat.default](args = ([%div, %div_1], 1), kwargs = {})
triton_poi_fused_cat_0 = async_compile.triton('triton_poi_fused_cat_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=[512], 
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 2), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_cat_0', 'mutated_arg_names': [], 'no_x_dim': False, 'num_load': 9, 'num_reduction': 0, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', '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_cat_0(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 320
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = (xindex // 16) % 5
    x0 = xindex % 16
    x2 = (xindex // 80)
    x3 = xindex
    tmp0 = x1
    tmp1 = tl.full([1], 0, tl.int64)
    tmp2 = tmp0 >= tmp1
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (x0 + (16*x1) + (64*x2)), tmp4 & xmask, other=0.0)
    tmp6 = 2.0
    tmp7 = tmp5 * tmp6
    tmp8 = tl.load(in_ptr0 + (x0 + (64*x2)), tmp4 & xmask, eviction_policy='evict_last', other=0.0)
    tmp9 = tmp8 * tmp8
    tmp10 = tl.load(in_ptr0 + (16 + x0 + (64*x2)), tmp4 & xmask, eviction_policy='evict_last', other=0.0)
    tmp11 = tmp10 * tmp10
    tmp12 = tmp9 + tmp11
    tmp13 = tl.load(in_ptr0 + (32 + x0 + (64*x2)), tmp4 & xmask, eviction_policy='evict_last', other=0.0)
    tmp14 = tmp13 * tmp13
    tmp15 = tmp12 + tmp14
    tmp16 = tl.load(in_ptr0 + (48 + x0 + (64*x2)), tmp4 & xmask, eviction_policy='evict_last', other=0.0)
    tmp17 = tmp16 * tmp16
    tmp18 = tmp15 + tmp17
    tmp19 = libdevice.sqrt(tmp18)
    tmp20 = tmp19 * tmp19
    tmp21 = 1.0
    tmp22 = tmp20 + tmp21
    tmp23 = tmp7 / tmp22
    tmp24 = tl.full(tmp23.shape, 0.0, tmp23.dtype)
    tmp25 = tl.where(tmp4, tmp23, tmp24)
    tmp26 = tmp0 >= tmp3
    tmp27 = tl.full([1], 5, tl.int64)
    tmp28 = tmp0 < tmp27
    tmp29 = tl.load(in_ptr0 + (x0 + (64*x2)), tmp26 & xmask, eviction_policy='evict_last', other=0.0)
    tmp30 = tmp29 * tmp29
    tmp31 = tl.load(in_ptr0 + (16 + x0 + (64*x2)), tmp26 & xmask, eviction_policy='evict_last', other=0.0)
    tmp32 = tmp31 * tmp31
    tmp33 = tmp30 + tmp32
    tmp34 = tl.load(in_ptr0 + (32 + x0 + (64*x2)), tmp26 & xmask, eviction_policy='evict_last', other=0.0)
    tmp35 = tmp34 * tmp34
    tmp36 = tmp33 + tmp35
    tmp37 = tl.load(in_ptr0 + (48 + x0 + (64*x2)), tmp26 & xmask, eviction_policy='evict_last', other=0.0)
    tmp38 = tmp37 * tmp37
    tmp39 = tmp36 + tmp38
    tmp40 = libdevice.sqrt(tmp39)
    tmp41 = tmp40 * tmp40
    tmp42 = tmp41 - tmp21
    tmp43 = tmp41 + tmp21
    tmp44 = tmp42 / tmp43
    tmp45 = tl.full(tmp44.shape, 0.0, tmp44.dtype)
    tmp46 = tl.where(tmp26, tmp44, tmp45)
    tmp47 = tl.where(tmp4, tmp25, tmp46)
    tl.store(out_ptr0 + (x3), tmp47, xmask)
''', device_str='cuda')


# kernel path: runs/run_shard_1/inductor_cache/4s/c4sf274gpf37u37symxy5dtiezp6lkdsditc3vzvonvuszsmn3yr.py
# Topologically Sorted Source Nodes: [mul_1, norm_1, q_norm, add_2, loss, acos, mean], Original ATen: [aten.mul, aten.linalg_vector_norm, aten.pow, aten.add, aten.div, aten.acos, aten.mean]
# Source node to ATen node mapping:
#   acos => acos
#   add_2 => add_2
#   loss => div_2
#   mean => mean
#   mul_1 => mul_1
#   norm_1 => pow_4, pow_5, sum_2
#   q_norm => pow_6
# Graph fragment:
#   %mul_1 : [num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%select, 2), kwargs = {})
#   %pow_4 : [num_users=1] = call_function[target=torch.ops.aten.pow.Tensor_Scalar](args = (%cat, 2), kwargs = {})
#   %sum_2 : [num_users=1] = call_function[target=torch.ops.aten.sum.dim_IntList](args = (%pow_4, [1]), kwargs = {})
#   %pow_5 : [num_users=1] = call_function[target=torch.ops.aten.pow.Tensor_Scalar](args = (%sum_2, 0.5), kwargs = {})
#   %pow_6 : [num_users=1] = call_function[target=torch.ops.aten.pow.Tensor_Scalar](args = (%pow_5, 2), kwargs = {})
#   %add_2 : [num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%pow_6, 1), kwargs = {})
#   %div_2 : [num_users=1] = call_function[target=torch.ops.aten.div.Tensor](args = (%mul_1, %add_2), kwargs = {})
#   %acos : [num_users=1] = call_function[target=torch.ops.aten.acos.default](args = (%div_2,), kwargs = {})
#   %mean : [num_users=1] = call_function[target=torch.ops.aten.mean.default](args = (%acos,), kwargs = {})
triton_per_fused_acos_add_div_linalg_vector_norm_mean_mul_pow_1 = async_compile.triton('triton_per_fused_acos_add_div_linalg_vector_norm_mean_mul_pow_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.persistent_reduction(
    size_hints=[1, 64],
    reduction_hint=ReductionHint.INNER,
    filename=__file__,
    triton_meta={'signature': {0: '*fp32', 1: '*fp32', 2: 'i32', 3: 'i32'}, 'device': DeviceProperties(type='cuda', index=0, cc=90, major=9, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, multi_processor_count=132), 'constants': {2: 1}, 'configs': [AttrsDescriptor(divisible_by_16=(0, 1, 3), equal_to_1=(2,))]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_acos_add_div_linalg_vector_norm_mean_mul_pow_1', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 5, 'num_reduction': 1, 'backend_hash': 'B098E03CDA7B8ADC90DAFFDF24A2956451D1B13F297756A5DCC209498AA53705', '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}
)
@triton.jit
def triton_per_fused_acos_add_div_linalg_vector_norm_mean_mul_pow_1(in_out_ptr0, in_ptr0, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 1
    rnumel = 64
    RBLOCK: tl.constexpr = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r0 = rindex % 16
    r1 = (rindex // 16)
    tmp0 = tl.load(in_ptr0 + (64 + r0 + (80*r1)), None)
    tmp3 = tl.load(in_ptr0 + (r0 + (80*r1)), None)
    tmp5 = tl.load(in_ptr0 + (16 + r0 + (80*r1)), None)
    tmp8 = tl.load(in_ptr0 + (32 + r0 + (80*r1)), None)
    tmp11 = tl.load(in_ptr0 + (48 + r0 + (80*r1)), None)
    tmp1 = 2.0
    tmp2 = tmp0 * tmp1
    tmp4 = tmp3 * tmp3
    tmp6 = tmp5 * tmp5
    tmp7 = tmp4 + tmp6
    tmp9 = tmp8 * tmp8
    tmp10 = tmp7 + tmp9
    tmp12 = tmp11 * tmp11
    tmp13 = tmp10 + tmp12
    tmp14 = tmp0 * tmp0
    tmp15 = tmp13 + tmp14
    tmp16 = libdevice.sqrt(tmp15)
    tmp17 = tmp16 * tmp16
    tmp18 = 1.0
    tmp19 = tmp17 + tmp18
    tmp20 = tmp2 / tmp19
    tmp21 = libdevice.acos(tmp20)
    tmp22 = tl.broadcast_to(tmp21, [XBLOCK, RBLOCK])
    tmp24 = tl.sum(tmp22, 1)[:, None]
    tmp25 = 64.0
    tmp26 = tmp24 / tmp25
    tl.debug_barrier()
    tl.store(in_out_ptr0 + (tl.full([XBLOCK, 1], 0, tl.int32)), tmp26, None)
''', device_str='cuda')


async_compile.wait(globals())
del async_compile

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)
        buf0 = empty_strided_cuda((4, 5, 4, 4), (80, 16, 4, 1), torch.float32)
        # Topologically Sorted Source Nodes: [q], Original ATen: [aten.cat]
        stream0 = get_raw_stream(0)
        triton_poi_fused_cat_0.run(arg0_1, buf0, 320, grid=grid(320), stream=stream0)
        del arg0_1
        buf1 = empty_strided_cuda((), (), torch.float32)
        buf2 = buf1; del buf1  # reuse
        # Topologically Sorted Source Nodes: [mul_1, norm_1, q_norm, add_2, loss, acos, mean], Original ATen: [aten.mul, aten.linalg_vector_norm, aten.pow, aten.add, aten.div, aten.acos, aten.mean]
        triton_per_fused_acos_add_div_linalg_vector_norm_mean_mul_pow_1.run(buf2, buf0, 1, 64, grid=grid(1), stream=stream0)
        del buf0
    return (buf2, )


def benchmark_compiled_module(times=10, repeat=10):
    from torch._dynamo.testing import rand_strided
    from torch._inductor.utils import print_performance
    arg0_1 = rand_strided((4, 4, 4, 4), (64, 16, 4, 1), device='cuda:0', dtype=torch.float32)
    fn = lambda: call([arg0_1])
    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
import torch.nn as nn


class HyperSphereLoss(nn.Module):

    def forward(self, input):
        """
        Calcuate distance between input and N(North Pole) using hypersphere metrics.

        Woo Park, Sung, and Junseok Kwon.
        "Sphere Generative Adversarial Network Based on Geometric Moment Matching."
        Proceedings of the IEEE Conference on Computer Vision and Pattern Recognition. 2019.
        """
        q = self.project_to_hypersphere(input)
        q_norm = torch.norm(q, dim=1) ** 2
        loss = 2 * q[:, -1] / (1 + q_norm)
        return torch.mean(torch.acos(loss))

    def project_to_hypersphere(self, v):
        v_norm = torch.norm(v, dim=1, keepdim=True) ** 2
        a = 2 * v / (v_norm + 1)
        b = (v_norm - 1) / (v_norm + 1)
        return torch.cat([a, b], dim=1)


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


def get_init_inputs():
    return [[], {}]

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
from torch._inductor.runtime.triton_helpers import libdevice
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_cat_0(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 320
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 16 % 5
    x0 = xindex % 16
    x2 = xindex // 80
    x3 = xindex
    tmp0 = x1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (x0 + 16 * x1 + 64 * x2), tmp4 & xmask, other=0.0)
    tmp6 = 2.0
    tmp7 = tmp5 * tmp6
    tmp8 = tl.load(in_ptr0 + (x0 + 64 * x2), tmp4 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp9 = tmp8 * tmp8
    tmp10 = tl.load(in_ptr0 + (16 + x0 + 64 * x2), tmp4 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp11 = tmp10 * tmp10
    tmp12 = tmp9 + tmp11
    tmp13 = tl.load(in_ptr0 + (32 + x0 + 64 * x2), tmp4 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp14 = tmp13 * tmp13
    tmp15 = tmp12 + tmp14
    tmp16 = tl.load(in_ptr0 + (48 + x0 + 64 * x2), tmp4 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp17 = tmp16 * tmp16
    tmp18 = tmp15 + tmp17
    tmp19 = libdevice.sqrt(tmp18)
    tmp20 = tmp19 * tmp19
    tmp21 = 1.0
    tmp22 = tmp20 + tmp21
    tmp23 = tmp7 / tmp22
    tmp24 = tl.full(tmp23.shape, 0.0, tmp23.dtype)
    tmp25 = tl.where(tmp4, tmp23, tmp24)
    tmp26 = tmp0 >= tmp3
    tl.full([1], 5, tl.int64)
    tmp29 = tl.load(in_ptr0 + (x0 + 64 * x2), tmp26 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp30 = tmp29 * tmp29
    tmp31 = tl.load(in_ptr0 + (16 + x0 + 64 * x2), tmp26 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp32 = tmp31 * tmp31
    tmp33 = tmp30 + tmp32
    tmp34 = tl.load(in_ptr0 + (32 + x0 + 64 * x2), tmp26 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp35 = tmp34 * tmp34
    tmp36 = tmp33 + tmp35
    tmp37 = tl.load(in_ptr0 + (48 + x0 + 64 * x2), tmp26 & xmask,
        eviction_policy='evict_last', other=0.0)
    tmp38 = tmp37 * tmp37
    tmp39 = tmp36 + tmp38
    tmp40 = libdevice.sqrt(tmp39)
    tmp41 = tmp40 * tmp40
    tmp42 = tmp41 - tmp21
    tmp43 = tmp41 + tmp21
    tmp44 = tmp42 / tmp43
    tmp45 = tl.full(tmp44.shape, 0.0, tmp44.dtype)
    tmp46 = tl.where(tmp26, tmp44, tmp45)
    tmp47 = tl.where(tmp4, tmp25, tmp46)
    tl.store(out_ptr0 + x3, tmp47, xmask)


@triton.jit
def triton_per_fused_acos_add_div_linalg_vector_norm_mean_mul_pow_1(in_out_ptr0
    , in_ptr0, xnumel, rnumel, XBLOCK: tl.constexpr):
    RBLOCK: tl.constexpr = 64
    xoffset = tl.program_id(0) * XBLOCK
    xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r0 = rindex % 16
    r1 = rindex // 16
    tmp0 = tl.load(in_ptr0 + (64 + r0 + 80 * r1), None)
    tmp3 = tl.load(in_ptr0 + (r0 + 80 * r1), None)
    tmp5 = tl.load(in_ptr0 + (16 + r0 + 80 * r1), None)
    tmp8 = tl.load(in_ptr0 + (32 + r0 + 80 * r1), None)
    tmp11 = tl.load(in_ptr0 + (48 + r0 + 80 * r1), None)
    tmp1 = 2.0
    tmp2 = tmp0 * tmp1
    tmp4 = tmp3 * tmp3
    tmp6 = tmp5 * tmp5
    tmp7 = tmp4 + tmp6
    tmp9 = tmp8 * tmp8
    tmp10 = tmp7 + tmp9
    tmp12 = tmp11 * tmp11
    tmp13 = tmp10 + tmp12
    tmp14 = tmp0 * tmp0
    tmp15 = tmp13 + tmp14
    tmp16 = libdevice.sqrt(tmp15)
    tmp17 = tmp16 * tmp16
    tmp18 = 1.0
    tmp19 = tmp17 + tmp18
    tmp20 = tmp2 / tmp19
    tmp21 = libdevice.acos(tmp20)
    tmp22 = tl.broadcast_to(tmp21, [XBLOCK, RBLOCK])
    tmp24 = tl.sum(tmp22, 1)[:, None]
    tmp25 = 64.0
    tmp26 = tmp24 / tmp25
    tl.debug_barrier()
    tl.store(in_out_ptr0 + tl.full([XBLOCK, 1], 0, tl.int32), tmp26, None)


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)
        buf0 = empty_strided_cuda((4, 5, 4, 4), (80, 16, 4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_cat_0[grid(320)](arg0_1, buf0, 320, XBLOCK=256,
            num_warps=4, num_stages=1)
        del arg0_1
        buf1 = empty_strided_cuda((), (), torch.float32)
        buf2 = buf1
        del buf1
        triton_per_fused_acos_add_div_linalg_vector_norm_mean_mul_pow_1[grid(1)
            ](buf2, buf0, 1, 64, XBLOCK=1, num_warps=2, num_stages=1)
        del buf0
    return buf2,


class HyperSphereLossNew(nn.Module):

    def project_to_hypersphere(self, v):
        v_norm = torch.norm(v, dim=1, keepdim=True) ** 2
        a = 2 * v / (v_norm + 1)
        b = (v_norm - 1) / (v_norm + 1)
        return torch.cat([a, b], dim=1)

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