import torch
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 import Tensor
from torch import nn as nn
assert_size_stride = torch._C._dynamo.guards.assert_size_stride
empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda


@triton.jit
def triton_per_fused_sort_0(in_ptr0, out_ptr0, xnumel, rnumel, XBLOCK: tl.
    constexpr):
    xnumel = 64
    RBLOCK: tl.constexpr = 4
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rindex = tl.arange(0, RBLOCK)[None, :]
    tl.full([XBLOCK, RBLOCK], True, tl.int1)
    r1 = rindex
    x0 = xindex
    tmp0 = tl.load(in_ptr0 + (r1 + 4 * x0), xmask, other=0.0)
    tmp1 = r1
    tmp2 = tmp1.to(tl.int16)
    tmp3 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp4 = tl.broadcast_to(tmp2, [XBLOCK, RBLOCK])
    _tmp5, tmp6 = triton_helpers.sort_with_index(tmp3, tmp4, None, 1,
        stable=False, descending=True)
    tl.store(out_ptr0 + (r1 + 4 * x0), tmp6, xmask)


@triton.jit
def triton_poi_fused_index_1(in_ptr0, in_ptr1, in_ptr2, out_ptr0, out_ptr1,
    xnumel, XBLOCK: tl.constexpr):
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    tl.full([XBLOCK], True, tl.int1)
    x1 = xindex // 64
    x0 = xindex % 64
    x2 = xindex
    tmp0 = tl.load(in_ptr0 + x1, None, eviction_policy='evict_last')
    tmp1 = tmp0.to(tl.int64)
    tmp2 = tl.full([XBLOCK], 4, tl.int32)
    tmp3 = tmp1 + tmp2
    tmp4 = tmp1 < 0
    tmp5 = tl.where(tmp4, tmp3, tmp1)
    tl.device_assert((0 <= tmp5) & (tmp5 < 4),
        'index out of bounds: 0 <= tmp5 < 4')
    tmp7 = tl.load(in_ptr1 + (x0 + 64 * tmp5), None)
    tmp8 = tl.load(in_ptr2 + (x0 + 64 * tmp5), None)
    tl.store(out_ptr0 + x2, tmp7, None)
    tl.store(out_ptr1 + x2, tmp8, None)


def call(args):
    arg0_1, arg1_1, arg2_1 = args
    args.clear()
    assert_size_stride(arg0_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(arg1_1, (4, 4, 4, 4), (64, 16, 4, 1))
    assert_size_stride(arg2_1, (4, 4, 4, 4), (64, 16, 4, 1))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf1 = empty_strided_cuda((4, 4, 4, 4), (64, 16, 4, 1), torch.int16)
        get_raw_stream(0)
        triton_per_fused_sort_0[grid(64)](arg0_1, buf1, 64, 4, XBLOCK=1,
            num_warps=2, num_stages=1)
        del arg0_1
        buf2 = empty_strided_cuda((4, 4, 4, 4, 4, 4, 4), (4096, 1024, 256, 
            64, 16, 4, 1), torch.float32)
        buf3 = empty_strided_cuda((4, 4, 4, 4, 4, 4, 4), (4096, 1024, 256, 
            64, 16, 4, 1), torch.float32)
        triton_poi_fused_index_1[grid(16384)](buf1, arg2_1, arg1_1, buf2,
            buf3, 16384, XBLOCK=128, num_warps=4, num_stages=1)
        del arg1_1
        del arg2_1
        del buf1
    return buf2, buf3


def cox_ph_loss_sorted(log_h: 'Tensor', events: 'Tensor', eps: 'float'=1e-07
    ) ->Tensor:
    """Requires the input to be sorted by descending duration time.
    See DatasetDurationSorted.
    We calculate the negative log of $(rac{h_i}{\\sum_{j \\in R_i} h_j})^d$,
    where h = exp(log_h) are the hazards and R is the risk set, and d is event.
    We just compute a cumulative sum, and not the true Risk sets. This is a
    limitation, but simple and fast.
    """
    events = events.float()
    events = events.view(-1)
    log_h = log_h.view(-1)
    gamma = log_h.max()
    log_cumsum_h = log_h.sub(gamma).exp().cumsum(0).add(eps).log().add(gamma)
    loss = -log_h.sub(log_cumsum_h).mul(events).sum()
    if events.sum().item() == 0:
        return loss
    return loss.div(events.sum())


def cox_ph_loss(log_h: 'Tensor', durations: 'Tensor', events: 'Tensor', eps:
    'float'=1e-07) ->Tensor:
    """Loss for CoxPH model. If data is sorted by descending duration, see `cox_ph_loss_sorted`.
    We calculate the negative log of $(rac{h_i}{\\sum_{j \\in R_i} h_j})^d$,
    where h = exp(log_h) are the hazards and R is the risk set, and d is event.
    We just compute a cumulative sum, and not the true Risk sets. This is a
    limitation, but simple and fast.
    """
    idx = durations.sort(descending=True)[1]
    events = events[idx]
    log_h = log_h[idx]
    return cox_ph_loss_sorted(log_h, events, eps)


class CoxPHLossNew(torch.nn.Module):
    """Loss for CoxPH model. If data is sorted by descending duration, see `cox_ph_loss_sorted`.
    We calculate the negative log of $(rac{h_i}{\\sum_{j \\in R_i} h_j})^d$,
    where h = exp(log_h) are the hazards and R is the risk set, and d is event.
    We just compute a cumulative sum, and not the true Risk sets. This is a
    limitation, but simple and fast.
    """

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


Measured latency_ms: 0.030359