Upload folder using huggingface_hub (part 3)

.inductor_cache/ym/cymwujume2azuvh7fbujnaekdocxf6szi57kpkkwcmuszidxouzv.py

@@ -0,0 +1,29 @@
+
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 1}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'out_ptr0': '*fp32', 'out_ptr1': '*fp32', 'xnumel': 'constexpr', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=3, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {'xnumel': 1}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_max_mul_unsqueeze_view_0', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 2, 'num_store': 2, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_max_mul_unsqueeze_view_0(in_ptr0, in_ptr1, out_ptr0, out_ptr1, xnumel, XBLOCK : tl.constexpr):
+    xnumel = 1
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = tl.full([XBLOCK], True, tl.int1)[:]
+    tmp0 = tl.load(in_ptr0 + (0))
+    tmp1 = tl.broadcast_to(tmp0, [XBLOCK])
+    tmp2 = tl.load(in_ptr1 + (0))
+    tmp3 = tl.broadcast_to(tmp2, [XBLOCK])
+    tmp4 = tmp1 * tmp3
+    tl.store(out_ptr0 + (tl.full([XBLOCK], 0, tl.int32).broadcast_to(XBLOCK)), tmp4, None)
+    tl.store(out_ptr1 + (tl.full([XBLOCK], 0, tl.int32).broadcast_to(XBLOCK)), tmp4, None)

.inductor_cache/yt/cytkblpmxtiyttie4sb6q24fmxvz3keyhvpd7ohvwhco6atydhtq.py

@@ -0,0 +1,24 @@
+
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 1}, 
+    filename=__file__,
+    triton_meta={'signature': {'out_ptr0': '*i64', 'ks0': 'i64', 'xnumel': 'constexpr', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=6, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {'xnumel': 1}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_max_view_3', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 0, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_add_max_view_3(out_ptr0, ks0, xnumel, XBLOCK : tl.constexpr):
+    xnumel = 1
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = tl.full([XBLOCK], True, tl.int1)[:]
+    tmp0 = ks0
+    tl.store(out_ptr0 + (tl.full([XBLOCK], 0, tl.int32)), tmp0, None)

.inductor_cache/yu/cyulr2ljflwhtei5asluhqjh5vawyptyitwiqbkohzzxhkthatlm.py

@@ -0,0 +1,29 @@
+
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 1}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'out_ptr0': '*fp32', 'out_ptr1': '*fp32', 'xnumel': 'constexpr', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=0, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {'xnumel': 1}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_max_mul_unsqueeze_view_0', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 2, 'num_store': 2, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_max_mul_unsqueeze_view_0(in_ptr0, in_ptr1, out_ptr0, out_ptr1, xnumel, XBLOCK : tl.constexpr):
+    xnumel = 1
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = tl.full([XBLOCK], True, tl.int1)[:]
+    tmp0 = tl.load(in_ptr0 + (0))
+    tmp1 = tl.broadcast_to(tmp0, [XBLOCK])
+    tmp2 = tl.load(in_ptr1 + (0))
+    tmp3 = tl.broadcast_to(tmp2, [XBLOCK])
+    tmp4 = tmp1 * tmp3
+    tl.store(out_ptr0 + (tl.full([XBLOCK], 0, tl.int32).broadcast_to(XBLOCK)), tmp4, None)
+    tl.store(out_ptr1 + (tl.full([XBLOCK], 0, tl.int32).broadcast_to(XBLOCK)), tmp4, None)

.inductor_cache/z4/cz4ihhmoiu4l5unvbgrsavh73dkatd6zjj7dj77vlmsov6yj6cdy.py

@@ -0,0 +1,137 @@
+# AOT ID: ['10_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 cmath import nanj
+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
+import triton
+import triton.language as tl
+from torch._inductor.runtime.triton_heuristics import 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
+assert_alignment = torch._C._dynamo.guards.assert_alignment
+empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
+empty_strided_cpu_pinned = torch._C._dynamo.guards._empty_strided_cpu_pinned
+empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
+empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
+empty_strided_mtia = torch._C._dynamo.guards._empty_strided_mtia
+reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
+alloc_from_pool = torch.ops.inductor._alloc_from_pool
+async_compile = AsyncCompile()
+empty_strided_p2p = torch._C._distributed_c10d._SymmetricMemory.empty_strided_p2p
+
+
+# kernel path: /mnt/glm5_2/.inductor_cache/yg/cygzxmweoklnvoy2z6e7xc72jmpynqu5ckcmk766j2dvm5llk3js.py
+# Topologically Sorted Source Nodes: [unsqueeze, mul], Original ATen: [aten.unsqueeze, aten.mul]
+# Source node to ATen node mapping:
+#   mul => mul_2
+#   unsqueeze => unsqueeze
+# Graph fragment:
+#   %arg1_1 : Tensor "f32[1, s3][s3, 1]cuda:6" = PlaceHolder[target=arg1_1]
+#   %arg2_1 : Tensor "f32[1, s3, 1][s3, 1, 1]cuda:6" = PlaceHolder[target=arg2_1]
+#   %arg3_1 : Tensor "f64[][]cpu" = PlaceHolder[target=arg3_1]
+#   %unsqueeze : Tensor "f32[1, s3, 1][s3, 1, 1]cuda:6"[num_users=1] = call_function[target=torch.ops.aten.unsqueeze.default](args = (%arg1_1, -1), kwargs = {})
+#   %mul_2 : Tensor "f32[1, s3, 1][s3, 1, 1]cuda:6"[num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%unsqueeze, %arg2_1), kwargs = {})
+#   %convert_element_type_default_1 : Tensor "f32[][]cpu"[num_users=1] = call_function[target=torch.ops.prims.convert_element_type.default](args = (%arg3_1, torch.float32), kwargs = {})
+#   %mul_tensor : Tensor "f32[1, s3, 1][s3, 1, 1]cuda:6"[num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_2, %convert_element_type_default_1), kwargs = {})
+#   return %mul_tensor
+triton_poi_fused_mul_unsqueeze_0 = async_compile.triton('triton_poi_fused_mul_unsqueeze_0', '''
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 32}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'in_ptr2': 'fp64', 'out_ptr0': '*fp32', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=6, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (3,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_mul_unsqueeze_0', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 3, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 512}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_mul_unsqueeze_0(in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, XBLOCK : tl.constexpr):
+    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_ptr1 + (x0), xmask)
+    tmp3 = in_ptr2
+    tmp2 = tmp0 * tmp1
+    tmp4 = tmp3.to(tl.float32)
+    tmp5 = tmp2 * tmp4
+    tl.store(out_ptr0 + (x0), tmp5, xmask)
+''', device_str='cuda')
+
+
+async_compile.wait(globals())
+del async_compile
+
+class Runner:
+    def __init__(self, partitions):
+        self.partitions = partitions
+
+    def recursively_apply_fns(self, fns):
+        new_callables = []
+        for fn, c in zip(fns, self.partitions):
+            new_callables.append(fn(c))
+        self.partitions = new_callables
+
+    def call(self, args):
+        arg0_1, arg1_1, arg2_1, arg3_1 = args
+        args.clear()
+        s3 = arg0_1
+        assert_size_stride(arg1_1, (1, s3), (s3, 1))
+        assert_size_stride(arg2_1, (1, s3, 1), (s3, 1, 1))
+        assert_size_stride(arg3_1, (), ())
+        with torch.cuda._DeviceGuard(6):
+            torch.cuda.set_device(6)
+            buf0 = empty_strided_cuda((1, s3, 1), (s3, 1, 1), torch.float32)
+            # Topologically Sorted Source Nodes: [unsqueeze, mul], Original ATen: [aten.unsqueeze, aten.mul]
+            stream6 = get_raw_stream(6)
+            triton_poi_fused_mul_unsqueeze_0.run(arg1_1, arg2_1, arg3_1.item(), buf0, s3, stream=stream6)
+            del arg1_1
+            del arg2_1
+            del arg3_1
+        return (buf0, )
+
+runner = Runner(partitions=[])
+call = runner.call
+recursively_apply_fns = runner.recursively_apply_fns
+
+
+def get_args():
+    from torch._dynamo.testing import rand_strided
+    arg0_1 = 32
+    arg1_1 = rand_strided((1, 32), (32, 1), device='cuda:6', dtype=torch.float32)
+    arg2_1 = rand_strided((1, 32, 1), (32, 1, 1), device='cuda:6', dtype=torch.float32)
+    arg3_1 = rand_strided((), (), device='cpu', dtype=torch.float64)
+    return [arg0_1, arg1_1, arg2_1, arg3_1]
+
+
+def benchmark_compiled_module(args, times=10, repeat=10):
+    from torch._inductor.utils import print_performance
+    fn = lambda: call(list(args))
+    return print_performance(fn, times=times, repeat=repeat)
+
+
+if __name__ == "__main__":
+    from torch._inductor.wrapper_benchmark import compiled_module_main
+    args = get_args()
+    compiled_module_main('None', lambda times, repeat: benchmark_compiled_module(args, times=times, repeat=repeat))

.inductor_cache/zb/czbs2fm46j6qhtd2znnofgeavkuhyl72vdcahggz7dhvxp5zvnma.py

@@ -0,0 +1,24 @@
+
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 8}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*i64', 'out_ptr0': '*i64', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=5, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_gather_max_view_1', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 1, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 192}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_gather_max_view_1(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = xindex < xnumel
+    x0 = xindex
+    tmp0 = tl.load(in_ptr0 + (x0), xmask)
+    tl.store(out_ptr0 + (x0), tmp0, xmask)

.inductor_cache/ze/1124bb1f7f39db53a57e49f4d36648a0d2fb0d8c5d95047e9e79c4f5ff6aa834.best_config

@@ -0,0 +1 @@
+{"XBLOCK": 128, "num_warps": 4, "num_stages": 1, "configs_hash": "1b2cc4dbebb9680d3ce31843331593b159e4046c056f195ca1ccf2464d5b37d1", "found_by_coordesc": false, "time_taken_ms": 82, "triton_cache_hash": "LBOXSWDAX2GIZFOXRIZQ4S3246ZSKZAMDEVIL3FUHRGF3W6BJYLQ"}

.inductor_cache/ze/czeardixathkscotfnp37pocvapthzwtdabgtpsto2evnxhd2wi7.py

@@ -0,0 +1,26 @@
+
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 256}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_out_ptr0': '*fp32', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=2, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_mul_0', 'mutated_arg_names': ['in_out_ptr0'], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 1, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 3072}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_mul_0(in_out_ptr0, xnumel, XBLOCK : tl.constexpr):
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = xindex < xnumel
+    x0 = xindex
+    tmp0 = tl.load(in_out_ptr0 + (x0), xmask)
+    tmp1 = tl.full([1], 0.1767766952966369, tl.float32)
+    tmp2 = tmp0 * tmp1
+    tl.store(in_out_ptr0 + (x0), tmp2, xmask)

.inductor_cache/zi/czi2gebq2ez4fkm4twycyxki3nxsr5q5asfo44xgpsj5txniaoey.py

@@ -0,0 +1,29 @@
+
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 32}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'in_ptr2': 'fp64', 'out_ptr0': '*fp32', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=3, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (3,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_mul_unsqueeze_0', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 3, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 512}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_mul_unsqueeze_0(in_ptr0, in_ptr1, in_ptr2, out_ptr0, xnumel, XBLOCK : tl.constexpr):
+    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_ptr1 + (x0), xmask)
+    tmp3 = in_ptr2
+    tmp2 = tmp0 * tmp1
+    tmp4 = tmp3.to(tl.float32)
+    tmp5 = tmp2 * tmp4
+    tl.store(out_ptr0 + (x0), tmp5, xmask)

.inductor_cache/zm/czmqgy6soqztyipifu6p2jfwi4je2rv5dgpxsqvvy3vdykavyemd.py

@@ -0,0 +1,24 @@
+
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 65536}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*bf16', 'out_ptr0': '*bf16', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=5, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 1, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 294912}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = xindex < xnumel
+    x2 = xindex
+    tmp0 = tl.load(in_ptr0 + (x2), xmask, eviction_policy='evict_last').to(tl.float32)
+    tl.store(out_ptr0 + (x2), tmp0, xmask)

.inductor_cache/zm/faeba8505e0ed9390c51709b1a4837ad201756af55501c251609791851cf944b.best_config

@@ -0,0 +1 @@
+{"XBLOCK": 256, "num_warps": 4, "num_stages": 1, "configs_hash": "7cced77f371acaa5aa7d90332a90e0c907727cfefb71d9cc9d997c24557fc44f", "found_by_coordesc": false, "time_taken_ms": 22, "triton_cache_hash": "QLXJTPJWIG7LHGEN76P7EYFUJIJLXLDECRO3I6NRHSQ4NBZC2KMA"}

.inductor_cache/zp/czpgtwy5of2i6b7aikxrgrez66yzismukpkcmrttlp5pnenqldlz.py

@@ -0,0 +1,150 @@
+# AOT ID: ['3_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 cmath import nanj
+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
+import triton
+import triton.language as tl
+from torch._inductor.runtime.triton_heuristics import 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
+assert_alignment = torch._C._dynamo.guards.assert_alignment
+empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
+empty_strided_cpu_pinned = torch._C._dynamo.guards._empty_strided_cpu_pinned
+empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
+empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
+empty_strided_mtia = torch._C._dynamo.guards._empty_strided_mtia
+reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
+alloc_from_pool = torch.ops.inductor._alloc_from_pool
+async_compile = AsyncCompile()
+empty_strided_p2p = torch._C._distributed_c10d._SymmetricMemory.empty_strided_p2p
+
+
+# kernel path: /mnt/glm5_2/.inductor_cache/4l/c4ljp5q37q7ktbyo35672uzle4llt4iwt4be37ikkwlit6j7fnco.py
+# Topologically Sorted Source Nodes: [weights_1, unsqueeze, mul_1], Original ATen: [aten.mul, aten.unsqueeze]
+# Source node to ATen node mapping:
+#   mul_1 => mul_7
+#   unsqueeze => unsqueeze
+#   weights_1 => mul_2
+# Graph fragment:
+#   %mm : Tensor "f32[s0, 32][32, 1]cuda:7" = PlaceHolder[target=mm]
+#   %arg3_1 : Tensor "f32[s0, 32, 1][32, 1, 1]cuda:7" = PlaceHolder[target=arg3_1]
+#   %arg4_1 : Tensor "f64[][]cpu" = PlaceHolder[target=arg4_1]
+#   %mul_2 : Tensor "f32[s0, 32][32, 1]cuda:7"[num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mm, 0.1767766952966369), kwargs = {})
+#   %unsqueeze : Tensor "f32[s0, 32, 1][32, 1, 1]cuda:7"[num_users=1] = call_function[target=torch.ops.aten.unsqueeze.default](args = (%mul_2, -1), kwargs = {})
+#   %mul_7 : Tensor "f32[s0, 32, 1][32, 1, 1]cuda:7"[num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%unsqueeze, %arg3_1), kwargs = {})
+#   %convert_element_type_default_1 : Tensor "f32[][]cpu"[num_users=1] = call_function[target=torch.ops.prims.convert_element_type.default](args = (%arg4_1, torch.float32), kwargs = {})
+#   %mul_tensor : Tensor "f32[s0, 32, 1][32, 1, 1]cuda:7"[num_users=1] = call_function[target=torch.ops.aten.mul.Tensor](args = (%mul_7, %convert_element_type_default_1), kwargs = {})
+#   return %mul_tensor
+triton_poi_fused_mul_unsqueeze_0 = async_compile.triton('triton_poi_fused_mul_unsqueeze_0', '''
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 2048}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_out_ptr0': '*fp32', 'in_ptr0': '*fp32', 'in_ptr1': 'fp64', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=7, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (3,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_mul_unsqueeze_0', 'mutated_arg_names': ['in_out_ptr0'], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 3, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 28672}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_mul_unsqueeze_0(in_out_ptr0, in_ptr0, in_ptr1, xnumel, XBLOCK : tl.constexpr):
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = xindex < xnumel
+    x0 = xindex
+    tmp0 = tl.load(in_out_ptr0 + (x0), xmask)
+    tmp3 = tl.load(in_ptr0 + (x0), xmask)
+    tmp5 = in_ptr1
+    tmp1 = tl.full([1], 0.1767766952966369, tl.float32)
+    tmp2 = tmp0 * tmp1
+    tmp4 = tmp2 * tmp3
+    tmp6 = tmp5.to(tl.float32)
+    tmp7 = tmp4 * tmp6
+    tl.store(in_out_ptr0 + (x0), tmp7, xmask)
+''', device_str='cuda')
+
+
+async_compile.wait(globals())
+del async_compile
+
+class Runner:
+    def __init__(self, partitions):
+        self.partitions = partitions
+
+    def recursively_apply_fns(self, fns):
+        new_callables = []
+        for fn, c in zip(fns, self.partitions):
+            new_callables.append(fn(c))
+        self.partitions = new_callables
+
+    def call(self, args):
+        arg0_1, arg1_1, arg2_1, arg3_1, arg4_1 = args
+        args.clear()
+        s77 = arg1_1
+        arg2_1_size = arg2_1.size()
+        s0 = arg2_1_size[0]
+        assert_size_stride(arg0_1, (32, 6144), (6144, 1))
+        assert_size_stride(arg2_1, (s0, 6144), (6144, 1))
+        assert_size_stride(arg3_1, (s0, 32, 1), (32, 1, 1))
+        assert_size_stride(arg4_1, (), ())
+        with torch.cuda._DeviceGuard(7):
+            torch.cuda.set_device(7)
+            buf0 = empty_strided_cuda((s0, 32), (32, 1), torch.float32)
+            # Topologically Sorted Source Nodes: [t, weights], Original ATen: [aten.t, aten.mm]
+            extern_kernels.mm_dtype(arg2_1, reinterpret_tensor(arg0_1, (6144, 32), (1, 6144), 0), out_dtype=torch.float32, out=buf0)
+            del arg0_1
+            del arg2_1
+            buf1 = reinterpret_tensor(buf0, (s0, 32, 1), (32, 1, 1), 0); del buf0  # reuse
+            # Topologically Sorted Source Nodes: [weights_1, unsqueeze, mul_1], Original ATen: [aten.mul, aten.unsqueeze]
+            triton_poi_fused_mul_unsqueeze_0_xnumel = 32*s0
+            stream7 = get_raw_stream(7)
+            triton_poi_fused_mul_unsqueeze_0.run(buf1, arg3_1, arg4_1.item(), triton_poi_fused_mul_unsqueeze_0_xnumel, stream=stream7)
+            del arg3_1
+            del arg4_1
+        return (buf1, )
+
+runner = Runner(partitions=[])
+call = runner.call
+recursively_apply_fns = runner.recursively_apply_fns
+
+
+def get_args():
+    from torch._dynamo.testing import rand_strided
+    arg0_1 = rand_strided((32, 6144), (6144, 1), device='cuda:7', dtype=torch.bfloat16)
+    arg1_1 = 56
+    arg2_1 = rand_strided((56, 6144), (6144, 1), device='cuda:7', dtype=torch.bfloat16)
+    arg3_1 = rand_strided((56, 32, 1), (32, 1, 1), device='cuda:7', dtype=torch.float32)
+    arg4_1 = rand_strided((), (), device='cpu', dtype=torch.float64)
+    return [arg0_1, arg1_1, arg2_1, arg3_1, arg4_1]
+
+
+def benchmark_compiled_module(args, times=10, repeat=10):
+    from torch._inductor.utils import print_performance
+    fn = lambda: call(list(args))
+    return print_performance(fn, times=times, repeat=repeat)
+
+
+if __name__ == "__main__":
+    from torch._inductor.wrapper_benchmark import compiled_module_main
+    args = get_args()
+    compiled_module_main('None', lambda times, repeat: benchmark_compiled_module(args, times=times, repeat=repeat))

.inductor_cache/zv/czvm6uj6zs2duwbjqoq4cmqrfz7ejc3igihrhqvth7uov36wgnuo.py

@@ -0,0 +1,289 @@
+# AOT ID: ['6_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 cmath import nanj
+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
+import triton
+import triton.language as tl
+from torch._inductor.runtime.triton_heuristics import 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
+assert_alignment = torch._C._dynamo.guards.assert_alignment
+empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
+empty_strided_cpu_pinned = torch._C._dynamo.guards._empty_strided_cpu_pinned
+empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
+empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
+empty_strided_mtia = torch._C._dynamo.guards._empty_strided_mtia
+reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
+alloc_from_pool = torch.ops.inductor._alloc_from_pool
+async_compile = AsyncCompile()
+empty_strided_p2p = torch._C._distributed_c10d._SymmetricMemory.empty_strided_p2p
+
+
+# kernel path: /mnt/glm5_2/.inductor_cache/r3/cr333mq2yzr7rixiiwaacpyfykcrh3oe4bimo3cngi42juonzohi.py
+# Topologically Sorted Source Nodes: [unsqueeze, view, expand_scores, flatten, max_1], Original ATen: [aten.unsqueeze, aten.view, aten.mul, aten.max]
+# Source node to ATen node mapping:
+#   expand_scores => mul_3
+#   flatten => view_1
+#   max_1 => getitem, max_1
+#   unsqueeze => unsqueeze
+#   view => view
+# Graph fragment:
+#   %arg2_1 : Tensor "f32[s97, 1][1, 1]cuda:1" = PlaceHolder[target=arg2_1]
+#   %arg3_1 : Tensor "f32[s97, 1][1, 1]cuda:1" = PlaceHolder[target=arg3_1]
+#   %mul_3 : Tensor "f32[s97, 1, 1][1, 1, 1]cuda:1" = PlaceHolder[target=mul_3]
+#   %unsqueeze : Tensor "f32[s97, 1, 1][1, 1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.unsqueeze.default](args = (%arg2_1, 2), kwargs = {})
+#   %view : Tensor "f32[s97, 1, 1][1, 1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.reshape.default](args = (%arg3_1, [-1, 1, 1]), kwargs = {})
+#   %mul_3 : Tensor "f32[s97, 1, 1][1, 1, 1]cuda:1"[num_users=2] = call_function[target=torch.ops.aten.mul.Tensor](args = (%unsqueeze, %view), kwargs = {})
+#   %view_1 : Tensor "f32[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.reshape.default](args = (%mul_3, [%arg1_1, 1]), kwargs = {})
+#   %max_1 : [num_users=2] = call_function[target=torch.ops.aten.max.dim](args = (%view_1, -1, True), kwargs = {})
+#   %getitem : Tensor "f32[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=operator.getitem](args = (%max_1, 0), kwargs = {})
+#   return %mul_3,%getitem
+triton_poi_fused_max_mul_unsqueeze_view_0 = async_compile.triton('triton_poi_fused_max_mul_unsqueeze_view_0', '''
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 8}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*fp32', 'in_ptr1': '*fp32', 'out_ptr0': '*fp32', 'out_ptr1': '*fp32', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=1, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]], (2,): [['tt.divisibility', 16]], (3,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_max_mul_unsqueeze_view_0', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 2, 'num_store': 2, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 192}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_max_mul_unsqueeze_view_0(in_ptr0, in_ptr1, out_ptr0, out_ptr1, xnumel, XBLOCK : tl.constexpr):
+    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_ptr1 + (x0), xmask)
+    tmp2 = tmp0 * tmp1
+    tl.store(out_ptr0 + (x0), tmp2, xmask)
+    tl.store(out_ptr1 + (x0), tmp2, xmask)
+''', device_str='cuda')
+
+
+# kernel path: /mnt/glm5_2/.inductor_cache/un/cun3h5xetskkgqj652kojes6opfqbaexbkzn5bq2lhsitna2pl4x.py
+# Topologically Sorted Source Nodes: [flatten, max_1, gather], Original ATen: [aten.view, aten.max, aten.gather]
+# Source node to ATen node mapping:
+#   flatten => view_1
+#   gather => gather
+#   max_1 => max_1
+# Graph fragment:
+#   %arg4_1 : Tensor "i64[s97, 1][1, 1]cuda:1" = PlaceHolder[target=arg4_1]
+#   %view_1 : Tensor "f32[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.reshape.default](args = (%mul_3, [%arg1_1, 1]), kwargs = {})
+#   %max_1 : [num_users=2] = call_function[target=torch.ops.aten.max.dim](args = (%view_1, -1, True), kwargs = {})
+#   %gather : Tensor "i64[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.gather.default](args = (%arg4_1, 1, %getitem_1), kwargs = {})
+#   return %gather
+triton_poi_fused_gather_max_view_1 = async_compile.triton('triton_poi_fused_gather_max_view_1', '''
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 8}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*i64', 'out_ptr0': '*i64', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=1, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_gather_max_view_1', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 1, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 192}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_gather_max_view_1(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = xindex < xnumel
+    x0 = xindex
+    tmp0 = tl.load(in_ptr0 + (x0), xmask)
+    tl.store(out_ptr0 + (x0), tmp0, xmask)
+''', device_str='cuda')
+
+
+# kernel path: /mnt/glm5_2/.inductor_cache/f4/cf4amfzddy5edtg5hhvqyznxzmqz3sy5hooeuillcupuqjry5zoh.py
+# Topologically Sorted Source Nodes: [flatten, max_1, flat_cs, floordiv, batch_offsets, repeat_interleave, selected_input_index, hidden_states], Original ATen: [aten.view, aten.max, aten.floor_divide, aten.arange, aten.unsqueeze, aten.add, aten.index]
+# Source node to ATen node mapping:
+#   batch_offsets => iota
+#   flat_cs => view_4
+#   flatten => view_1
+#   floordiv => div
+#   hidden_states => index
+#   max_1 => max_1
+#   repeat_interleave => unsqueeze_1, view_5
+#   selected_input_index => add_47
+# Graph fragment:
+#   %arg6_1 : Tensor "bf16[s97, s87][s87, 1]cuda:1" = PlaceHolder[target=arg6_1]
+#   %view_1 : Tensor "f32[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.reshape.default](args = (%mul_3, [%arg1_1, 1]), kwargs = {})
+#   %max_1 : [num_users=2] = call_function[target=torch.ops.aten.max.dim](args = (%view_1, -1, True), kwargs = {})
+#   %view_4 : Tensor "i64[s97][1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.reshape.default](args = (%getitem_1, [%arg1_1]), kwargs = {})
+#   %div : Tensor "i64[s97][1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.div.Tensor_mode](args = (%view_4, 1), kwargs = {rounding_mode: floor})
+#   %iota : Tensor "i64[s97][1]cuda:1"[num_users=1] = call_function[target=torch.ops.prims.iota.default](args = (%arg1_1,), kwargs = {start: 0, step: 1, dtype: torch.int64, device: cuda:1, requires_grad: False})
+#   %unsqueeze_1 : Tensor "i64[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.unsqueeze.default](args = (%iota, 1), kwargs = {})
+#   %view_5 : Tensor "i64[s97][1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.reshape.default](args = (%expand, [%arg1_1]), kwargs = {})
+#   %add_47 : Tensor "i64[s97][1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%div, %view_5), kwargs = {})
+#   %index : Tensor "bf16[s97, s87][s87, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.index.Tensor](args = (%arg6_1, [%add_47]), kwargs = {})
+#   return %index
+triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2 = async_compile.triton('triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2', '''
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 65536}, 
+    filename=__file__,
+    triton_meta={'signature': {'in_ptr0': '*bf16', 'out_ptr0': '*bf16', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=1, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]], (1,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 1, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 294912}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2(in_ptr0, out_ptr0, xnumel, XBLOCK : tl.constexpr):
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = xindex < xnumel
+    x2 = xindex
+    tmp0 = tl.load(in_ptr0 + (x2), xmask, eviction_policy='evict_last').to(tl.float32)
+    tl.store(out_ptr0 + (x2), tmp0, xmask)
+''', device_str='cuda')
+
+
+# kernel path: /mnt/glm5_2/.inductor_cache/tk/ctkv5ncbcvcsa2sqv7fy6waloep7rvxvcgauk5pfktnei5g2ypuk.py
+# Topologically Sorted Source Nodes: [flatten, max_1, add_2], Original ATen: [aten.view, aten.max, aten.add]
+# Source node to ATen node mapping:
+#   add_2 => add_53
+#   flatten => view_1
+#   max_1 => max_1
+# Graph fragment:
+#   %view_1 : Tensor "f32[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.reshape.default](args = (%mul_3, [%arg1_1, 1]), kwargs = {})
+#   %max_1 : [num_users=2] = call_function[target=torch.ops.aten.max.dim](args = (%view_1, -1, True), kwargs = {})
+#   %add_53 : Tensor "i64[s97, 1][1, 1]cuda:1"[num_users=1] = call_function[target=torch.ops.aten.add.Tensor](args = (%getitem_1, %arg7_1), kwargs = {})
+#   return %add_53
+triton_poi_fused_add_max_view_3 = async_compile.triton('triton_poi_fused_add_max_view_3', '''
+import triton
+import triton.language as tl
+
+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, DeviceProperties
+triton_helpers.set_driver_to_gpu()
+
+@triton_heuristics.pointwise(
+    size_hints={'x': 8}, 
+    filename=__file__,
+    triton_meta={'signature': {'out_ptr0': '*i64', 'ks0': 'i64', 'xnumel': 'i32', 'XBLOCK': 'constexpr'}, 'device': DeviceProperties(type='cuda', index=1, multi_processor_count=148, cc=100, major=10, regs_per_multiprocessor=65536, max_threads_per_multi_processor=2048, max_threads_per_block=1024, warp_size=32), 'constants': {}, 'native_matmul': False, 'enable_fp_fusion': True, 'launch_pdl': False, 'disable_ftz': False, 'configs': [{(0,): [['tt.divisibility', 16]]}]},
+    inductor_meta={'grid_type': 'Grid1D', 'autotune_hints': set(), 'kernel_name': 'triton_poi_fused_add_max_view_3', 'mutated_arg_names': [], 'optimize_mem': True, 'no_x_dim': False, 'atomic_add_found': False, 'num_load': 0, 'num_store': 1, 'num_reduction': 0, 'backend_hash': 'B9347ED8E82563C0DF6E4C45E8A40BBC93F818749C1452A3CD0A7365095FE2BD', '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, 'deterministic': False, 'force_filter_reduction_configs': False, 'mix_order_reduction_allow_multi_stages': False, 'are_deterministic_algorithms_enabled': False, 'tiling_scores': {'x': 128}},
+    min_elem_per_thread=0
+)
+@triton.jit
+def triton_poi_fused_add_max_view_3(out_ptr0, ks0, xnumel, XBLOCK : tl.constexpr):
+    xoffset = tl.program_id(0) * XBLOCK
+    xindex = xoffset + tl.arange(0, XBLOCK)[:]
+    xmask = xindex < xnumel
+    x0 = xindex
+    tmp0 = ks0
+    tl.store(out_ptr0 + (x0), tmp0, xmask)
+''', device_str='cuda')
+
+
+async_compile.wait(globals())
+del async_compile
+
+class Runner:
+    def __init__(self, partitions):
+        self.partitions = partitions
+
+    def recursively_apply_fns(self, fns):
+        new_callables = []
+        for fn, c in zip(fns, self.partitions):
+            new_callables.append(fn(c))
+        self.partitions = new_callables
+
+    def call(self, args):
+        arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1 = args
+        args.clear()
+        s53 = arg0_1
+        s97 = arg1_1
+        s87 = arg5_1
+        s77 = arg7_1
+        assert_size_stride(arg2_1, (s97, 1), (1, 1))
+        assert_size_stride(arg3_1, (s97, 1), (1, 1))
+        assert_size_stride(arg4_1, (s97, 1), (1, 1))
+        assert_size_stride(arg6_1, (s97, s87), (s87, 1))
+        with torch.cuda._DeviceGuard(1):
+            torch.cuda.set_device(1)
+            buf0 = empty_strided_cuda((s97, 1, 1), (1, 1, 1), torch.float32)
+            buf1 = empty_strided_cuda((s97, 1), (1, 1), torch.float32)
+            # Topologically Sorted Source Nodes: [unsqueeze, view, expand_scores, flatten, max_1], Original ATen: [aten.unsqueeze, aten.view, aten.mul, aten.max]
+            stream1 = get_raw_stream(1)
+            triton_poi_fused_max_mul_unsqueeze_view_0.run(arg2_1, arg3_1, buf0, buf1, s97, stream=stream1)
+            del arg2_1
+            del arg3_1
+            buf2 = empty_strided_cuda((s97, 1), (1, 1), torch.int64)
+            # Topologically Sorted Source Nodes: [flatten, max_1, gather], Original ATen: [aten.view, aten.max, aten.gather]
+            stream1 = get_raw_stream(1)
+            triton_poi_fused_gather_max_view_1.run(arg4_1, buf2, s97, stream=stream1)
+            buf3 = empty_strided_cuda((s97, s87), (s87, 1), torch.bfloat16)
+            # Topologically Sorted Source Nodes: [flatten, max_1, flat_cs, floordiv, batch_offsets, repeat_interleave, selected_input_index, hidden_states], Original ATen: [aten.view, aten.max, aten.floor_divide, aten.arange, aten.unsqueeze, aten.add, aten.index]
+            triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2_xnumel = s87*s97
+            stream1 = get_raw_stream(1)
+            triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2.run(arg6_1, buf3, triton_poi_fused_add_arange_floor_divide_index_max_unsqueeze_view_2_xnumel, stream=stream1)
+            del arg6_1
+            buf4 = empty_strided_cuda((s97, 1), (1, 1), torch.int64)
+            # Topologically Sorted Source Nodes: [flatten, max_1, add_2], Original ATen: [aten.view, aten.max, aten.add]
+            stream1 = get_raw_stream(1)
+            triton_poi_fused_add_max_view_3.run(buf4, s77, s97, stream=stream1)
+        return (reinterpret_tensor(buf2, (s97, ), (1, ), 0), buf3, buf1, buf0, arg4_1, buf4, )
+
+runner = Runner(partitions=[])
+call = runner.call
+recursively_apply_fns = runner.recursively_apply_fns
+
+
+def get_args():
+    from torch._dynamo.testing import rand_strided
+    arg0_1 = 1
+    arg1_1 = 8
+    arg2_1 = rand_strided((8, 1), (1, 1), device='cuda:1', dtype=torch.float32)
+    arg3_1 = rand_strided((8, 1), (1, 1), device='cuda:1', dtype=torch.float32)
+    arg4_1 = rand_strided((8, 1), (1, 1), device='cuda:1', dtype=torch.int64)
+    arg5_1 = 6144
+    arg6_1 = rand_strided((8, 6144), (6144, 1), device='cuda:1', dtype=torch.bfloat16)
+    arg7_1 = 1
+    return [arg0_1, arg1_1, arg2_1, arg3_1, arg4_1, arg5_1, arg6_1, arg7_1]
+
+
+def benchmark_compiled_module(args, times=10, repeat=10):
+    from torch._inductor.utils import print_performance
+    fn = lambda: call(list(args))
+    return print_performance(fn, times=times, repeat=repeat)
+
+
+if __name__ == "__main__":
+    from torch._inductor.wrapper_benchmark import compiled_module_main
+    args = get_args()
+    compiled_module_main('None', lambda times, repeat: benchmark_compiled_module(args, times=times, repeat=repeat))