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 libdevice, 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_add_bitwise_or_eq_exp_expm1_log_max_mul_sub_sum_where_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
    x2 = xindex
    x1 = xindex // 4
    tmp0 = tl.load(in_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + 4 * x1, xmask, eviction_policy='evict_last')
    tmp2 = tl.load(in_ptr0 + (1 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp4 = tl.load(in_ptr0 + (2 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr0 + (3 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp3 = triton_helpers.maximum(tmp1, tmp2)
    tmp5 = triton_helpers.maximum(tmp3, tmp4)
    tmp7 = triton_helpers.maximum(tmp5, tmp6)
    tmp8 = float('inf')
    tmp9 = tmp7 == tmp8
    tmp10 = float('-inf')
    tmp11 = tmp7 == tmp10
    tmp12 = tmp9 | tmp11
    tmp13 = tmp1 - tmp7
    tmp14 = tl_math.exp(tmp13)
    tmp15 = tmp2 - tmp7
    tmp16 = tl_math.exp(tmp15)
    tmp17 = tmp14 + tmp16
    tmp18 = tmp4 - tmp7
    tmp19 = tl_math.exp(tmp18)
    tmp20 = tmp17 + tmp19
    tmp21 = tmp6 - tmp7
    tmp22 = tl_math.exp(tmp21)
    tmp23 = tmp20 + tmp22
    tmp24 = tl_math.log(tmp23)
    tmp25 = tmp7 + tmp24
    tmp26 = tl.where(tmp12, tmp7, tmp25)
    tmp27 = tmp0 - tmp26
    tmp28 = -3.0
    tmp29 = tmp27 * tmp28
    tmp30 = libdevice.expm1(tmp29)
    tl.store(out_ptr0 + x2, tmp30, xmask)


@triton.jit
def triton_per_fused_div_gather_neg_sum_1(in_out_ptr0, in_ptr0, in_ptr1,
    xnumel, rnumel, XBLOCK: tl.constexpr):
    RBLOCK: tl.constexpr = 4
    xoffset = tl.program_id(0) * XBLOCK
    xoffset + tl.arange(0, XBLOCK)[:, None]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r0 = rindex
    tmp0 = tl.load(in_ptr0 + r0, None)
    tmp1 = tl.full([XBLOCK, RBLOCK], 4, tl.int32)
    tmp2 = tmp0 + tmp1
    tmp3 = tmp0 < 0
    tmp4 = tl.where(tmp3, tmp2, tmp0)
    tl.device_assert((0 <= tmp4) & (tmp4 < 4),
        'index out of bounds: 0 <= tmp4 < 4')
    tmp6 = tl.load(in_ptr1 + (tmp4 + 4 * r0), None, eviction_policy=
        'evict_last')
    tmp7 = -0.08333333333333333
    tmp8 = tmp6 * tmp7
    tmp9 = tl.broadcast_to(tmp8, [XBLOCK, RBLOCK])
    tmp11 = tl.sum(tmp9, 1)[:, None]
    tmp12 = -tmp11
    tmp13 = 0.25
    tmp14 = tmp12 * tmp13
    tl.debug_barrier()
    tl.store(in_out_ptr0 + tl.full([XBLOCK, 1], 0, tl.int32), tmp14, None)


def call(args):
    arg0_1, arg1_1, arg2_1, arg3_1 = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4), (4, 1))
    assert_size_stride(arg1_1, (4, 4), (4, 1))
    assert_size_stride(arg2_1, (4, 4), (4, 1))
    assert_size_stride(arg3_1, (4,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        extern_kernels.addmm(arg1_1, arg0_1, reinterpret_tensor(arg2_1, (4,
            4), (1, 4), 0), alpha=1, beta=1, out=buf0)
        del arg0_1
        del arg1_1
        del arg2_1
        buf1 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_add_bitwise_or_eq_exp_expm1_log_max_mul_sub_sum_where_0[
            grid(16)](buf0, buf1, 16, XBLOCK=16, num_warps=1, num_stages=1)
        del buf0
        buf2 = empty_strided_cuda((), (), torch.float32)
        buf3 = buf2
        del buf2
        triton_per_fused_div_gather_neg_sum_1[grid(1)](buf3, arg3_1, buf1, 
            1, 4, XBLOCK=1, num_warps=2, num_stages=1)
        del arg3_1
        del buf1
    return buf3,


def logsumexp(x, dim=None, keepdim=False):
    if dim is None:
        x, dim = x.view(-1), 0
    xm, _ = torch.max(x, dim, keepdim=True)
    x = torch.where((xm == float('inf')) | (xm == float('-inf')), xm, xm +
        torch.log(torch.sum(torch.exp(x - xm), dim, keepdim=True)))
    return x if keepdim else x.squeeze(dim)


class SplitCrossEntropyLossNew(nn.Module):
    """SplitCrossEntropyLoss calculates an approximate softmax"""

    def __init__(self, hidden_size, q, b, g):
        super(SplitCrossEntropyLossNew, self).__init__()
        self.hidden_size = hidden_size
        self.q = q
        self.b = b
        self.g = g

    def log_b(self, logits):
        lnorm = logsumexp(logits, dim=-1, keepdim=True)
        lsum = torch.exp(self.b * logits).sum(dim=1, keepdim=True) * torch.exp(
            -self.b * lnorm)
        return torch.exp((self.b - 1.0) * (logits - lnorm)) / (self.b - 1.0
            ) - lsum / self.b - 1.0 / ((self.b - 1.0) * self.b), lsum

    def log_g(self, logits):
        lnorm = logsumexp(logits, dim=-1, keepdim=True)
        lsum = torch.exp(self.g * logits).sum(dim=1, keepdim=True)
        return logits - torch.log(lsum) / self.g, lsum * torch.exp(-self.g *
            lnorm)

    def log_q(self, logits):
        lnorm = logsumexp(logits, dim=-1, keepdim=True)
        return torch.expm1((1.0 - self.q) * (logits - lnorm)) / ((1.0 -
            self.q) * self.q), torch.exp(lnorm)

    def forward(self, input_0, input_1, input_2, input_3):
        arg0_1 = input_0
        arg1_1 = input_1
        arg2_1 = input_2
        arg3_1 = input_3
        output = call([arg0_1, arg1_1, arg2_1, arg3_1])
        return output[0]


Measured latency_ms: 0.038352