```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 kernel that computes mean(clamp(1 - predictions * targets, min=0))
# in a single pass over the data (multiply + hinge + sum reduction), followed by
# a tiny scaling kernel that divides by N to produce the mean.
hinge_loss_cpp_source = (
    "torch::Tensor hinge_loss_cuda(torch::Tensor predictions, torch::Tensor targets);"
)

hinge_loss_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>
#include <limits>

#define THREADS 256
#define WARPS (THREADS / 32)

// Warp-shuffle + shared-memory block reduction, then one atomicAdd per block.
__device__ __forceinline__ void atomic_block_reduce(float val, float* out) {
    __shared__ float sdata[WARPS];
    const int lane = threadIdx.x & 31;
    const int wid = threadIdx.x >> 5;
    #pragma unroll
    for (int o = 16; o > 0; o >>= 1) {
        val += __shfl_down_sync(0xffffffffu, val, o);
    }
    if (lane == 0) sdata[wid] = val;
    __syncthreads();
    if (wid == 0) {
        val = (lane < WARPS) ? sdata[lane] : 0.0f;
        #pragma unroll
        for (int o = 16; o > 0; o >>= 1) {
            val += __shfl_down_sync(0xffffffffu, val, o);
        }
        if (lane == 0) atomicAdd(out, val);
    }
}

// Vectorized (float4) fused multiply -> hinge -> partial-sum kernel.
__global__ void hinge_loss_vec_kernel(const float* __restrict__ p,
                                      const float* __restrict__ t,
                                      float* __restrict__ out,
                                      int64_t n) {
    const int64_t n4 = n >> 2;
    const float4* p4 = reinterpret_cast<const float4*>(p);
    const float4* t4 = reinterpret_cast<const float4*>(t);
    float sum = 0.0f;
    const int64_t tid = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    const int64_t stride = (int64_t)gridDim.x * blockDim.x;
    for (int64_t i = tid; i < n4; i += stride) {
        const float4 pv = p4[i];
        const float4 tv = t4[i];
        sum += fmaxf(1.0f - pv.x * tv.x, 0.0f);
        sum += fmaxf(1.0f - pv.y * tv.y, 0.0f);
        sum += fmaxf(1.0f - pv.z * tv.z, 0.0f);
        sum += fmaxf(1.0f - pv.w * tv.w, 0.0f);
    }
    // Handle the (up to 3) tail elements.
    const int64_t rem = n4 << 2;
    if (tid < n - rem) {
        sum += fmaxf(1.0f - p[rem + tid] * t[rem + tid], 0.0f);
    }
    atomic_block_reduce(sum, out);
}

// Scalar fallback for misaligned pointers.
__global__ void hinge_loss_scalar_kernel(const float* __restrict__ p,
                                         const float* __restrict__ t,
                                         float* __restrict__ out,
                                         int64_t n) {
    float sum = 0.0f;
    const int64_t tid = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    const int64_t stride = (int64_t)gridDim.x * blockDim.x;
    for (int64_t i = tid; i < n; i += stride) {
        sum += fmaxf(1.0f - p[i] * t[i], 0.0f);
    }
    atomic_block_reduce(sum, out);
}

__global__ void scale_kernel(float* out, float inv_n) {
    out[0] *= inv_n;
}

torch::Tensor hinge_loss_cuda(torch::Tensor predictions, torch::Tensor targets) {
    TORCH_CHECK(predictions.is_cuda() && targets.is_cuda(), "inputs must be CUDA tensors");
    TORCH_CHECK(predictions.scalar_type() == at::kFloat && targets.scalar_type() == at::kFloat,
                "inputs must be float32");

    auto p = predictions.contiguous();
    auto t = targets.contiguous();
    // General broadcast support (fast path is equal shapes).
    if (p.numel() != t.numel()) {
        auto b = at::broadcast_tensors({p, t});
        p = b[0].contiguous();
        t = b[1].contiguous();
    }

    const int64_t n = p.numel();
    auto out = torch::zeros({}, p.options());
    if (n == 0) {
        out.fill_(std::numeric_limits<float>::quiet_NaN());
        return out;
    }

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

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

    const bool aligned = ((reinterpret_cast<uintptr_t>(p_ptr) & 0xF) == 0) &&
                         ((reinterpret_cast<uintptr_t>(t_ptr) & 0xF) == 0);

    if (aligned) {
        const int64_t n4 = n >> 2;
        const int64_t wanted = (n4 + THREADS - 1) / THREADS;
        const int blocks = (int)std::min(std::max<int64_t>(wanted, (int64_t)1), max_blocks);
        hinge_loss_vec_kernel<<<blocks, THREADS, 0, stream>>>(p_ptr, t_ptr, out_ptr, n);
    } else {
        const int64_t wanted = (n + THREADS - 1) / THREADS;
        const int blocks = (int)std::min(std::max<int64_t>(wanted, (int64_t)1), max_blocks);
        hinge_loss_scalar_kernel<<<blocks, THREADS, 0, stream>>>(p_ptr, t_ptr, out_ptr, n);
    }
    scale_kernel<<<1, 1, 0, stream>>>(out_ptr, 1.0f / (float)n);
    return out;
}
"""

# Compile the fused hinge-loss CUDA kernel.
hinge_loss = load_inline(
    name="hinge_loss_fused",
    cpp_sources=hinge_loss_cpp_source,
    cuda_sources=hinge_loss_cuda_source,
    functions=["hinge_loss_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Hinge loss computed with a single fused custom CUDA kernel:
    multiply + clamp + sum-reduction + mean in one pass over the data.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss

    def forward(self, predictions, targets):
        if (predictions.is_cuda and targets.is_cuda
                and predictions.dtype == torch.float32
                and targets.dtype == torch.float32):
            return self.hinge_loss.hinge_loss_cuda(predictions, targets)
        # Fallback to the reference PyTorch implementation for other cases.
        return torch.mean(torch.clamp(1 - predictions * targets, min=0))
```