```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 max(0, 1 - p*t) elementwise and reduces to the
# mean in a single kernel launch, avoiding all intermediate tensors.
hinge_loss_source = """
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__global__ void hinge_loss_kernel(const float* __restrict__ preds,
                                  const float* __restrict__ targets,
                                  double* __restrict__ out,
                                  long long numel,
                                  long long num_targets) {
    long long idx = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    const long long stride = (long long)gridDim.x * blockDim.x;

    float acc = 0.0f;
    for (; idx < numel; idx += stride) {
        float p = preds[idx];
        // targets broadcasts along the trailing dimension of preds
        float t = targets[idx % num_targets];
        acc += fmaxf(0.0f, 1.0f - p * t);
    }

    // Warp-level reduction
    for (int offset = 16; offset > 0; offset >>= 1) {
        acc += __shfl_down_sync(0xffffffffu, acc, offset);
    }

    __shared__ float sdata[32];
    const int lane = threadIdx.x & 31;
    const int wid = threadIdx.x >> 5;
    if (lane == 0) {
        sdata[wid] = acc;
    }
    __syncthreads();

    if (wid == 0) {
        const int nwarps = (blockDim.x + 31) >> 5;
        acc = (lane < nwarps) ? sdata[lane] : 0.0f;
        for (int offset = 16; offset > 0; offset >>= 1) {
            acc += __shfl_down_sync(0xffffffffu, acc, offset);
        }
        if (lane == 0) {
            // Accumulate in double for accuracy; scaled later by 1/numel.
            atomicAdd(out, (double)acc);
        }
    }
}

torch::Tensor hinge_loss_cuda(torch::Tensor preds, torch::Tensor targets) {
    TORCH_CHECK(preds.is_cuda() && targets.is_cuda(), "Inputs must be CUDA tensors");
    TORCH_CHECK(preds.scalar_type() == torch::kFloat, "preds must be float32");
    TORCH_CHECK(targets.scalar_type() == torch::kFloat, "targets must be float32");
    TORCH_CHECK(preds.is_contiguous() && targets.is_contiguous(), "Inputs must be contiguous");

    const long long numel = preds.numel();
    const long long num_targets = targets.numel();

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

    if (numel > 0 && num_targets > 0) {
        const int block = 256;
        long long grid = (numel + block - 1) / block;
        const long long max_grid = 8192;
        if (grid > max_grid) grid = max_grid;

        cudaStream_t stream = at::cuda::getCurrentCUDAStream();
        hinge_loss_kernel<<<grid, block, 0, stream>>>(
            preds.data_ptr<float>(),
            targets.data_ptr<float>(),
            out.data_ptr<double>(),
            numel,
            num_targets);
    }

    // numel == 0 yields NaN, matching torch.mean of an empty tensor.
    return (out / (double)numel).to(torch::kFloat);
}
"""

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

hinge_loss = load_inline(
    name="hinge_loss_fused",
    cpp_sources=hinge_loss_cpp_source,
    cuda_sources=hinge_loss_source,
    functions=["hinge_loss_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


class ModelNew(nn.Module):
    """
    Optimized Hinge Loss using a single fused CUDA kernel that computes
    mean(clamp(1 - predictions * targets, min=0)) without intermediates.
    """

    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.contiguous(), targets.contiguous()
        )
```