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 import triton_helpers
from torch._inductor.runtime.triton_helpers import math as tl_math
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
reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor


@triton.jit
def triton_poi_fused_convolution_relu_0(in_ptr0, in_ptr1, out_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 1906624
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x4 = xindex
    x2 = xindex // 238328 % 2
    x0 = xindex % 3844
    x5 = xindex // 3844
    tmp0 = tl.load(in_ptr0 + x4, xmask)
    tmp1 = tl.load(in_ptr1 + x2, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tl.store(out_ptr0 + (x0 + 3872 * x5), tmp4, xmask)


@triton.jit
def triton_poi_fused_convolution_relu_1(in_ptr0, in_ptr1, out_ptr0, xnumel,
    XBLOCK: tl.constexpr):
    xnumel = 390224
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x3 = xindex
    x1 = xindex // 24389 % 4
    x0 = xindex % 24389
    x4 = xindex // 24389
    tmp0 = tl.load(in_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr1 + x1, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tl.store(out_ptr0 + (x0 + 24416 * x4), tmp4, xmask)


@triton.jit
def triton_poi_fused_relu_2(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl.constexpr
    ):
    xnumel = 87808
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x0 = xindex % 8
    tmp0 = tl.load(in_out_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + x0, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tl.store(in_out_ptr0 + x2, tmp4, xmask)


@triton.jit
def triton_per_fused__softmax_3(in_ptr0, out_ptr2, xnumel, rnumel, XBLOCK:
    tl.constexpr):
    xnumel = 10976
    rnumel = 10
    RBLOCK: tl.constexpr = 16
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    rmask = rindex < rnumel
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + 10 * x0), rmask & xmask, other=0.0)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.where(rmask & xmask, tmp1, float('-inf'))
    tmp4 = triton_helpers.max2(tmp3, 1)[:, None]
    tmp5 = tmp0 - tmp4
    tmp6 = tl_math.exp(tmp5)
    tmp7 = tl.broadcast_to(tmp6, [XBLOCK, RBLOCK])
    tmp9 = tl.where(rmask & xmask, tmp7, 0)
    tmp10 = tl.sum(tmp9, 1)[:, None]
    tmp11 = tmp6 / tmp10
    tl.store(out_ptr2 + (r1 + 10 * x0), tmp11, rmask & xmask)


def call(args):
    (primals_1, primals_2, primals_3, primals_4, primals_5, primals_6,
        primals_7, primals_8, primals_9) = args
    args.clear()
    assert_size_stride(primals_1, (2, 1, 3, 3, 3), (27, 27, 9, 3, 1))
    assert_size_stride(primals_2, (2,), (1,))
    assert_size_stride(primals_3, (4, 1, 64, 64, 64), (262144, 262144, 4096,
        64, 1))
    assert_size_stride(primals_4, (4, 2, 3, 3, 3), (54, 27, 9, 3, 1))
    assert_size_stride(primals_5, (4,), (1,))
    assert_size_stride(primals_6, (8, 4), (4, 1))
    assert_size_stride(primals_7, (8,), (1,))
    assert_size_stride(primals_8, (10, 8), (8, 1))
    assert_size_stride(primals_9, (10,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = extern_kernels.convolution(primals_3, primals_1, stride=(1, 
            1, 1), padding=(0, 0, 0), dilation=(1, 1, 1), transposed=False,
            output_padding=(0, 0, 0), groups=1, bias=None)
        assert_size_stride(buf0, (4, 2, 62, 62, 62), (476656, 238328, 3844,
            62, 1))
        buf1 = empty_strided_cuda((4, 2, 62, 62, 62), (480128, 240064, 3872,
            62, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_convolution_relu_0[grid(1906624)](buf0, primals_2,
            buf1, 1906624, XBLOCK=1024, num_warps=4, num_stages=1)
        del buf0
        del primals_2
        buf2 = torch.ops.aten.max_pool3d_with_indices.default(buf1, [2, 2, 
            2], [2, 2, 2])
        buf3 = buf2[0]
        buf4 = buf2[1]
        del buf2
        buf5 = extern_kernels.convolution(buf3, primals_4, stride=(1, 1, 1),
            padding=(0, 0, 0), dilation=(1, 1, 1), transposed=False,
            output_padding=(0, 0, 0), groups=1, bias=None)
        assert_size_stride(buf5, (4, 4, 29, 29, 29), (97556, 24389, 841, 29, 1)
            )
        buf6 = empty_strided_cuda((4, 4, 29, 29, 29), (97664, 24416, 841, 
            29, 1), torch.float32)
        triton_poi_fused_convolution_relu_1[grid(390224)](buf5, primals_5,
            buf6, 390224, XBLOCK=1024, num_warps=4, num_stages=1)
        del buf5
        del primals_5
        buf7 = torch.ops.aten.max_pool3d_with_indices.default(buf6, [2, 2, 
            2], [2, 2, 2])
        buf8 = buf7[0]
        buf9 = buf7[1]
        del buf7
        buf10 = empty_strided_cuda((10976, 8), (8, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(buf8, (10976, 4), (4, 1), 0),
            reinterpret_tensor(primals_6, (4, 8), (1, 4), 0), out=buf10)
        buf11 = buf10
        del buf10
        triton_poi_fused_relu_2[grid(87808)](buf11, primals_7, 87808,
            XBLOCK=1024, num_warps=4, num_stages=1)
        del primals_7
        buf12 = empty_strided_cuda((10976, 10), (10, 1), torch.float32)
        extern_kernels.addmm(primals_9, buf11, reinterpret_tensor(primals_8,
            (8, 10), (1, 8), 0), alpha=1, beta=1, out=buf12)
        del primals_9
        buf15 = empty_strided_cuda((10976, 10), (10, 1), torch.float32)
        triton_per_fused__softmax_3[grid(10976)](buf12, buf15, 10976, 10,
            XBLOCK=128, num_warps=8, num_stages=1)
        del buf12
    return (buf15, primals_1, primals_3, primals_4, buf1, buf3, buf4, buf6,
        buf9, reinterpret_tensor(buf8, (10976, 4), (4, 1), 0), buf11, buf15,
        primals_8, primals_6)


class BasicModel_ConvNet_MaxPool3dNew(nn.Module):
    """Same as above, but with the MaxPool1d replaced
    with a MaxPool3d. This is useful because the MaxPool modules
    behave differently to other modules from the perspective
    of the DeepLift Attributions
    """

    def __init__(self) ->None:
        super().__init__()
        self.conv1 = nn.Conv3d(1, 2, 3)
        self.relu1 = nn.ReLU()
        self.pool1 = nn.MaxPool3d(2)
        self.conv2 = nn.Conv3d(2, 4, 3)
        self.relu2 = nn.ReLU()
        self.pool2 = nn.MaxPool3d(2)
        self.fc1 = nn.Linear(4, 8)
        self.relu3 = nn.ReLU()
        self.fc2 = nn.Linear(8, 10)
        self.softmax = nn.Softmax(dim=1)
        self.fc1.weight = nn.Parameter(torch.ones(8, 4))
        self.fc2.weight = nn.Parameter(torch.ones(10, 8))

    def forward(self, input_0):
        primals_1 = self.conv1.weight
        primals_2 = self.conv1.bias
        primals_4 = self.conv2.weight
        primals_5 = self.conv2.bias
        primals_6 = self.fc1.weight
        primals_7 = self.fc1.bias
        primals_8 = self.fc2.weight
        primals_9 = self.fc2.bias
        primals_3 = input_0
        output = call([primals_1, primals_2, primals_3, primals_4,
            primals_5, primals_6, primals_7, primals_8, primals_9])
        return output[0]


Measured latency_ms: 0.025752