```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. Each thread computes the hinge value for its elements,
# accumulates a local sum, performs a block-level reduction, and atomically
# accumulates into the output scalar (pre-scaled by 1/N to produce the mean).
hinge_loss_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <limits>

#define MAX_DIMS 8
#define BLOCK_SIZE 256

struct BroadcastInfo {
    int ndim;
    long long shape[MAX_DIMS];
    long long strideP[MAX_DIMS];
    long long strideT[MAX_DIMS];
};

__device__ __forceinline__ float warpReduceSum(float v) {
    for (int offset = 16; offset > 0; offset >>= 1)
        v += __shfl_down_sync(0xffffffff, v, offset);
    return v;
}

__device__ __forceinline__ float blockReduceSum(float v) {
    __shared__ float shared[BLOCK_SIZE / 32];
    int lane = threadIdx.x & 31;
    int wid  = threadIdx.x >> 5;
    v = warpReduceSum(v);
    if (lane == 0) shared[wid] = v;
    __syncthreads();
    int nwarps = BLOCK_SIZE / 32;
    v = (threadIdx.x < nwarps) ? shared[threadIdx.x] : 0.0f;
    if (wid == 0) v = warpReduceSum(v);
    return v;
}

// Fast path: both tensors contiguous with identical layout.
__global__ void hinge_loss_contig_kernel(const float* __restrict__ p,
                                         const float* __restrict__ t,
                                         float* __restrict__ out,
                                         long long numel,
                                         float inv_n) {
    long long idx    = blockIdx.x * (long long)BLOCK_SIZE + threadIdx.x;
    long long stride = (long long)gridDim.x * BLOCK_SIZE;
    float local = 0.0f;
    for (long long i = idx; i < numel; i += stride) {
        float v = fmaf(-p[i], t[i], 1.0f);   // 1 - p * t
        v = v > 0.0f ? v : 0.0f;             // clamp(min=0)
        local += v;
    }
    local *= inv_n;                          // scale to mean contribution
    local = blockReduceSum(local);
    if (threadIdx.x == 0 && local != 0.0f) atomicAdd(out, local);
}

// Generic path: handles arbitrary broadcasting / strides.
__global__ void hinge_loss_strided_kernel(const float* __restrict__ p,
                                          const float* __restrict__ t,
                                          float* __restrict__ out,
                                          long long numel,
                                          float inv_n,
                                          BroadcastInfo info) {
    long long idx    = blockIdx.x * (long long)BLOCK_SIZE + threadIdx.x;
    long long stride = (long long)gridDim.x * BLOCK_SIZE;
    float local = 0.0f;
    for (long long i = idx; i < numel; i += stride) {
        long long offP = 0, offT = 0;
        long long rem = i;
        #pragma unroll
        for (int d = MAX_DIMS - 1; d >= 0; --d) {
            if (d < info.ndim) {
                long long coord = rem % info.shape[d];
                rem /= info.shape[d];
                offP += coord * info.strideP[d];
                offT += coord * info.strideT[d];
            }
        }
        float v = fmaf(-p[offP], t[offT], 1.0f);
        v = v > 0.0f ? v : 0.0f;
        local += v;
    }
    local *= inv_n;
    local = blockReduceSum(local);
    if (threadIdx.x == 0 && local != 0.0f) atomicAdd(out, local);
}

torch::Tensor hinge_loss_cuda(torch::Tensor predictions, torch::Tensor targets) {
    TORCH_CHECK(predictions.is_cuda() && targets.is_cuda(),
                "inputs must be CUDA tensors");

    auto p_in = predictions.scalar_type() == at::kFloat ? predictions
                                                        : predictions.to(at::kFloat);
    auto t_in = targets.scalar_type() == at::kFloat ? targets
                                                    : targets.to(at::kFloat);

    // Match PyTorch broadcasting semantics exactly.
    auto bc = at::broadcast_tensors({p_in, t_in});
    auto p = bc[0];
    auto t = bc[1];

    long long numel = p.numel();
    auto out = at::zeros({}, p_in.options().dtype(at::kFloat));
    if (numel == 0) {
        // torch.mean of an empty tensor returns NaN.
        out.fill_(std::numeric_limits<float>::quiet_NaN());
        return out;
    }
    float inv_n = 1.0f / (float)numel;

    long long blocks_ll = (numel + BLOCK_SIZE - 1) / BLOCK_SIZE;
    int blocks = (int)std::min<long long>(blocks_ll, 65535);
    auto stream = at::cuda::getCurrentCUDAStream();

    if (p.is_contiguous() && t.is_contiguous()) {
        hinge_loss_contig_kernel<<<blocks, BLOCK_SIZE, 0, stream>>>(
            p.data_ptr<float>(), t.data_ptr<float>(),
            out.data_ptr<float>(), numel, inv_n);
    } else {
        BroadcastInfo info;
        info.ndim = (int)p.dim();
        TORCH_CHECK(info.ndim <= MAX_DIMS, "too many dimensions");
        for (int d = 0; d < info.ndim; ++d) {
            info.shape[d]   = p.size(d);
            info.strideP[d] = p.stride(d);
            info.strideT[d] = t.stride(d);
        }
        hinge_loss_strided_kernel<<<blocks, BLOCK_SIZE, 0, stream>>>(
            p.data_ptr<float>(), t.data_ptr<float>(),
            out.data_ptr<float>(), numel, inv_n, info);
    }
    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_op = 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_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Hinge Loss model backed by 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_op

    def forward(self, predictions, targets):
        return self.hinge_loss.hinge_loss_cuda(predictions, targets)
```