```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) element-wise and performs the
# mean reduction in a single kernel launch (block-level warp-shuffle reduction
# + atomic accumulation into the scalar output).
hinge_loss_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

__global__ void hinge_loss_kernel(const float* __restrict__ preds,
                                  const float* __restrict__ targets,
                                  float* __restrict__ out,
                                  int64_t n,
                                  float inv_n) {
    float local = 0.0f;
    int64_t stride = (int64_t)blockDim.x * gridDim.x;
    for (int64_t i = (int64_t)blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride) {
        float v = 1.0f - preds[i] * targets[i];
        local += fmaxf(v, 0.0f);
    }

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

    __shared__ float shared[32];
    int lane = threadIdx.x & 31;
    int wid = threadIdx.x >> 5;
    if (lane == 0) {
        shared[wid] = local;
    }
    __syncthreads();

    if (wid == 0) {
        int num_warps = (blockDim.x + 31) >> 5;
        local = (lane < num_warps) ? shared[lane] : 0.0f;
        for (int offset = 16; offset > 0; offset >>= 1) {
            local += __shfl_down_sync(0xffffffffu, local, offset);
        }
        if (lane == 0) {
            atomicAdd(out, local * inv_n);
        }
    }
}

torch::Tensor hinge_loss_cuda(torch::Tensor predictions, torch::Tensor targets) {
    TORCH_CHECK(predictions.is_cuda(), "predictions must be a CUDA tensor");
    TORCH_CHECK(targets.is_cuda(), "targets must be a CUDA tensor");
    TORCH_CHECK(predictions.scalar_type() == torch::kFloat32, "predictions must be float32");
    TORCH_CHECK(targets.scalar_type() == torch::kFloat32, "targets must be float32");

    auto p = predictions.contiguous();
    auto t = targets.contiguous();
    int64_t n = (int64_t)p.numel();

    auto out = torch::zeros({}, p.options());
    if (n == 0) {
        return out;
    }

    const int block_size = 256;
    int64_t grid = (n + block_size - 1) / block_size;
    if (grid > 4096) {
        grid = 4096;  // grid-stride loop handles the rest
    }
    float inv_n = 1.0f / (float)n;

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    hinge_loss_kernel<<<(int)grid, block_size, 0, stream>>>(
        p.data_ptr<float>(), t.data_ptr<float>(), out.data_ptr<float>(), n, inv_n);

    return out;
}
"""

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

# Compile the inline CUDA fused hinge-loss kernel
hinge_loss = load_inline(
    name="hinge_loss_fused",
    cpp_sources=hinge_loss_cpp_source,
    cuda_sources=hinge_loss_source,
    functions=["hinge_loss_cuda"],
    verbose=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized Hinge Loss model using a single fused CUDA kernel that computes
    mean(clamp(1 - predictions * targets, min=0)) in one pass.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss

    def forward(self, predictions, targets):
        # Handle broadcastable targets (e.g. scalar) by materializing them.
        if targets.numel() != predictions.numel():
            targets = torch.broadcast_to(targets, predictions.shape)
        return self.hinge_loss.hinge_loss_cuda(predictions, targets)
```