```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/ExpandUtils.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <algorithm>
#include <cstdint>

struct BroadcastInfo {
    int ndim;
    int64_t sizes[64];
    int64_t pred_strides[64];
    int64_t target_strides[64];
};

__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();

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

template<bool Contiguous>
__global__ void hinge_partial_kernel(
    const float* __restrict__ predictions,
    const float* __restrict__ targets,
    float* __restrict__ partials,
    int64_t n,
    BroadcastInfo info
) {
    float sum = 0.0f;
    const int64_t step = static_cast<int64_t>(blockDim.x) * gridDim.x;

    for (int64_t i = static_cast<int64_t>(blockIdx.x) * blockDim.x
                     + threadIdx.x;
         i < n;
         i += step) {
        int64_t p_offset = i;
        int64_t t_offset = i;

        if (!Contiguous) {
            p_offset = 0;
            t_offset = 0;
            int64_t remaining = i;
            for (int d = info.ndim - 1; d >= 0; --d) {
                const int64_t coordinate = remaining % info.sizes[d];
                remaining /= info.sizes[d];
                p_offset += coordinate * info.pred_strides[d];
                t_offset += coordinate * info.target_strides[d];
            }
        }

        const float product = __fmul_rn(
            predictions[p_offset], targets[t_offset]);
        const float margin = __fsub_rn(1.0f, product);
        sum += margin < 0.0f ? 0.0f : margin;
    }

    sum = block_sum(sum);
    if (threadIdx.x == 0)
        partials[blockIdx.x] = sum;
}

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

    sum = block_sum(sum);
    if (threadIdx.x == 0)
        output[0] = n == 0 ? nanf("") : sum / static_cast<float>(n);
}

torch::Tensor hinge_loss_cuda(
    torch::Tensor predictions,
    torch::Tensor targets
) {
    TORCH_CHECK(predictions.is_cuda() && targets.is_cuda(),
                "Both inputs must be CUDA tensors");
    TORCH_CHECK(predictions.device() == targets.device(),
                "Inputs must be on the same device");
    TORCH_CHECK(predictions.scalar_type() == at::kFloat &&
                targets.scalar_type() == at::kFloat,
                "Inputs must have dtype float32");

    const c10::cuda::CUDAGuard device_guard(predictions.device());
    auto expanded = at::broadcast_tensors({predictions, targets});
    const auto& p = expanded[0];
    const auto& t = expanded[1];

    TORCH_CHECK(p.dim() <= 64, "At most 64 dimensions are supported");

    BroadcastInfo info{};
    info.ndim = static_cast<int>(p.dim());
    for (int d = 0; d < info.ndim; ++d) {
        info.sizes[d] = p.size(d);
        info.pred_strides[d] = p.stride(d);
        info.target_strides[d] = t.stride(d);
    }

    const int64_t n = p.numel();
    const int blocks = static_cast<int>(
        std::max<int64_t>(1, std::min<int64_t>(1024, (n + 255) / 256)));
    auto output = torch::empty({}, predictions.options());
    auto partials = torch::empty({blocks}, predictions.options());
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (p.is_contiguous() && t.is_contiguous()) {
        hinge_partial_kernel<true><<<blocks, 256, 0, stream>>>(
            p.data_ptr<float>(), t.data_ptr<float>(),
            partials.data_ptr<float>(), n, info);
    } else {
        hinge_partial_kernel<false><<<blocks, 256, 0, stream>>>(
            p.data_ptr<float>(), t.data_ptr<float>(),
            partials.data_ptr<float>(), n, info);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    hinge_finish_kernel<<<1, 256, 0, stream>>>(
        partials.data_ptr<float>(), output.data_ptr<float>(), blocks, n);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return output;
}
"""

_hinge_extension = load_inline(
    name="fused_hinge_loss_cuda",
    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):
        return _hinge_extension.hinge_loss_cuda(predictions, targets)
```