```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 that computes mean hinge loss:
#   mean(clamp(1 - predictions * targets, min=0))
# in a single pass (multiply + clamp + reduction fused), with full support
# for PyTorch broadcasting semantics between predictions and targets.
hinge_loss_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <limits>

#define BLOCK_SIZE 256
#define MAX_DIMS 8

struct BroadcastInfo {
    int64_t shape[MAX_DIMS];
    int64_t p_strides[MAX_DIMS];
    int64_t t_strides[MAX_DIMS];
    int ndim;
};

// Fast path: both tensors are contiguous with identical number of elements.
__global__ void hinge_sum_contig_kernel(const float* __restrict__ p,
                                        const float* __restrict__ t,
                                        double* __restrict__ acc,
                                        int64_t N) {
    __shared__ double sdata[BLOCK_SIZE];
    double local = 0.0;
    int64_t tid = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    int64_t stride = (int64_t)gridDim.x * blockDim.x;
    for (int64_t i = tid; i < N; i += stride) {
        float v = 1.0f - p[i] * t[i];
        local += (double)(v > 0.0f ? v : 0.0f);
    }
    sdata[threadIdx.x] = local;
    __syncthreads();
    #pragma unroll
    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (threadIdx.x < s) sdata[threadIdx.x] += sdata[threadIdx.x + s];
        __syncthreads();
    }
    if (threadIdx.x == 0) atomicAdd(acc, sdata[0]);
}

// Generic path: handles arbitrary broadcasting between the two inputs.
__global__ void hinge_sum_broadcast_kernel(const float* __restrict__ p,
                                           const float* __restrict__ t,
                                           double* __restrict__ acc,
                                           BroadcastInfo info,
                                           int64_t N) {
    __shared__ double sdata[BLOCK_SIZE];
    double local = 0.0;
    int64_t tid = (int64_t)blockIdx.x * blockDim.x + threadIdx.x;
    int64_t stride = (int64_t)gridDim.x * blockDim.x;
    for (int64_t i = tid; i < N; i += stride) {
        int64_t rem = i;
        int64_t po = 0, to = 0;
        for (int d = info.ndim - 1; d >= 0; --d) {
            int64_t c = rem % info.shape[d];
            rem /= info.shape[d];
            po += c * info.p_strides[d];
            to += c * info.t_strides[d];
        }
        float v = 1.0f - p[po] * t[to];
        local += (double)(v > 0.0f ? v : 0.0f);
    }
    sdata[threadIdx.x] = local;
    __syncthreads();
    #pragma unroll
    for (int s = BLOCK_SIZE / 2; s > 0; s >>= 1) {
        if (threadIdx.x < s) sdata[threadIdx.x] += sdata[threadIdx.x + s];
        __syncthreads();
    }
    if (threadIdx.x == 0) atomicAdd(acc, sdata[0]);
}

// Finalize: divide accumulated sum by N and cast to float32.
__global__ void finalize_kernel(const double* __restrict__ acc,
                                float* __restrict__ out,
                                double invN) {
    out[0] = (float)(acc[0] * invN);
}

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 p = predictions.contiguous();
    auto t = targets.contiguous();

    auto out_shape = at::infer_size(p.sizes(), t.sizes());
    int64_t N = 1;
    for (auto s : out_shape) N *= s;

    auto out = torch::empty({}, p.options());
    auto acc = torch::zeros({}, p.options().dtype(torch::kDouble));

    int64_t blocks64 = (N + BLOCK_SIZE - 1) / BLOCK_SIZE;
    int blocks = (int)std::min<int64_t>(std::max<int64_t>(blocks64, 1), 65535);

    auto stream = at::cuda::getCurrentCUDAStream();

    bool fast = (p.numel() == N) && (t.numel() == N) &&
                p.is_contiguous() && t.is_contiguous();

    if (fast) {
        hinge_sum_contig_kernel<<<blocks, BLOCK_SIZE, 0, stream>>>(
            p.data_ptr<float>(), t.data_ptr<float>(),
            acc.data_ptr<double>(), N);
    } else {
        TORCH_CHECK((int)out_shape.size() <= MAX_DIMS,
                    "broadcast rank too large");
        BroadcastInfo info;
        info.ndim = (int)out_shape.size();
        auto pe = p.expand(out_shape);
        auto te = t.expand(out_shape);
        for (int d = 0; d < info.ndim; ++d) {
            info.shape[d] = out_shape[d];
            info.p_strides[d] = pe.strides()[d];
            info.t_strides[d] = te.strides()[d];
        }
        hinge_sum_broadcast_kernel<<<blocks, BLOCK_SIZE, 0, stream>>>(
            p.data_ptr<float>(), t.data_ptr<float>(),
            acc.data_ptr<double>(), info, N);
    }

    double invN = (N > 0) ? (1.0 / (double)N)
                          : std::numeric_limits<double>::quiet_NaN();
    finalize_kernel<<<1, 1, 0, stream>>>(
        acc.data_ptr<double>(), out.data_ptr<float>(), invN);

    return out;
}
"""

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

# Compile the inline CUDA extension for the fused hinge loss.
hinge_loss = load_inline(
    name="fused_hinge_loss",
    cpp_sources=hinge_loss_cpp_source,
    cuda_sources=hinge_loss_cuda_source,
    functions=["hinge_loss_cuda"],
    verbose=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    """
    Optimized model that computes mean Hinge Loss with a single fused
    custom CUDA kernel (multiply + clamp + mean reduction in one pass).
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss

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