```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(clamp(1 - predictions * targets, min=0))
# in a single pass using a grid-stride loop + block-level tree reduction
# + atomic accumulation of per-block partial means.
hinge_loss_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <limits>

#define BLOCK_SIZE 256

__global__ void hinge_loss_kernel(const float* __restrict__ preds,
                                  const float* __restrict__ targets,
                                  float* __restrict__ out,
                                  const long size,
                                  const float inv_size) {
    const long stride = (long)blockDim.x * gridDim.x;
    float local_sum = 0.0f;

    // Grid-stride loop: each thread accumulates clamp(1 - p*t, 0) locally.
    for (long i = (long)blockIdx.x * blockDim.x + threadIdx.x; i < size; i += stride) {
        float val = 1.0f - preds[i] * targets[i];
        local_sum += fmaxf(val, 0.0f);
    }

    // Block-level tree reduction in shared memory.
    __shared__ float shared[BLOCK_SIZE];
    const int tid = threadIdx.x;
    shared[tid] = local_sum;
    __syncthreads();

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

    // Accumulate this block's contribution to the mean.
    if (tid == 0) {
        atomicAdd(out, shared[0] * inv_size);
    }
}

torch::Tensor hinge_loss_cuda(torch::Tensor preds, torch::Tensor targets) {
    const long size = preds.numel();
    auto out = torch::zeros({}, preds.options());

    if (size == 0) {
        // Match torch.mean behavior on an empty tensor (NaN).
        out.fill_(std::numeric_limits<float>::quiet_NaN());
        return out;
    }

    int num_blocks = (int)((size + BLOCK_SIZE - 1) / BLOCK_SIZE);
    if (num_blocks > 512) num_blocks = 512;

    const float inv_size = 1.0f / (float)size;

    hinge_loss_kernel<<<num_blocks, BLOCK_SIZE, 0, at::cuda::getCurrentCUDAStream()>>>(
        preds.data_ptr<float>(),
        targets.data_ptr<float>(),
        out.data_ptr<float>(),
        size,
        inv_size);

    return out;
}
"""

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

# 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_source,
    functions=["hinge_loss_cuda"],
    verbose=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized Hinge Loss: fuses multiply, subtract, clamp, and mean-reduction
    into a single custom CUDA kernel.
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss

    def forward(self, predictions, targets):
        # Match PyTorch broadcasting semantics, then feed contiguous
        # buffers to the fused kernel.
        p, t = torch.broadcast_tensors(predictions, targets)
        p = p.contiguous()
        t = t.contiguous()
        return self.hinge_loss.hinge_loss_cuda(p, t)
```