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 import nn
import torch.nn
import torch.optim
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_cat_0(in_ptr0, out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 80
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x0 = xindex % 5
    x3 = xindex // 5
    x1 = xindex // 5 % 4
    x4 = xindex
    tmp0 = x0
    tl.full([1], 0, tl.int64)
    tmp3 = tl.full([1], 4, tl.int64)
    tmp4 = tmp0 < tmp3
    tmp5 = tl.load(in_ptr0 + (4 * x3 + x0), tmp4 & xmask, eviction_policy=
        'evict_last', other=0.0)
    tmp6 = tmp0 >= tmp3
    tl.full([1], 5, tl.int64)
    tmp9 = x1
    tmp10 = tmp9.to(tl.float32)
    tmp11 = tl.full(tmp10.shape, 0.0, tmp10.dtype)
    tmp12 = tl.where(tmp6, tmp10, tmp11)
    tmp13 = tl.where(tmp4, tmp5, tmp12)
    tl.store(out_ptr0 + x4, tmp13, xmask)


@triton.jit
def triton_poi_fused_relu_threshold_backward_1(in_out_ptr0, in_ptr0,
    out_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 64
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = xindex < xnumel
    x2 = xindex
    x0 = xindex % 4
    tmp0 = tl.load(in_out_ptr0 + x2, xmask)
    tmp1 = tl.load(in_ptr0 + x0, xmask, eviction_policy='evict_last')
    tmp2 = tmp0 + tmp1
    tmp3 = tl.full([1], 0, tl.int32)
    tmp4 = triton_helpers.maximum(tmp3, tmp2)
    tmp5 = 0.0
    tmp6 = tmp4 <= tmp5
    tl.store(in_out_ptr0 + x2, tmp4, xmask)
    tl.store(out_ptr0 + x2, tmp6, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4, 4), (16, 4, 1))
    assert_size_stride(primals_2, (4, 5), (5, 1))
    assert_size_stride(primals_3, (4,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4, 5), (20, 5, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_cat_0[grid(80)](primals_1, buf0, 80, XBLOCK=128,
            num_warps=4, num_stages=1)
        del primals_1
        buf1 = empty_strided_cuda((16, 4), (4, 1), torch.float32)
        extern_kernels.mm(reinterpret_tensor(buf0, (16, 5), (5, 1), 0),
            reinterpret_tensor(primals_2, (5, 4), (1, 5), 0), out=buf1)
        del primals_2
        buf2 = reinterpret_tensor(buf1, (4, 4, 4), (16, 4, 1), 0)
        del buf1
        buf3 = empty_strided_cuda((4, 4, 4), (16, 4, 1), torch.bool)
        triton_poi_fused_relu_threshold_backward_1[grid(64)](buf2,
            primals_3, buf3, 64, XBLOCK=64, num_warps=1, num_stages=1)
        del primals_3
    return buf2, reinterpret_tensor(buf0, (16, 5), (5, 1), 0), buf3


class PositionalEncodingNew(nn.Module):
    """
    A special, non-learnable positional encoding for handling variable (possibly longer)
    lengths of inputs. We simply add an ordinal number as an additional dimension for
    the input embeddings, and then project them back to the original number of dimensions
    """

    def __init__(self, dim_model):
        super().__init__()
        self.pos_embed = nn.Linear(dim_model + 1, dim_model)
        self.activation = nn.ReLU()

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


Measured latency_ms: 0.036141