```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 mean(max(0, 1 - predictions * targets)) in a
# single pass using vectorized float4 loads, warp-shuffle block reductions,
# and atomic accumulation into a scalar.
hinge_loss_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>

#define BLOCK_SIZE 256

__global__ void hinge_loss_kernel(const float* __restrict__ p,
                                  const float* __restrict__ t,
                                  float* __restrict__ out,
                                  long long N) {
    const long long N4 = N >> 2;
    const float4* p4 = reinterpret_cast<const float4*>(p);
    const float4* t4 = reinterpret_cast<const float4*>(t);

    float acc = 0.0f;
    const long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    const long long stride = (long long)gridDim.x * blockDim.x;

    // Vectorized main loop (4 elements per load).
    for (long long i = idx; i < N4; i += stride) {
        float4 pv = p4[i];
        float4 tv = t4[i];
        acc += fmaxf(1.0f - pv.x * tv.x, 0.0f);
        acc += fmaxf(1.0f - pv.y * tv.y, 0.0f);
        acc += fmaxf(1.0f - pv.z * tv.z, 0.0f);
        acc += fmaxf(1.0f - pv.w * tv.w, 0.0f);
    }
    // Remainder elements (N not divisible by 4).
    for (long long i = (N4 << 2) + idx; i < N; i += stride) {
        acc += fmaxf(1.0f - p[i] * t[i], 0.0f);
    }

    // Warp-level reduction.
    for (int offset = 16; offset > 0; offset >>= 1) {
        acc += __shfl_down_sync(0xffffffff, acc, offset);
    }

    __shared__ float warp_sums[BLOCK_SIZE / 32];
    const int lane = threadIdx.x & 31;
    const int wid  = threadIdx.x >> 5;
    if (lane == 0) {
        warp_sums[wid] = acc;
    }
    __syncthreads();

    if (wid == 0) {
        float v = (threadIdx.x < (BLOCK_SIZE / 32)) ? warp_sums[lane] : 0.0f;
        for (int offset = 16; offset > 0; offset >>= 1) {
            v += __shfl_down_sync(0xffffffff, v, offset);
        }
        if (threadIdx.x == 0) {
            atomicAdd(out, v);
        }
    }
}

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 = p.numel();
    auto out = torch::zeros({}, p.options());
    if (N == 0) {
        return out;
    }

    auto stream = at::cuda::getCurrentCUDAStream();
    const long long N4 = (N + 3) >> 2;
    const long long max_blocks =
        (long long)at::cuda::getCurrentDeviceProperties()->multiProcessorCount * 16;
    long long blocks = (N4 + BLOCK_SIZE - 1) / BLOCK_SIZE;
    blocks = std::min(blocks, max_blocks);
    if (blocks < 1) {
        blocks = 1;
    }

    hinge_loss_kernel<<<(int)blocks, BLOCK_SIZE, 0, stream>>>(
        p.data_ptr<float>(), t.data_ptr<float>(), out.data_ptr<float>(), N);

    out.div_(static_cast<double>(N));
    return out;
}
"""

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

# Compile the fused hinge loss CUDA kernel.
hinge_loss_ext = 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):
    """
    Optimized Hinge Loss using a single fused CUDA kernel that computes
    mean(clamp(1 - predictions * targets, min=0)) with one reduction pass.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss_ext

    def forward(self, predictions, targets):
        p = predictions.contiguous()
        t = targets.contiguous()
        if p.numel() != t.numel():
            # Match PyTorch broadcasting semantics for the elementwise product.
            t = torch.broadcast_to(t, p.shape).contiguous()
        return self.hinge_loss.hinge_loss_cuda(p, t)
```