```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

# Fused CUDA implementation of: mean(clamp(1 - p * t, min=0))
# A single kernel computes the elementwise hinge term on the fly and reduces it
# (warp shuffle + shared memory + one atomicAdd per block into a double
# accumulator for high accuracy), then a tiny finalize kernel divides by N.
hinge_loss_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <limits>

#define BLOCK_SIZE 256

__device__ __forceinline__ double warp_reduce_sum(double val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffffu, val, offset);
    }
    return val;
}

__device__ __forceinline__ double block_reduce_sum(double val) {
    __shared__ double shared[BLOCK_SIZE / 32];
    const int lane = threadIdx.x & 31;
    const int wid  = threadIdx.x >> 5;
    val = warp_reduce_sum(val);
    if (lane == 0) shared[wid] = val;
    __syncthreads();
    val = (threadIdx.x < BLOCK_SIZE / 32) ? shared[lane] : 0.0;
    if (wid == 0) val = warp_reduce_sum(val);
    return val;
}

__global__ void hinge_loss_sum_kernel(const float* __restrict__ p,
                                      const float* __restrict__ t,
                                      double* __restrict__ out,
                                      long long n) {
    long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    const long long stride = (long long)gridDim.x * blockDim.x;
    double acc = 0.0;
    for (; i < n; i += stride) {
        const float v = fmaxf(fmaf(-p[i], t[i], 1.0f), 0.0f);
        acc += (double)v;
    }
    acc = block_reduce_sum(acc);
    if (threadIdx.x == 0) atomicAdd(out, acc);
}

__global__ void hinge_loss_sum_vec4_kernel(const float4* __restrict__ p,
                                           const float4* __restrict__ t,
                                           double* __restrict__ out,
                                           long long n4) {
    long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    const long long stride = (long long)gridDim.x * blockDim.x;
    double acc = 0.0;
    for (; i < n4; i += stride) {
        const float4 pv = p[i];
        const float4 tv = t[i];
        acc += (double)fmaxf(fmaf(-pv.x, tv.x, 1.0f), 0.0f);
        acc += (double)fmaxf(fmaf(-pv.y, tv.y, 1.0f), 0.0f);
        acc += (double)fmaxf(fmaf(-pv.z, tv.z, 1.0f), 0.0f);
        acc += (double)fmaxf(fmaf(-pv.w, tv.w, 1.0f), 0.0f);
    }
    acc = block_reduce_sum(acc);
    if (threadIdx.x == 0) atomicAdd(out, acc);
}

__global__ void hinge_loss_finalize_kernel(const double* __restrict__ sum,
                                           float* __restrict__ out,
                                           long long n) {
    *out = (float)(*sum / (double)n);
}

torch::Tensor hinge_loss_cuda(torch::Tensor p, torch::Tensor t) {
    TORCH_CHECK(p.is_cuda() && t.is_cuda(), "inputs must be CUDA tensors");
    TORCH_CHECK(p.scalar_type() == at::kFloat && t.scalar_type() == at::kFloat,
                "inputs must be float32");
    TORCH_CHECK(p.is_contiguous() && t.is_contiguous(), "inputs must be contiguous");
    TORCH_CHECK(p.numel() == t.numel(), "inputs must have the same number of elements");

    const long long n = (long long)p.numel();
    auto opts = torch::TensorOptions().dtype(torch::kFloat32).device(p.device());
    auto out = torch::empty({1}, opts);
    if (n == 0) {
        out.fill_(std::numeric_limits<float>::quiet_NaN());
        return out.squeeze(0);
    }

    auto sum_buf = torch::zeros({1},
        torch::TensorOptions().dtype(torch::kFloat64).device(p.device()));

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const int sms = at::cuda::getCurrentDeviceProperties()->multiProcessorCount;
    const long long max_blocks = 8LL * sms;

    const float* p_ptr = p.data_ptr<float>();
    const float* t_ptr = t.data_ptr<float>();

    const bool aligned16 = ((reinterpret_cast<uintptr_t>(p_ptr) & 15) == 0) &&
                           ((reinterpret_cast<uintptr_t>(t_ptr) & 15) == 0);

    if ((n & 3) == 0 && aligned16) {
        const long long n4 = n >> 2;
        long long blocks = (n4 + BLOCK_SIZE - 1) / BLOCK_SIZE;
        if (blocks > max_blocks) blocks = max_blocks;
        hinge_loss_sum_vec4_kernel<<<(unsigned int)blocks, BLOCK_SIZE, 0, stream>>>(
            reinterpret_cast<const float4*>(p_ptr),
            reinterpret_cast<const float4*>(t_ptr),
            sum_buf.data_ptr<double>(), n4);
    } else {
        long long blocks = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;
        if (blocks > max_blocks) blocks = max_blocks;
        hinge_loss_sum_kernel<<<(unsigned int)blocks, BLOCK_SIZE, 0, stream>>>(
            p_ptr, t_ptr, sum_buf.data_ptr<double>(), n);
    }

    hinge_loss_finalize_kernel<<<1, 1, 0, stream>>>(
        sum_buf.data_ptr<double>(), out.data_ptr<float>(), n);

    return out.squeeze(0);
}
"""

hinge_loss_cpp_source = "torch::Tensor hinge_loss_cuda(torch::Tensor p, torch::Tensor t);"

# Compile the inline CUDA extension
hinge_loss_ext = load_inline(
    name="hinge_loss_fused_v1",
    cpp_sources=hinge_loss_cpp_source,
    cuda_sources=hinge_loss_cuda_source,
    functions=["hinge_loss_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


class ModelNew(nn.Module):
    """
    Optimized Hinge Loss using a single fused CUDA kernel that computes
    max(0, 1 - predictions * targets) on the fly and reduces it to the mean
    in one pass over the data (no intermediate tensors are materialized).
    """
    def __init__(self):
        super(ModelNew, self).__init__()

    def forward(self, predictions, targets):
        if (predictions.is_cuda and targets.is_cuda
                and predictions.dtype == torch.float32
                and targets.dtype == torch.float32):
            # Match PyTorch broadcasting semantics exactly, then flatten to
            # contiguous buffers (no-op copies when shapes already match).
            p, t = torch.broadcast_tensors(predictions, targets)
            p = p.contiguous()
            t = t.contiguous()
            return hinge_loss_ext.hinge_loss_cuda(p, t)
        # Safe fallback for non-CUDA / non-float32 inputs.
        return torch.mean(torch.clamp(1 - predictions * targets, min=0))
```