```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(clamp(1 - predictions * targets, min=0))
# in a single pass: the multiply + clamp are fused into a block-level
# reduction (shared memory tree reduction) followed by one atomicAdd per
# block into a scalar accumulator, which is finally divided by the element
# count to produce the mean.
hinge_loss_cuda_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__global__ void hinge_loss_kernel(const float* __restrict__ preds,
                                  const float* __restrict__ targets,
                                  float* __restrict__ out,
                                  long size,
                                  bool has_targets) {
    __shared__ float shared[256];

    // Grid-stride loop: each thread accumulates its partial hinge sum.
    float local_sum = 0.0f;
    long stride = (long)gridDim.x * blockDim.x;
    for (long idx = (long)blockIdx.x * blockDim.x + threadIdx.x; idx < size; idx += stride) {
        float p = preds[idx];
        float val = has_targets ? (1.0f - p * targets[idx]) : (1.0f - p);
        local_sum += fmaxf(val, 0.0f);
    }

    // Block-level tree reduction in shared memory.
    shared[threadIdx.x] = local_sum;
    __syncthreads();

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

    // One atomic add per block into the global scalar accumulator.
    if (threadIdx.x == 0) {
        atomicAdd(out, shared[0]);
    }
}

torch::Tensor hinge_loss_cuda(torch::Tensor preds, torch::Tensor targets, bool has_targets) {
    const long size = preds.numel();

    // Scalar (0-dim) accumulator on the same device/dtype as the input.
    auto out = torch::zeros({}, preds.options());

    if (size > 0) {
        const int block_size = 256;
        long num_blocks = (size + block_size - 1) / block_size;
        // Cap the grid; the grid-stride loop covers the remaining elements
        // and this keeps atomic contention low.
        if (num_blocks > 4096) {
            num_blocks = 4096;
        }

        const float* targets_ptr = has_targets ? targets.data_ptr<float>() : nullptr;

        hinge_loss_kernel<<<num_blocks, block_size>>>(
            preds.data_ptr<float>(),
            targets_ptr,
            out.data_ptr<float>(),
            size,
            has_targets);
    }

    // Convert the sum into a mean (empty input yields NaN, matching torch.mean).
    out.div_((double)size);
    return out;
}
"""

hinge_loss_cpp_source = (
    "torch::Tensor hinge_loss_cuda(torch::Tensor preds, torch::Tensor targets, bool has_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,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Hinge Loss for binary classification, computed with a single fused
    CUDA kernel (multiply + clamp + mean reduction in one launch).

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

    def forward(self, predictions, targets):
        predictions = predictions.contiguous()
        targets = targets.contiguous()

        if predictions.shape == targets.shape:
            # Fast path: fuse the elementwise multiply, clamp, and mean
            # reduction entirely inside one CUDA kernel.
            return self.hinge_loss.hinge_loss_cuda(
                predictions.view(-1), targets.view(-1), True
            )

        # General broadcasting fallback: materialize the broadcast product,
        # then let the kernel fuse the clamp + mean reduction.
        prod = torch.mul(predictions, targets).contiguous().view(-1)
        return self.hinge_loss.hinge_loss_cuda(prod, prod, False)
```