```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 (fused multiply -> hinge -> clamp -> sum reduction),
# followed by a tiny finalize kernel that divides by N.
hinge_loss_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <limits>

template <bool DIRECT>
__global__ void hinge_loss_sum_kernel(const float* __restrict__ preds,
                                      const float* __restrict__ targets,
                                      double* __restrict__ acc,
                                      const long long N,
                                      const long long T) {
    __shared__ double warp_sums[32];

    const long long tid = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    const long long stride = (long long)gridDim.x * blockDim.x;

    // For the broadcast case, track the target column index without
    // performing an expensive 64-bit modulo per element.
    long long col = 0;
    long long stride_mod = 0;
    if (!DIRECT) {
        col = tid % T;
        stride_mod = stride % T;
    }

    double local = 0.0;
    for (long long idx = tid; idx < N; idx += stride) {
        const float t = DIRECT ? targets[idx] : targets[col];
        float v = 1.0f - preds[idx] * t;
        v = v > 0.0f ? v : 0.0f;   // clamp(min=0)
        local += (double)v;
        if (!DIRECT) {
            col += stride_mod;
            if (col >= T) col -= T;
        }
    }

    // Warp-level reduction (double precision for accuracy).
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1) {
        local += __shfl_down_sync(0xffffffffu, local, offset);
    }

    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;
    if (lane == 0) {
        warp_sums[warp] = local;
    }
    __syncthreads();

    if (warp == 0) {
        const int num_warps = (int)((blockDim.x + 31) >> 5);
        double val = (lane < num_warps) ? warp_sums[lane] : 0.0;
        #pragma unroll
        for (int offset = 16; offset > 0; offset >>= 1) {
            val += __shfl_down_sync(0xffffffffu, val, offset);
        }
        if (lane == 0) {
            atomicAdd(acc, val);
        }
    }
}

__global__ void hinge_loss_finalize_kernel(const double* __restrict__ acc,
                                           float* __restrict__ out,
                                           const double inv_n) {
    out[0] = (float)(acc[0] * inv_n);
}

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");

    auto preds = predictions.to(torch::kFloat).contiguous();
    auto targs_raw = targets.to(torch::kFloat);

    long long N = (long long)preds.numel();

    torch::Tensor targs;
    bool direct = true;
    long long T = N;

    if (preds.sizes() == targs_raw.sizes()) {
        // Identical shapes: pure element-wise.
        targs = targs_raw.contiguous();
        direct = true;
        T = N;
    } else {
        // Common case: targets broadcasts along the last dimension
        // (e.g. predictions (B, C) with targets (C,), or scalar targets).
        bool simple_lastdim = false;
        long long C = (long long)targs_raw.numel();
        if (C == 1) {
            simple_lastdim = true;
        } else if (preds.dim() >= 1 && targs_raw.dim() >= 1 &&
                   targs_raw.dim() <= preds.dim() &&
                   targs_raw.size(-1) == preds.size(-1)) {
            simple_lastdim = true;
            for (int64_t d = 0; d + 1 < targs_raw.dim(); ++d) {
                if (targs_raw.size(d) != 1) { simple_lastdim = false; break; }
            }
            C = (long long)preds.size(-1);
        }

        if (simple_lastdim) {
            targs = targs_raw.contiguous();
            direct = false;
            T = C;
        } else {
            // Fully general broadcasting fallback.
            auto both = at::broadcast_tensors({preds, targs_raw});
            preds = both[0].contiguous();
            targs = both[1].contiguous();
            N = (long long)preds.numel();
            direct = true;
            T = N;
        }
    }

    auto options_f = torch::TensorOptions().dtype(torch::kFloat).device(preds.device());
    auto options_d = torch::TensorOptions().dtype(torch::kDouble).device(preds.device());
    auto out = torch::empty({}, options_f);

    if (N == 0) {
        out.fill_(std::numeric_limits<float>::quiet_NaN());
        return out;
    }

    auto acc = torch::zeros({}, options_d);

    const int block = 256;
    const long long needed = (N + block - 1) / block;
    const int grid = (int)std::min<long long>(needed, 8192);

    auto stream = at::cuda::getCurrentCUDAStream();
    if (direct) {
        hinge_loss_sum_kernel<true><<<grid, block, 0, stream>>>(
            preds.data_ptr<float>(), targs.data_ptr<float>(),
            acc.data_ptr<double>(), N, T);
    } else {
        hinge_loss_sum_kernel<false><<<grid, block, 0, stream>>>(
            preds.data_ptr<float>(), targs.data_ptr<float>(),
            acc.data_ptr<double>(), N, T);
    }
    hinge_loss_finalize_kernel<<<1, 1, 0, stream>>>(
        acc.data_ptr<double>(), out.data_ptr<float>(), 1.0 / (double)N);

    return out;
}
"""

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

# Compile the inline CUDA extension.
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=False,
)


class ModelNew(nn.Module):
    """
    Optimized hinge loss model using a fused custom CUDA kernel.
    """
    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)
```