# 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_8/inductor_cache/tu/ctuej2j6f3oxr5p43q7juhagc3r3ncgs2ikvxemutunlnxlnvl24.py
# Topologically Sorted Source Nodes: [adaptive_avg_pool3d], Original ATen: [aten.mean]
# Source node to ATen node mapping:
#   adaptive_avg_pool3d => mean
# Graph fragment:
#   %mean : [num_users=1] = call_function[target=torch.ops.aten.mean.dim](args = (%arg0_1, [-1, -2, -3], True), kwargs = {})
triton_per_fused_mean_0 = async_compile.triton('triton_per_fused_mean_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.persistent_reduction(
    size_hints=[4, 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=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, 3), equal_to_1=())]},
    inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_per_fused_mean_0', 'mutated_arg_names': ['in_out_ptr0'], 'no_x_dim': False, 'num_load': 1, 'num_reduction': 1, '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}
)
@triton.jit
def triton_per_fused_mean_0(in_out_ptr0, in_ptr0, xnumel, rnumel, XBLOCK : tl.constexpr):
    xnumel = 4
    rnumel = 64
    RBLOCK: tl.constexpr = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    roffset = 0
    rmask = tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + (64*x0)), xmask, other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.where(xmask, tmp1, 0)
    tmp4 = tl.sum(tmp3, 1)[:, None]
    tmp5 = 64.0
    tmp6 = tmp4 / tmp5
    tl.debug_barrier()
    tl.store(in_out_ptr0 + (x0), tmp6, xmask)
''', 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, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf1 = reinterpret_tensor(buf0, (4, 1, 1, 1), (1, 1, 1, 1), 0); del buf0  # reuse
        # Topologically Sorted Source Nodes: [adaptive_avg_pool3d], Original ATen: [aten.mean]
        stream0 = get_raw_stream(0)
        triton_per_fused_mean_0.run(buf1, arg0_1, 4, 64, grid=grid(4), stream=stream0)
        del arg0_1
    return (buf1, )


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
from abc import abstractmethod
from typing import Tuple
import torch.utils.data
import torch.nn


class EfficientBlockBase(nn.Module):
    """
    PyTorchVideo/accelerator provides a set of efficient blocks
    that have optimal efficiency for each target hardware device.

    Each efficient block has two forms:
    - original form: this form is for training. When efficient block is instantiated,
        it is in this original form.
    - deployable form: this form is for deployment. Once the network is ready for
        deploy, it can be converted into deployable form for efficient execution
        on target hardware. One block is transformed into deployable form by calling
        convert() method. By conversion to deployable form,
        various optimization (operator fuse, kernel optimization, etc.) are applied.

    EfficientBlockBase is the base class for efficient blocks.
    All efficient blocks should inherit this base class
    and implement following methods:
    - forward(): same as required by nn.Module
    - convert(): called to convert block into deployable form
    """

    @abstractmethod
    def convert(self):
        pass

    @abstractmethod
    def forward(self):
        pass


class AdaptiveAvgPool3dOutSize1(EfficientBlockBase):
    """
    Implements AdaptiveAvgPool3d with output (T, H, W) = (1, 1, 1). This operator has
    better efficiency than AdaptiveAvgPool for mobile CPU.
    """

    def __init__(self):
        super().__init__()
        self.pool = nn.AdaptiveAvgPool3d(1)
        self.convert_flag = False

    def convert(self, input_blob_size: 'Tuple', **kwargs):
        """
        Converts AdaptiveAvgPool into AvgPool with constant kernel size for better
        efficiency.
        Args:
            input_blob_size (tuple): blob size at the input of
                AdaptiveAvgPool3dOutSize1 instance during forward.
            kwargs (any): any keyword argument (unused).
        """
        assert self.convert_flag is False, 'AdaptiveAvgPool3dOutSize1: already converted, cannot be converted again'
        kernel_size = input_blob_size[2:]
        self.pool = nn.AvgPool3d(kernel_size)
        self.convert_flag = True

    def forward(self, x):
        return self.pool(x)


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
import torch.nn as nn
from abc import abstractmethod
from typing import Tuple
import torch.utils.data
import torch.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_per_fused_mean_0(in_out_ptr0, in_ptr0, xnumel, rnumel, XBLOCK:
    tl.constexpr):
    xnumel = 4
    RBLOCK: tl.constexpr = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + 64 * x0), xmask, other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.where(xmask, tmp1, 0)
    tmp4 = tl.sum(tmp3, 1)[:, None]
    tmp5 = 64.0
    tmp6 = tmp4 / tmp5
    tl.debug_barrier()
    tl.store(in_out_ptr0 + x0, tmp6, 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)
        buf0 = empty_strided_cuda((4, 1, 1, 1), (1, 4, 4, 4), torch.float32)
        buf1 = reinterpret_tensor(buf0, (4, 1, 1, 1), (1, 1, 1, 1), 0)
        del buf0
        get_raw_stream(0)
        triton_per_fused_mean_0[grid(4)](buf1, arg0_1, 4, 64, XBLOCK=1,
            num_warps=2, num_stages=1)
        del arg0_1
    return buf1,


class EfficientBlockBase(nn.Module):
    """
    PyTorchVideo/accelerator provides a set of efficient blocks
    that have optimal efficiency for each target hardware device.

    Each efficient block has two forms:
    - original form: this form is for training. When efficient block is instantiated,
        it is in this original form.
    - deployable form: this form is for deployment. Once the network is ready for
        deploy, it can be converted into deployable form for efficient execution
        on target hardware. One block is transformed into deployable form by calling
        convert() method. By conversion to deployable form,
        various optimization (operator fuse, kernel optimization, etc.) are applied.

    EfficientBlockBase is the base class for efficient blocks.
    All efficient blocks should inherit this base class
    and implement following methods:
    - forward(): same as required by nn.Module
    - convert(): called to convert block into deployable form
    """

    @abstractmethod
    def convert(self):
        pass

    @abstractmethod
    def forward(self):
        pass


class AdaptiveAvgPool3dOutSize1New(EfficientBlockBase):
    """
    Implements AdaptiveAvgPool3d with output (T, H, W) = (1, 1, 1). This operator has
    better efficiency than AdaptiveAvgPool for mobile CPU.
    """

    def __init__(self):
        super().__init__()
        self.pool = nn.AdaptiveAvgPool3d(1)
        self.convert_flag = False

    def convert(self, input_blob_size: 'Tuple', **kwargs):
        """
        Converts AdaptiveAvgPool into AvgPool with constant kernel size for better
        efficiency.
        Args:
            input_blob_size (tuple): blob size at the input of
                AdaptiveAvgPool3dOutSize1 instance during forward.
            kwargs (any): any keyword argument (unused).
        """
        assert self.convert_flag is False, 'AdaptiveAvgPool3dOutSize1: already converted, cannot be converted again'
        kernel_size = input_blob_size[2:]
        self.pool = nn.AvgPool3d(kernel_size)
        self.convert_flag = True

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