import torch
import triton
import triton.language as tl
from torch._inductor.runtime.triton_heuristics import grid, split_scan_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
from torch import Tensor
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_red_fused_max_1(in_ptr0, in_ptr1, out_ptr0, xnumel, rnumel,
    XBLOCK: tl.constexpr, RBLOCK: tl.constexpr):
    xnumel = 2
    rnumel = 8192
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    _tmp9 = tl.full([XBLOCK, RBLOCK], float('-inf'), tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r1 = rindex
        tmp0 = tl.load(in_ptr0 + (128 * x0 + r1 // 64), rmask & xmask,
            eviction_policy='evict_last', other=0.0)
        tmp1 = tmp0.to(tl.int64)
        tmp2 = tl.full([XBLOCK, RBLOCK], 4, tl.int32)
        tmp3 = tmp1 + tmp2
        tmp4 = tmp1 < 0
        tmp5 = tl.where(tmp4, tmp3, tmp1)
        tl.device_assert((0 <= tmp5) & (tmp5 < 4) | ~(rmask & xmask),
            'index out of bounds: 0 <= tmp5 < 4')
        tmp7 = tl.load(in_ptr1 + (64 * tmp5 + r1 % 64), rmask & xmask,
            eviction_policy='evict_last', other=0.0)
        tmp8 = tl.broadcast_to(tmp7, [XBLOCK, RBLOCK])
        tmp10 = triton_helpers.maximum(_tmp9, tmp8)
        _tmp9 = tl.where(rmask & xmask, tmp10, _tmp9)
    tmp9 = triton_helpers.max2(_tmp9, 1)[:, None]
    tl.store(out_ptr0 + x0, tmp9, xmask)


@triton.jit
def triton_per_fused_max_2(in_ptr0, out_ptr0, xnumel, rnumel, XBLOCK: tl.
    constexpr):
    RBLOCK: tl.constexpr = 2
    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.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = triton_helpers.max2(tmp1, 1)[:, None]
    tl.store(out_ptr0 + tl.full([XBLOCK, 1], 0, tl.int32), tmp3, None)


@triton.jit
def _triton_helper_fn_add0(arg0_0, arg1_0):
    tmp0 = arg0_0 + arg1_0
    return tmp0


@triton.jit
def triton_spl_fused_cumsum_exp_sub_3(in_ptr0, in_ptr1, in_ptr2, out_ptr0,
    ws_ptr, xnumel, rnumel, RBLOCK: tl.constexpr):
    XBLOCK: tl.constexpr = 1
    rnumel = 16384
    xoffset = tl.program_id(1) * XBLOCK
    tl.full([1], xoffset, tl.int32)
    tl.full([RBLOCK], True, tl.int1)
    roffset = tl.program_id(0) * RBLOCK
    rindex = roffset + tl.arange(0, RBLOCK)[:]
    rmask = rindex < rnumel
    r0 = rindex
    tmp0 = tl.load(in_ptr0 + r0 // 64, rmask, eviction_policy='evict_last',
        other=0.0)
    tmp8 = tl.load(in_ptr2 + 0)
    tmp9 = tl.broadcast_to(tmp8, [RBLOCK])
    tmp12 = tl.num_programs(0)
    tmp13 = ws_ptr.to(tl.pointer_type(tl.uint64)) + xoffset * 1 * tmp12
    tmp1 = tmp0.to(tl.int64)
    tmp2 = tl.full([RBLOCK], 4, tl.int32)
    tmp3 = tmp1 + tmp2
    tmp4 = tmp1 < 0
    tmp5 = tl.where(tmp4, tmp3, tmp1)
    tl.device_assert((0 <= tmp5) & (tmp5 < 4) | ~rmask,
        'index out of bounds: 0 <= tmp5 < 4')
    tmp7 = tl.load(in_ptr1 + (64 * tmp5 + r0 % 64), rmask, other=0.0)
    tmp10 = tmp7 - tmp9
    tmp11 = tl_math.exp(tmp10)
    tmp14 = tmp11.to(tl.float32)
    tmp15 = tl.broadcast_to(tmp14, [RBLOCK])
    tmp16 = tl.reduce(tmp15, 0, _triton_helper_fn_add0)
    tmp17 = triton_helpers.exclusive_scan_decoupled_lookback(tmp13, tmp16,
        tl.program_id(0), _triton_helper_fn_add0, DTYPE_VALUE_AS_UINT=tl.
        uint32, DTYPE_PACK=tl.uint64)
    tmp18 = tl.associative_scan(tmp15, 0, _triton_helper_fn_add0)
    tmp19 = _triton_helper_fn_add0(tmp17, tmp18)
    tmp20 = tl.where(roffset == 0, tmp18, tmp19)
    tl.store(out_ptr0 + tl.broadcast_to(r0, [RBLOCK]), tmp20, rmask)


@triton.jit
def triton_red_fused_add_log_mul_sub_sum_4(in_ptr0, in_ptr1, in_ptr2,
    in_ptr3, in_ptr4, out_ptr0, out_ptr1, xnumel, rnumel, XBLOCK: tl.
    constexpr, RBLOCK: tl.constexpr):
    xnumel = 2
    rnumel = 8192
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:, None]
    xmask = xindex < xnumel
    rbase = tl.arange(0, RBLOCK)[None, :]
    x0 = xindex
    tmp12 = tl.load(in_ptr3 + 0)
    tmp13 = tl.broadcast_to(tmp12, [XBLOCK, RBLOCK])
    _tmp19 = tl.full([XBLOCK, RBLOCK], 0, tl.float32)
    _tmp22 = tl.full([XBLOCK, RBLOCK], 0, tl.float32)
    for roffset in range(0, rnumel, RBLOCK):
        rindex = roffset + rbase
        rmask = rindex < rnumel
        r1 = rindex
        tmp0 = tl.load(in_ptr0 + (128 * x0 + r1 // 64), rmask & xmask,
            eviction_policy='evict_last', other=0.0)
        tmp8 = tl.load(in_ptr2 + (r1 + 8192 * x0), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp1 = tmp0.to(tl.int64)
        tmp2 = tl.full([XBLOCK, RBLOCK], 4, tl.int32)
        tmp3 = tmp1 + tmp2
        tmp4 = tmp1 < 0
        tmp5 = tl.where(tmp4, tmp3, tmp1)
        tl.device_assert((0 <= tmp5) & (tmp5 < 4) | ~(rmask & xmask),
            'index out of bounds: 0 <= tmp5 < 4')
        tmp7 = tl.load(in_ptr1 + (64 * tmp5 + r1 % 64), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp9 = 1e-07
        tmp10 = tmp8 + tmp9
        tmp11 = tl_math.log(tmp10)
        tmp14 = tmp11 + tmp13
        tmp15 = tmp7 - tmp14
        tmp16 = tl.load(in_ptr4 + (64 * tmp5 + r1 % 64), rmask & xmask,
            eviction_policy='evict_first', other=0.0)
        tmp17 = tmp15 * tmp16
        tmp18 = tl.broadcast_to(tmp17, [XBLOCK, RBLOCK])
        tmp20 = _tmp19 + tmp18
        _tmp19 = tl.where(rmask & xmask, tmp20, _tmp19)
        tmp21 = tl.broadcast_to(tmp16, [XBLOCK, RBLOCK])
        tmp23 = _tmp22 + tmp21
        _tmp22 = tl.where(rmask & xmask, tmp23, _tmp22)
    tmp19 = tl.sum(_tmp19, 1)[:, None]
    tl.store(out_ptr0 + x0, tmp19, xmask)
    tmp22 = tl.sum(_tmp22, 1)[:, None]
    tl.store(out_ptr1 + x0, tmp22, xmask)


@triton.jit
def triton_per_fused_add_div_log_mul_neg_sub_sum_5(in_out_ptr0, in_ptr0,
    in_ptr1, xnumel, rnumel, XBLOCK: tl.constexpr):
    RBLOCK: tl.constexpr = 2
    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)
    tmp4 = tl.load(in_ptr1 + r0, None)
    tmp1 = tl.broadcast_to(tmp0, [XBLOCK, RBLOCK])
    tmp3 = tl.sum(tmp1, 1)[:, None]
    tmp5 = tl.broadcast_to(tmp4, [XBLOCK, RBLOCK])
    tmp7 = tl.sum(tmp5, 1)[:, None]
    tmp8 = tmp3 / tmp7
    tmp9 = -tmp8
    tl.debug_barrier()
    tl.store(in_out_ptr0 + tl.full([XBLOCK, 1], 0, tl.int32), tmp9, 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)](arg1_1, buf1, 64, 4, XBLOCK=8,
            num_warps=2, num_stages=1)
        del arg1_1
        buf2 = empty_strided_cuda((2,), (1,), torch.float32)
        triton_red_fused_max_1[grid(2)](buf1, arg0_1, buf2, 2, 8192, XBLOCK
            =1, RBLOCK=2048, num_warps=16, num_stages=1)
        buf3 = empty_strided_cuda((), (), torch.float32)
        triton_per_fused_max_2[grid(1)](buf2, buf3, 1, 2, XBLOCK=1,
            num_warps=2, num_stages=1)
        buf4 = empty_strided_cuda((16384,), (1,), torch.float32)
        workspace = empty_strided_cuda((512,), (1,), torch.uint8)
        workspace.zero_()
        triton_spl_fused_cumsum_exp_sub_3[split_scan_grid(1, 16384)](buf1,
            arg0_1, buf3, buf4, workspace, 1, 16384, RBLOCK=2048, num_warps
            =16, num_stages=1)
        del workspace
        buf5 = buf2
        del buf2
        buf7 = empty_strided_cuda((2,), (1,), torch.float32)
        triton_red_fused_add_log_mul_sub_sum_4[grid(2)](buf1, arg0_1, buf4,
            buf3, arg2_1, buf5, buf7, 2, 8192, XBLOCK=1, RBLOCK=2048,
            num_warps=16, num_stages=1)
        del arg0_1
        del arg2_1
        del buf1
        del buf4
        buf6 = buf3
        del buf3
        buf9 = buf6
        del buf6
        triton_per_fused_add_div_log_mul_neg_sub_sum_5[grid(1)](buf9, buf5,
            buf7, 1, 2, XBLOCK=1, num_warps=2, num_stages=1)
        del buf5
        del buf7
    return buf9,


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.
    """
    if events.dtype is torch.bool:
        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)
    return -log_h.sub(log_cumsum_h).mul(events).sum().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.015133