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_relu_threshold_backward_0(in_out_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
    tmp0 = tl.load(in_out_ptr0 + x0, xmask)
    tmp1 = tl.full([1], 0, tl.int32)
    tmp2 = triton_helpers.maximum(tmp1, tmp0)
    tmp3 = 0.0
    tmp4 = tmp2 <= tmp3
    tl.store(in_out_ptr0 + x0, tmp2, xmask)
    tl.store(out_ptr0 + x0, tmp4, xmask)


@triton.jit
def triton_poi_fused__log_softmax_1(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
    x3 = xindex
    x0 = xindex % 16
    x2 = xindex // 64
    tmp0 = tl.load(in_ptr0 + x3, xmask)
    tmp1 = tl.load(in_ptr0 + (x0 + 64 * x2), xmask, eviction_policy=
        'evict_last')
    tmp2 = tl.load(in_ptr0 + (16 + x0 + 64 * x2), xmask, eviction_policy=
        'evict_last')
    tmp4 = tl.load(in_ptr0 + (32 + x0 + 64 * x2), xmask, eviction_policy=
        'evict_last')
    tmp6 = tl.load(in_ptr0 + (48 + x0 + 64 * x2), xmask, eviction_policy=
        'evict_last')
    tmp3 = triton_helpers.maximum(tmp1, tmp2)
    tmp5 = triton_helpers.maximum(tmp3, tmp4)
    tmp7 = triton_helpers.maximum(tmp5, tmp6)
    tmp8 = tmp0 - tmp7
    tl.store(out_ptr0 + x3, tmp8, xmask)


@triton.jit
def triton_per_fused__log_softmax_div_mul_neg_sum_2(in_out_ptr0, in_ptr0,
    in_ptr1, xnumel, rnumel):
    XBLOCK: tl.constexpr = 1
    RBLOCK: tl.constexpr = 256
    xoffset = tl.program_id(0) * XBLOCK
    tl.full([1], xoffset, tl.int32)
    tl.full([RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[:]
    tl.full([RBLOCK], True, tl.int1)
    r3 = rindex
    r0 = rindex % 16
    r2 = rindex // 64
    tmp0 = tl.load(in_ptr0 + r3, None)
    tmp1 = tl.load(in_ptr0 + (r0 + 64 * r2), None, eviction_policy='evict_last'
        )
    tmp3 = tl.load(in_ptr0 + (16 + r0 + 64 * r2), None, eviction_policy=
        'evict_last')
    tmp6 = tl.load(in_ptr0 + (32 + r0 + 64 * r2), None, eviction_policy=
        'evict_last')
    tmp9 = tl.load(in_ptr0 + (48 + r0 + 64 * r2), None, eviction_policy=
        'evict_last')
    tmp14 = tl.load(in_ptr1 + r3, None)
    tmp2 = tl_math.exp(tmp1)
    tmp4 = tl_math.exp(tmp3)
    tmp5 = tmp2 + tmp4
    tmp7 = tl_math.exp(tmp6)
    tmp8 = tmp5 + tmp7
    tmp10 = tl_math.exp(tmp9)
    tmp11 = tmp8 + tmp10
    tmp12 = tl_math.log(tmp11)
    tmp13 = tmp0 - tmp12
    tmp15 = tmp13 * tmp14
    tmp16 = tl.broadcast_to(tmp15, [RBLOCK])
    tmp18 = triton_helpers.promote_to_tensor(tl.sum(tmp16, 0))
    tmp19 = -tmp18
    tmp20 = 0.015625
    tmp21 = tmp19 * tmp20
    tl.debug_barrier()
    tl.store(in_out_ptr0 + tl.full([1], 0, tl.int32), tmp21, None)


def call(args):
    primals_1, primals_2, primals_3, primals_4, primals_5 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (4, 1))
    assert_size_stride(primals_2, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(primals_3, (4, 4), (4, 1))
    assert_size_stride(primals_4, (4, 4), (4, 1))
    assert_size_stride(primals_5, (4, 4, 4, 4), (64, 16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((64, 4), (4, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(primals_2, (64, 4), (4, 1), 0),
            reinterpret_tensor(primals_1, (4, 4), (1, 4), 0), out=buf0)
        del primals_1
        buf1 = reinterpret_tensor(buf0, (4, 4, 4, 4), (64, 16, 4, 1), 0)
        del buf0
        buf6 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.bool)
        get_raw_stream(0)
        triton_poi_fused_relu_threshold_backward_0[grid(256)](buf1, buf6, 
            256, XBLOCK=128, num_warps=4, num_stages=1)
        buf2 = empty_strided_cuda((64, 4), (4, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(buf1, (64, 4), (4, 1), 0),
            reinterpret_tensor(primals_3, (4, 4), (1, 4), 0), out=buf2)
        buf3 = empty_strided_cuda((64, 4), (4, 1), torch.float32)
        extern_kernels.mm(buf2, reinterpret_tensor(primals_4, (4, 4), (1, 4
            ), 0), out=buf3)
        buf4 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.float32)
        triton_poi_fused__log_softmax_1[grid(256)](buf3, buf4, 256, XBLOCK=
            256, num_warps=4, num_stages=1)
        buf5 = empty_strided_cuda((), (), torch.float32)
        buf7 = buf5
        del buf5
        triton_per_fused__log_softmax_div_mul_neg_sum_2[grid(1)](buf7, buf4,
            primals_5, 1, 256, num_warps=2, num_stages=1)
        del buf4
    return buf7, reinterpret_tensor(buf3, (4, 4, 4, 4), (64, 16, 4, 1), 0
        ), primals_5, reinterpret_tensor(primals_2, (64, 4), (4, 1), 0
        ), reinterpret_tensor(buf1, (64, 4), (4, 1), 0
        ), buf2, reinterpret_tensor(buf3, (4, 4, 4, 4), (64, 16, 4, 1), 0
        ), primals_4, primals_3, buf6


class FFN(nn.Module):
    """Feed Forward Network."""

    def __init__(self, num_features: 'int', ffn_dim_1: 'int', ffn_dim_2: 'int'
        ) ->None:
        """Initialize the class."""
        super().__init__()
        self.gemm1 = nn.Linear(num_features, ffn_dim_1, bias=False)
        self.relu = nn.ReLU()
        self.gemm2 = nn.Linear(ffn_dim_1, ffn_dim_2, bias=False)

    def forward(self, x: 'torch.Tensor') ->torch.Tensor:
        """Step forward."""
        out = self.gemm1(x)
        out = self.relu(out)
        out = self.gemm2(out)
        return out


class LogReg(nn.Module):
    """Logreg class."""

    def __init__(self, num_features: 'int', num_classes: 'int'):
        """Initialize the class."""
        super().__init__()
        self.lin_layer = nn.Linear(in_features=num_features, out_features=
            num_classes, bias=False)
        self.criterion = nn.CrossEntropyLoss()

    def forward(self, inputs: 'torch.Tensor', targets: 'torch.Tensor'
        ) ->torch.Tensor:
        """Step forward."""
        out = self.lin_layer(inputs)
        loss = self.criterion(out, targets)
        return loss, out


class FFNLogRegNew(nn.Module):
    """Feed Forward Network + LogReg."""

    def __init__(self, num_features: 'int', ffn_embedding_size: 'int',
        embedding_size: 'int', num_classes: 'int') ->None:
        """Initialize the class."""
        super().__init__()
        self.ffn = FFN(num_features, ffn_embedding_size, embedding_size)
        self.logreg = LogReg(embedding_size, num_classes)
        self._init_params()

    def _init_params(self) ->None:
        for p in self.parameters():
            nn.init.xavier_normal_(p)

    @staticmethod
    def get_fake_inputs(args):
        """
        Get fake inputs.

        The size of the inputs are required for the SambaNova compiler.

        Args:
            args: CLI arguments.

        Outputs:
            X_randn: A Samba tensor representing the correct shape for model inputs.
            Y_randint: A Samba tensor representing the correct shape for model outputs.
        """
        X_randn = samba.randn(args.batch_size, args.num_features, name=
            'image', batch_dim=0).bfloat16().float()
        low_inclusive = 0
        high_exclusive = args.num_classes
        size = args.batch_size,
        Y_randint = samba.randint(low_inclusive, high_exclusive, size, name
            ='label', batch_dim=0)
        return X_randn, Y_randint

    def forward(self, input_0, input_1):
        primals_1 = self.ffn.gemm1.weight
        primals_3 = self.ffn.gemm2.weight
        primals_4 = self.logreg.lin_layer.weight
        primals_2 = input_0
        primals_5 = input_1
        output = call([primals_1, primals_2, primals_3, primals_4, primals_5])
        return output[0], output[1]


Measured latency_ms: 0.037273