```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 kernel launch (elementwise hinge + reduction fused together).
hinge_loss_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <ATen/cuda/CUDAContext.h>

template <typename scalar_t>
__global__ void hinge_loss_kernel(const scalar_t* __restrict__ preds,
                                  const scalar_t* __restrict__ tgts,
                                  scalar_t* __restrict__ out,
                                  const long n,
                                  const scalar_t inv_n) {
    __shared__ scalar_t sdata[256];
    const int tid = threadIdx.x;
    const long stride = (long)gridDim.x * blockDim.x;

    // Grid-stride loop: each thread accumulates its local hinge-loss sum.
    scalar_t local = 0;
    for (long i = (long)blockIdx.x * blockDim.x + tid; i < n; i += stride) {
        scalar_t v = (scalar_t)1 - preds[i] * tgts[i];
        local += v > (scalar_t)0 ? v : (scalar_t)0;
    }

    // Block-level reduction in shared memory.
    sdata[tid] = local;
    __syncthreads();

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

    // Each block adds its (already normalized) partial sum. Accumulating
    // sum * inv_n directly yields the final mean without a second kernel.
    if (tid == 0) {
        atomicAdd(out, sdata[0] * inv_n);
    }
}

torch::Tensor hinge_loss_cuda(torch::Tensor predictions, torch::Tensor targets) {
    // Match PyTorch broadcasting semantics exactly, then flatten.
    auto bc = at::broadcast_tensors({predictions, targets});
    auto p = bc[0].contiguous();
    auto t = bc[1].contiguous();

    const long n = p.numel();
    // 0-dim tensor, matching the shape returned by torch.mean(...).
    auto out = torch::zeros({}, p.options());
    if (n == 0) {
        return out;
    }

    const int block = 256;
    long grid = (n + block - 1) / block;
    if (grid > 4096) grid = 4096;

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

    AT_DISPATCH_FLOATING_TYPES(p.scalar_type(), "hinge_loss_cuda", ([&] {
        hinge_loss_kernel<scalar_t><<<(int)grid, block, 0, stream>>>(
            p.data_ptr<scalar_t>(),
            t.data_ptr<scalar_t>(),
            out.data_ptr<scalar_t>(),
            n,
            (scalar_t)(1.0 / (double)n));
    }));

    return out;
}
"""

hinge_loss_cpp_source = (
    "torch::Tensor hinge_loss_cuda(torch::Tensor predictions, 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=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    A model that computes Hinge Loss for binary classification tasks using a
    single fused CUDA kernel (hinge elementwise op + mean reduction in one pass).

    Parameters:
        None
    """
    def __init__(self):
        super(ModelNew, self).__init__()
        self.hinge_loss = hinge_loss

    def forward(self, predictions, targets):
        """
        Computes mean hinge loss with a fused custom CUDA kernel.

        Args:
            predictions (torch.Tensor): Predictions tensor of shape (batch_size, num_classes), float32.
            targets (torch.Tensor): Targets tensor of shape (batch_size,), float32.

        Returns:
            torch.Tensor: Scalar mean hinge loss.
        """
        return self.hinge_loss.hinge_loss_cuda(predictions, targets)
```