```python
import torch
import torch.nn as nn
from torch.utils.cpp_extension import load_inline

_cpp_source = """
#include <torch/extension.h>
torch::Tensor hinge_loss_cuda(torch::Tensor predictions, torch::Tensor targets);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

__device__ __forceinline__ float block_sum(float value) {
    __shared__ float warp_sums[8];

    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;

    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1)
        value += __shfl_down_sync(0xffffffff, value, offset);

    if (lane == 0)
        warp_sums[warp] = value;
    __syncthreads();

    if (warp == 0) {
        value = lane < 8 ? warp_sums[lane] : 0.0f;
        #pragma unroll
        for (int offset = 16; offset > 0; offset >>= 1)
            value += __shfl_down_sync(0xffffffff, value, offset);
    }
    return value;
}

template<int TargetMode>
__global__ void hinge_reduce(
    const float* __restrict__ predictions,
    const float* __restrict__ targets,
    float* __restrict__ result,
    int64_t count,
    int64_t target_count,
    bool final_output
) {
    float sum = 0.0f;
    const int64_t step = static_cast<int64_t>(gridDim.x) * blockDim.x;

    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < count;
         i += step) {
        int64_t target_index;
        if (TargetMode == 0)
            target_index = i;
        else if (TargetMode == 1)
            target_index = 0;
        else
            target_index = i % target_count;

        // Separate operations preserve PyTorch's float32 rounding.
        const float product = __fmul_rn(predictions[i], targets[target_index]);
        const float margin = __fsub_rn(1.0f, product);
        sum += margin < 0.0f ? 0.0f : margin;
    }

    sum = block_sum(sum);
    if (threadIdx.x == 0)
        result[blockIdx.x] = final_output
            ? (count == 0 ? nanf("") : sum / static_cast<float>(count))
            : sum;
}

__global__ void finish_reduce(
    const float* __restrict__ partials,
    float* __restrict__ output,
    int partial_count,
    int64_t count
) {
    float sum = 0.0f;
    for (int i = threadIdx.x; i < partial_count; i += blockDim.x)
        sum += partials[i];

    sum = block_sum(sum);
    if (threadIdx.x == 0)
        output[0] = sum / static_cast<float>(count);
}

torch::Tensor hinge_loss_cuda(
    torch::Tensor predictions,
    torch::Tensor targets
) {
    TORCH_CHECK(predictions.is_cuda() && targets.is_cuda(),
                "Inputs must be CUDA tensors");
    TORCH_CHECK(predictions.device() == targets.device(),
                "Inputs must be on the same device");
    TORCH_CHECK(predictions.scalar_type() == torch::kFloat32 &&
                targets.scalar_type() == torch::kFloat32,
                "Inputs must be float32");
    TORCH_CHECK(predictions.is_contiguous() && targets.is_contiguous(),
                "Inputs must be contiguous");

    const c10::cuda::CUDAGuard guard(predictions.device());
    const int64_t count = predictions.numel();
    const int64_t target_count = targets.numel();

    TORCH_CHECK(count == 0 ||
                (target_count > 0 && count % target_count == 0),
                "Invalid target broadcasting");

    auto output = torch::empty({}, predictions.options());
    const int blocks = static_cast<int>(
        std::min<int64_t>(1024, std::max<int64_t>(1, (count + 1023) / 1024))
    );

    auto partials = blocks == 1
        ? output
        : torch::empty({blocks}, predictions.options());

    const auto stream = at::cuda::getCurrentCUDAStream();
    const float* p = predictions.data_ptr<float>();
    const float* t = targets.data_ptr<float>();
    float* r = partials.data_ptr<float>();

    if (target_count == count) {
        hinge_reduce<0><<<blocks, 256, 0, stream>>>(
            p, t, r, count, target_count, blocks == 1);
    } else if (target_count == 1) {
        hinge_reduce<1><<<blocks, 256, 0, stream>>>(
            p, t, r, count, target_count, blocks == 1);
    } else {
        hinge_reduce<2><<<blocks, 256, 0, stream>>>(
            p, t, r, count, target_count, blocks == 1);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    if (blocks > 1) {
        finish_reduce<<<1, 256, 0, stream>>>(
            r, output.data_ptr<float>(), blocks, count);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }
    return output;
}
"""

_hinge_extension = load_inline(
    name="fused_hinge_loss_cuda_v1",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["hinge_loss_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3", "--fmad=false"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()

    def forward(self, predictions, targets):
        if (
            not predictions.is_cuda
            or not targets.is_cuda
            or predictions.dtype != torch.float32
            or targets.dtype != torch.float32
        ):
            return torch.mean(torch.clamp(1 - predictions * targets, min=0))

        if (
            predictions.shape == targets.shape
            or targets.numel() == 1
            or (
                targets.ndim == 1
                and predictions.ndim >= 1
                and targets.shape[0] == predictions.shape[-1]
            )
        ):
            return _hinge_extension.hinge_loss_cuda(
                predictions.contiguous(), targets.contiguous()
            )

        predictions, targets = torch.broadcast_tensors(predictions, targets)
        return _hinge_extension.hinge_loss_cuda(
            predictions.contiguous(), targets.contiguous()
        )
```