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
import math
from torch import nn
from torch.nn.parameter import Parameter
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_mul_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 + x0, xmask)
    tmp1 = 0.0001
    tmp2 = tmp0 * tmp1
    tl.store(out_ptr0 + x0, tmp2, xmask)


@triton.jit
def triton_poi_fused_cat_1(in_ptr0, in_ptr1, in_ptr2, in_ptr3, out_ptr0,
    xnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x1 = xindex // 16
    x0 = xindex % 16
    x2 = xindex
    tmp0 = x1
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 3, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.full([1], 2, tl.int64)
    tmp6 = tmp0 < tmp5
    tmp7 = tmp6 & tmp4
    tmp8 = tl.full([1], 1, tl.int64)
    tmp9 = tmp0 < tmp8
    tmp10 = tmp9 & tmp7
    tmp11 = tl.load(in_ptr0 + x0, tmp10 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp12 = tmp0 >= tmp8
    tmp13 = tmp12 & tmp7
    tmp14 = tl.load(in_ptr1 + x0, tmp13 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp15 = tl.where(tmp9, tmp11, tmp14)
    tmp16 = tl.full(tmp15.shape, 0.0, tmp15.dtype)
    tmp17 = tl.where(tmp7, tmp15, tmp16)
    tmp18 = tmp0 >= tmp5
    tmp19 = tmp18 & tmp4
    tmp20 = tl.load(in_ptr2 + x0, tmp19 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp21 = tl.where(tmp6, tmp17, tmp20)
    tmp22 = tl.full(tmp21.shape, 0.0, tmp21.dtype)
    tmp23 = tl.where(tmp4, tmp21, tmp22)
    tmp24 = tmp0 >= tmp3
    tl.full([1], 4, tl.int64)
    tmp27 = tl.load(in_ptr3 + x0, tmp24 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp28 = tl.where(tmp4, tmp23, tmp27)
    tl.store(out_ptr0 + x2, tmp28, xmask)


@triton.jit
def triton_poi_fused_clone_2(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 4
    x1 = xindex // 4 % 4
    x2 = xindex // 16 % 4
    x4 = xindex
    tmp0 = tl.load(in_ptr0 + (x1 + 4 * x0 + 16 * x2), xmask,
        eviction_policy='evict_last')
    tl.store(out_ptr0 + x4, tmp0, xmask)


@triton.jit
def triton_poi_fused_clone_3(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 256
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 64
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + x0, xmask, eviction_policy='evict_last')
    tl.store(out_ptr0 + x2, tmp0, xmask)


@triton.jit
def triton_poi_fused_sqrt_4(in_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_out_ptr0 + x0, xmask)
    tmp1 = libdevice.sqrt(tmp0)
    tl.store(in_out_ptr0 + x0, tmp1, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (16, 4), (4, 1))
    assert_size_stride(primals_2, (4, 4), (4, 1))
    assert_size_stride(primals_3, (4, 4), (4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_mul_0[grid(16)](primals_2, buf0, 16, XBLOCK=16,
            num_warps=1, num_stages=1)
        del primals_2
        buf1 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        extern_kernels.addmm(buf0, reinterpret_tensor(primals_1, (4, 4), (4,
            1), 0), reinterpret_tensor(primals_1, (4, 4), (1, 4), 0), alpha
            =1, beta=1, out=buf1)
        buf2 = torch.ops.aten.linalg_inv_ex.default(buf1)
        buf3 = buf2[0]
        del buf2
        buf5 = buf1
        del buf1
        extern_kernels.addmm(buf0, reinterpret_tensor(primals_1, (4, 4), (4,
            1), 16), reinterpret_tensor(primals_1, (4, 4), (1, 4), 16),
            alpha=1, beta=1, out=buf5)
        buf6 = torch.ops.aten.linalg_inv_ex.default(buf5)
        buf7 = buf6[0]
        del buf6
        buf9 = buf5
        del buf5
        extern_kernels.addmm(buf0, reinterpret_tensor(primals_1, (4, 4), (4,
            1), 32), reinterpret_tensor(primals_1, (4, 4), (1, 4), 32),
            alpha=1, beta=1, out=buf9)
        buf10 = torch.ops.aten.linalg_inv_ex.default(buf9)
        buf11 = buf10[0]
        del buf10
        buf13 = buf9
        del buf9
        extern_kernels.addmm(buf0, reinterpret_tensor(primals_1, (4, 4), (4,
            1), 48), reinterpret_tensor(primals_1, (4, 4), (1, 4), 48),
            alpha=1, beta=1, out=buf13)
        del buf0
        buf14 = torch.ops.aten.linalg_inv_ex.default(buf13)
        buf15 = buf14[0]
        del buf14
        buf17 = empty_strided_cuda((4, 4, 4), (16, 1, 4), torch.float32)
        triton_poi_fused_cat_1[grid(64)](buf3, buf7, buf11, buf15, buf17, 
            64, XBLOCK=64, num_warps=1, num_stages=1)
        buf18 = empty_strided_cuda((4, 16), (16, 1), torch.float32)
        extern_kernels.mm(primals_3, reinterpret_tensor(primals_1, (4, 16),
            (1, 4), 0), out=buf18)
        buf19 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused_clone_2[grid(256)](buf17, buf19, 256, XBLOCK=256,
            num_warps=4, num_stages=1)
        buf20 = reinterpret_tensor(buf17, (16, 1, 4), (4, 4, 1), 0)
        del buf17
        extern_kernels.bmm(reinterpret_tensor(buf18, (16, 1, 4), (4, 4, 1),
            0), reinterpret_tensor(buf19, (16, 4, 4), (16, 4, 1), 0), out=buf20
            )
        buf21 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused_clone_3[grid(256)](primals_1, buf21, 256, XBLOCK=
            128, num_warps=4, num_stages=1)
        buf22 = empty_strided_cuda((16, 1, 4), (4, 4, 1), torch.float32)
        extern_kernels.bmm(buf20, reinterpret_tensor(buf21, (16, 4, 4), (16,
            4, 1), 0), out=buf22)
        buf23 = reinterpret_tensor(buf13, (4, 4, 1), (4, 1, 1), 0)
        del buf13
        extern_kernels.bmm(reinterpret_tensor(buf22, (4, 4, 4), (16, 4, 1),
            0), reinterpret_tensor(primals_3, (4, 4, 1), (4, 1, 1), 0), out
            =buf23)
        del buf22
        buf24 = reinterpret_tensor(buf23, (4, 4), (4, 1), 0)
        del buf23
        triton_poi_fused_sqrt_4[grid(16)](buf24, 16, XBLOCK=16, num_warps=1,
            num_stages=1)
    return buf24, primals_3, reinterpret_tensor(primals_1, (4, 4), (4, 1), 0
        ), reinterpret_tensor(primals_1, (4, 4), (4, 1), 16
        ), reinterpret_tensor(primals_1, (4, 4), (4, 1), 32
        ), reinterpret_tensor(primals_1, (4, 4), (4, 1), 48
        ), buf24, reinterpret_tensor(buf20, (16, 4, 1), (4, 1, 4), 0
        ), reinterpret_tensor(buf21, (16, 4, 4), (16, 1, 4), 0
        ), reinterpret_tensor(buf18, (16, 4, 1), (4, 1, 4), 0
        ), reinterpret_tensor(buf19, (16, 4, 4), (16, 1, 4), 0
        ), reinterpret_tensor(buf15, (4, 4), (4, 1), 0), reinterpret_tensor(
        buf11, (4, 4), (4, 1), 0), reinterpret_tensor(buf7, (4, 4), (4, 1), 0
        ), reinterpret_tensor(buf3, (4, 4), (4, 1), 0)


class LinearCapsProNew(nn.Module):

    def __init__(self, in_features, num_C, num_D, eps=0.0001):
        super(LinearCapsProNew, self).__init__()
        self.in_features = in_features
        self.num_C = num_C
        self.num_D = num_D
        self.eps = eps
        self.weight = Parameter(torch.Tensor(num_C * num_D, in_features))
        self.reset_parameters()

    def reset_parameters(self):
        stdv = 1.0 / math.sqrt(self.weight.size(1))
        self.weight.data.uniform_(-stdv, stdv)

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


Measured latency_ms: 0.03399