```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: computes max(0, 1 - p*t) elementwise and reduces to the
# mean in a single pass over the data (two kernel launches total, no
# intermediate tensors). Uses float4 vectorized loads plus a grid-stride loop,
# block-level shared-memory reductions, and exact non-FMA rounding for the
# per-element computation to match PyTorch's elementwise semantics closely.
hinge_loss_cuda_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <climits>

#define BLOCK_SIZE 256

__global__ void hinge_loss_partial_kernel(const float* __restrict__ preds,
                                          const float* __restrict__ targets,
                                          float* __restrict__ partial,
                                          int n, int n4) {
    __shared__ float shared[BLOCK_SIZE];
    const int tid = threadIdx.x;
    float sum = 0.0f;

    const float4* p4 = reinterpret_cast<const float4*>(preds);
    const float4* t4 = reinterpret_cast<const float4*>(targets);
    const int stride = gridDim.x * blockDim.x;

    for (int i = blockIdx.x * blockDim.x + tid; i < n4; i += stride) {
        float4 p = p4[i];
        float4 t = t4[i];
        sum += fmaxf(__fsub_rn(1.0f, __fmul_rn(p.x, t.x)), 0.0f);
        sum += fmaxf(__fsub_rn(1.0f, __fmul_rn(p.y, t.y)), 0.0f);
        sum += fmaxf(__fsub_rn(1.0f, __fmul_rn(p.z, t.z)), 0.0f);
        sum += fmaxf(__fsub_rn(1.0f, __fmul_rn(p.w, t.w)), 0.0f);
    }

    // Handle the tail (n % 4 elements) with block 0.
    if (blockIdx.x == 0) {
        for (int i = (n4 << 2) + tid; i < n; i += blockDim.x) {
            sum += fmaxf(__fsub_rn(1.0f, __fmul_rn(preds[i], targets[i])), 0.0f);
        }
    }

    shared[tid] = sum;
    __syncthreads();
    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (tid < s) shared[tid] += shared[tid + s];
        __syncthreads();
    }
    if (tid == 0) partial[blockIdx.x] = shared[0];
}

__global__ void hinge_loss_reduce_kernel(const float* __restrict__ partial,
                                         float* __restrict__ out,
                                         int num_partials, float inv_n) {
    __shared__ float shared[BLOCK_SIZE];
    const int tid = threadIdx.x;
    float sum = 0.0f;
    for (int i = tid; i < num_partials; i += blockDim.x) {
        sum += partial[i];
    }
    shared[tid] = sum;
    __syncthreads();
    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (tid < s) shared[tid] += shared[tid + s];
        __syncthreads();
    }
    if (tid == 0) out[0] = shared[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");
    TORCH_CHECK(predictions.numel() == targets.numel(), "inputs must have the same number of elements");

    const int64_t n = predictions.numel();
    TORCH_CHECK(n > 0, "input must be non-empty");
    TORCH_CHECK(n <= (int64_t)INT_MAX, "input too large");

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    const int64_t n4 = n >> 2;
    const int64_t blocks_needed = (n4 + BLOCK_SIZE - 1) / BLOCK_SIZE;
    const int num_blocks = (int)std::min<int64_t>(std::max<int64_t>(blocks_needed, 1), 4096);

    auto partial = torch::empty({num_blocks}, predictions.options());
    auto out = torch::empty({}, predictions.options());

    hinge_loss_partial_kernel<<<num_blocks, BLOCK_SIZE, 0, stream>>>(
        predictions.data_ptr<float>(),
        targets.data_ptr<float>(),
        partial.data_ptr<float>(),
        (int)n,
        (int)n4);

    hinge_loss_reduce_kernel<<<1, BLOCK_SIZE, 0, stream>>>(
        partial.data_ptr<float>(),
        out.data_ptr<float>(),
        num_blocks,
        1.0f / (float)n);

    return out;
}
"""

hinge_loss_cpp_source = (
    "torch::Tensor hinge_loss_cuda(torch::Tensor predictions, torch::Tensor targets);"
)

# Compile the inline CUDA extension
hinge_loss_module = 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_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized Hinge Loss using a single fused CUDA kernel
    (elementwise max(0, 1 - p*t) + mean reduction in one pass).
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss_module

    def forward(self, predictions, targets):
        if (predictions.is_cuda and targets.is_cuda
                and predictions.dtype == torch.float32
                and targets.dtype == torch.float32):
            if predictions.shape != targets.shape:
                predictions, targets = torch.broadcast_tensors(predictions, targets)
            p = predictions.contiguous()
            t = targets.contiguous()
            if p.numel() > 0:
                return self.hinge_loss.hinge_loss_cuda(p, t)
        # Fallback for non-CUDA / non-float32 / empty inputs
        return torch.mean(torch.clamp(1 - predictions * targets, min=0))
```