```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) and the mean reduction in a
# single pass, replacing the mul / rsub / clamp / mean chain of PyTorch ops.
hinge_loss_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>

#define BLOCK_SIZE 256
#define NUM_WARPS (BLOCK_SIZE / 32)

__device__ __forceinline__ void blockReduceAdd(float val, float* out) {
    __shared__ float shared[NUM_WARPS];
    const int lane = threadIdx.x & 31;
    const int wid = threadIdx.x >> 5;

    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffffu, val, offset);
    }
    if (lane == 0) {
        shared[wid] = val;
    }
    __syncthreads();

    if (wid == 0) {
        val = (lane < NUM_WARPS) ? shared[lane] : 0.0f;
        #pragma unroll
        for (int offset = 16; offset > 0; offset >>= 1) {
            val += __shfl_down_sync(0xffffffffu, val, offset);
        }
        if (lane == 0) {
            atomicAdd(out, val);
        }
    }
}

__global__ void hingeLossKernel(const float* __restrict__ preds,
                                const float* __restrict__ tgts,
                                float* __restrict__ out,
                                const long long n) {
    const long long tid = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    const long long stride = (long long)gridDim.x * blockDim.x;
    float acc = 0.0f;
    for (long long i = tid; i < n; i += stride) {
        acc += fmaxf(1.0f - preds[i] * tgts[i], 0.0f);
    }
    blockReduceAdd(acc, out);
}

__global__ void hingeLossKernelVec4(const float* __restrict__ preds,
                                    const float* __restrict__ tgts,
                                    float* __restrict__ out,
                                    const long long n) {
    const long long n4 = n >> 2;
    const float4* preds4 = reinterpret_cast<const float4*>(preds);
    const float4* tgts4 = reinterpret_cast<const float4*>(tgts);
    const long long tid = (long long)blockIdx.x * blockDim.x + threadIdx.x;
    const long long stride = (long long)gridDim.x * blockDim.x;
    float acc = 0.0f;
    for (long long i = tid; i < n4; i += stride) {
        const float4 p = preds4[i];
        const float4 t = tgts4[i];
        acc += fmaxf(1.0f - p.x * t.x, 0.0f);
        acc += fmaxf(1.0f - p.y * t.y, 0.0f);
        acc += fmaxf(1.0f - p.z * t.z, 0.0f);
        acc += fmaxf(1.0f - p.w * t.w, 0.0f);
    }
    // Handle the tail elements (n % 4).
    const long long i = (n4 << 2) + tid;
    if (i < n) {
        acc += fmaxf(1.0f - preds[i] * tgts[i], 0.0f);
    }
    blockReduceAdd(acc, out);
}

torch::Tensor hinge_loss_cuda(torch::Tensor preds, torch::Tensor tgts) {
    TORCH_CHECK(preds.is_cuda() && tgts.is_cuda(), "inputs must be CUDA tensors");
    TORCH_CHECK(preds.scalar_type() == torch::kFloat32 && tgts.scalar_type() == torch::kFloat32,
                "inputs must be float32");
    TORCH_CHECK(preds.is_contiguous() && tgts.is_contiguous(), "inputs must be contiguous");
    TORCH_CHECK(preds.numel() == tgts.numel(), "inputs must have the same number of elements");

    const long long n = preds.numel();
    auto out = torch::zeros({}, preds.options());

    if (n > 0) {
        const int sms = at::cuda::getCurrentDeviceProperties()->multiProcessorCount;
        const long long maxBlocks = (long long)sms * 8;
        const long long needed = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;
        const int blocks = (int)(needed < maxBlocks ? needed : maxBlocks);
        cudaStream_t stream = at::cuda::getCurrentCUDAStream().stream();
        const float* pPtr = preds.data_ptr<float>();
        const float* tPtr = tgts.data_ptr<float>();
        const bool aligned =
            ((reinterpret_cast<uintptr_t>(pPtr) & 0xF) == 0) &&
            ((reinterpret_cast<uintptr_t>(tPtr) & 0xF) == 0);
        if (aligned) {
            hingeLossKernelVec4<<<blocks, BLOCK_SIZE, 0, stream>>>(
                pPtr, tPtr, out.data_ptr<float>(), n);
        } else {
            hingeLossKernel<<<blocks, BLOCK_SIZE, 0, stream>>>(
                pPtr, tPtr, out.data_ptr<float>(), n);
        }
    }

    out.div_(static_cast<double>(n));  // n == 0 -> NaN, matching torch.mean of an empty tensor
    return out;
}
"""

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

# Compile the inline CUDA extension.
hinge_loss_ext = load_inline(
    name="fused_hinge_loss_ext",
    cpp_sources=hinge_loss_cpp_source,
    cuda_sources=hinge_loss_cuda_source,
    functions=["hinge_loss_cuda"],
    verbose=False,
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
)


def _sum_to_size(grad, shape):
    """Reduce a broadcasted gradient back to the original input shape."""
    while grad.dim() > len(shape):
        grad = grad.sum(dim=0)
    for i, s in enumerate(shape):
        if s == 1 and grad.shape[i] != 1:
            grad = grad.sum(dim=i, keepdim=True)
    return grad.reshape(shape)


class _FusedHingeLoss(torch.autograd.Function):
    @staticmethod
    def forward(ctx, predictions, targets):
        preds, tgts = torch.broadcast_tensors(predictions, targets)
        preds = preds.contiguous()
        tgts = tgts.contiguous()
        loss = hinge_loss_ext.hinge_loss_cuda(preds.view(-1), tgts.view(-1))
        ctx.save_for_backward(preds, tgts)
        ctx.pred_shape = predictions.shape
        ctx.tgt_shape = targets.shape
        return loss

    @staticmethod
    def backward(ctx, grad_output):
        preds, tgts = ctx.saved_tensors
        n = preds.numel()
        # d/dp mean(clamp(1 - p*t, min=0)) = -t / n where 1 - p*t >= 0 (inclusive,
        # matching torch.clamp's backward), else 0.
        active = (preds * tgts <= 1.0).to(preds.dtype)
        scale = grad_output / float(n)
        grad_pred = grad_tgt = None
        if ctx.needs_input_grad[0]:
            grad_pred = _sum_to_size((-tgts) * active * scale, ctx.pred_shape)
        if ctx.needs_input_grad[1]:
            grad_tgt = _sum_to_size((-preds) * active * scale, ctx.tgt_shape)
        return grad_pred, grad_tgt


class ModelNew(nn.Module):
    """
    Optimized Hinge Loss: the elementwise hinge computation and the mean
    reduction are fused into a single custom CUDA kernel.
    """

    def __init__(self):
        super(ModelNew, self).__init__()

    def forward(self, predictions, targets):
        if (predictions.is_cuda and targets.is_cuda
                and predictions.dtype == torch.float32
                and targets.dtype == torch.float32):
            return _FusedHingeLoss.apply(predictions, targets)
        # Fallback for CPU / non-float32 inputs.
        return torch.mean(torch.clamp(1 - predictions * targets, min=0))
```