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


@triton.jit
def triton_poi_fused_div_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')
    tmp3 = tl.load(in_ptr0 + (1 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp6 = tl.load(in_ptr0 + (2 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp9 = tl.load(in_ptr0 + (3 + 4 * x1), xmask, eviction_policy='evict_last')
    tmp2 = tmp1 * tmp1
    tmp4 = tmp3 * tmp3
    tmp5 = tmp2 + tmp4
    tmp7 = tmp6 * tmp6
    tmp8 = tmp5 + tmp7
    tmp10 = tmp9 * tmp9
    tmp11 = tmp8 + tmp10
    tmp12 = libdevice.sqrt(tmp11)
    tmp13 = 1e-12
    tmp14 = triton_helpers.maximum(tmp12, tmp13)
    tmp15 = tmp0 / tmp14
    tl.store(out_ptr0 + x2, tmp15, xmask)


@triton.jit
def triton_poi_fused_mul_1(in_out_ptr0, in_ptr0, xnumel, XBLOCK: tl.constexpr):
    xnumel = 16
    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.load(in_ptr0 + 0)
    tmp2 = tl.broadcast_to(tmp1, [XBLOCK])
    tmp3 = tmp0 * tmp2
    tl.store(in_out_ptr0 + x0, tmp3, xmask)


def call(args):
    primals_1, primals_2, primals_3 = args
    args.clear()
    assert_size_stride(primals_1, (4, 4), (4, 1))
    assert_size_stride(primals_2, (4, 4), (4, 1))
    assert_size_stride(primals_3, (1,), (1,))
    with torch.cuda._DeviceGuard(0):
        torch.cuda.set_device(0)
        buf0 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        get_raw_stream(0)
        triton_poi_fused_div_0[grid(16)](primals_1, buf0, 16, XBLOCK=16,
            num_warps=1, num_stages=1)
        del primals_1
        buf1 = empty_strided_cuda((4, 4), (1, 4), torch.float32)
        triton_poi_fused_div_0[grid(16)](primals_2, buf1, 16, XBLOCK=16,
            num_warps=1, num_stages=1)
        buf2 = empty_strided_cuda((4, 4), (4, 1), torch.float32)
        extern_kernels.mm(buf0, buf1, out=buf2)
        del buf1
        buf3 = buf2
        del buf2
        triton_poi_fused_mul_1[grid(16)](buf3, primals_3, 16, XBLOCK=16,
            num_warps=1, num_stages=1)
    return buf3, primals_2, primals_3, buf0


def cosine_fully_connected_layer(x_in, weight, scale=None, bias=None,
    normalize_x=True, normalize_w=True):
    assert x_in.dim() == 2
    assert weight.dim() == 2
    assert x_in.size(1) == weight.size(0)
    if normalize_x:
        x_in = F.normalize(x_in, p=2, dim=1, eps=1e-12)
    if normalize_w:
        weight = F.normalize(weight, p=2, dim=0, eps=1e-12)
    x_out = torch.mm(x_in, weight)
    if scale is not None:
        x_out = x_out * scale.view(1, -1)
    if bias is not None:
        x_out = x_out + bias.view(1, -1)
    return x_out


class CosineClassifierNew(nn.Module):

    def __init__(self, num_channels, num_classes, scale=1.0, learn_scale=
        False, bias=False, normalize_x=True, normalize_w=True):
        super().__init__()
        self.num_channels = num_channels
        self.num_classes = num_classes
        self.normalize_x = normalize_x
        self.normalize_w = normalize_w
        weight = torch.FloatTensor(num_classes, num_channels).normal_(0.0,
            np.sqrt(2.0 / num_channels))
        self.weight = nn.Parameter(weight, requires_grad=True)
        if bias:
            bias = torch.FloatTensor(num_classes).fill_(0.0)
            self.bias = nn.Parameter(bias, requires_grad=True)
        else:
            self.bias = None
        scale_cls = torch.FloatTensor(1).fill_(scale)
        self.scale_cls = nn.Parameter(scale_cls, requires_grad=learn_scale)

    def extra_repr(self):
        s = ('num_channels={}, num_classes={}, scale_cls={} (learnable={})'
            .format(self.num_channels, self.num_classes, self.scale_cls.
            item(), self.scale_cls.requires_grad))
        learnable = self.scale_cls.requires_grad
        s = (
            f'num_channels={self.num_channels}, num_classes={self.num_classes}, scale_cls={self.scale_cls.item()} (learnable={learnable}), normalize_x={self.normalize_x}, normalize_w={self.normalize_w}'
            )
        if self.bias is None:
            s += ', bias=False'
        return s

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


Measured latency_ms: 0.043764