```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 clamp(1 - p*t, 0) and the mean reduction
# in a single pass over the data (block reduction + atomic accumulation).
hinge_loss_cuda_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <cstdint>
#include <limits>

__global__ void hinge_loss_fused_kernel(const float* __restrict__ preds,
                                        const float* __restrict__ targs,
                                        float* __restrict__ out,
                                        int64_t size) {
    __shared__ float shared[256];
    int tid = threadIdx.x;
    int64_t idx = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    int64_t stride = (int64_t)blockDim.x * gridDim.x;

    // Grid-stride loop: each thread accumulates its local hinge-loss sum.
    float acc = 0.0f;
    for (int64_t i = idx; i < size; i += stride) {
        float v = 1.0f - preds[i] * targs[i];
        acc += v > 0.0f ? v : 0.0f;
    }

    shared[tid] = acc;
    __syncthreads();

    // Block-level reduction in shared memory.
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            shared[tid] += shared[tid + s];
        }
        __syncthreads();
    }

    // Accumulate per-block partial sums into the global scalar.
    if (tid == 0) {
        atomicAdd(out, shared[0]);
    }
}

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

    auto preds = predictions.contiguous();
    auto targs = targets.contiguous();
    const int64_t size = preds.numel();

    auto out = torch::zeros({}, preds.options());
    if (size == 0) {
        out.fill_(std::numeric_limits<float>::quiet_NaN());
        return out;
    }

    const int block_size = 256;
    int64_t num_blocks = (size + block_size - 1) / block_size;
    if (num_blocks > 4096) {
        num_blocks = 4096;
    }

    hinge_loss_fused_kernel<<<(int)num_blocks, block_size>>>(
        preds.data_ptr<float>(),
        targs.data_ptr<float>(),
        out.data_ptr<float>(),
        size);

    // Convert accumulated sum into the mean.
    out.div_(size);
    return out;
}
"""

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

# Compile the inline 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=True,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized Hinge Loss using a single fused CUDA kernel that performs
    the elementwise hinge computation and the mean reduction in one pass.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss

    def forward(self, predictions, targets):
        if predictions.dtype != torch.float32:
            predictions = predictions.float()
        if targets.dtype != torch.float32:
            targets = targets.float()
        # Match PyTorch broadcasting semantics of predictions * targets.
        if predictions.shape != targets.shape:
            predictions, targets = torch.broadcast_tensors(predictions, targets)
        return self.hinge_loss.hinge_loss_cuda(predictions, targets)
```