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

_cpp_source = """
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>

struct BroadcastLayout {
    int ndim;
    int64_t sizes[64];
    int64_t prediction_strides[64];
    int64_t target_strides[64];
};

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

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

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

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

    __syncthreads();

    value = threadIdx.x < 8 ? warp_sums[lane] : 0.0f;

    if (warp == 0) {
        for (int offset = 16; offset > 0; offset >>= 1)
            value += __shfl_down_sync(0xffffffff, value, offset);
    }
    return value;
}

template<bool P_CONTIGUOUS, bool T_CONTIGUOUS>
__global__ void hinge_partial_kernel(
    const float* __restrict__ predictions,
    const float* __restrict__ targets,
    float* __restrict__ result,
    int64_t count,
    BroadcastLayout layout,
    bool final_output
) {
    float sum = 0.0f;
    const int64_t step = static_cast<int64_t>(gridDim.x) * blockDim.x;

    for (int64_t index =
             static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
         index < count;
         index += step) {
        int64_t p_offset = P_CONTIGUOUS ? index : 0;
        int64_t t_offset = T_CONTIGUOUS ? index : 0;

        if (!P_CONTIGUOUS || !T_CONTIGUOUS) {
            int64_t remaining = index;
            for (int d = layout.ndim - 1; d >= 0; --d) {
                const int64_t coordinate = remaining % layout.sizes[d];
                remaining /= layout.sizes[d];
                if (!P_CONTIGUOUS)
                    p_offset += coordinate * layout.prediction_strides[d];
                if (!T_CONTIGUOUS)
                    t_offset += coordinate * layout.target_strides[d];
            }
        }

        const float margin =
            1.0f - predictions[p_offset] * targets[t_offset];
        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 hinge_finish_kernel(
    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() == at::kFloat &&
                targets.scalar_type() == at::kFloat,
                "Inputs must have dtype float32");
    TORCH_CHECK(predictions.sizes() == targets.sizes(),
                "Inputs must already be broadcast to the same shape");
    TORCH_CHECK(predictions.dim() <= 64, "Unsupported tensor rank");

    const c10::cuda::CUDAGuard device_guard(predictions.device());
    const auto stream =
        at::cuda::getCurrentCUDAStream(predictions.get_device());

    BroadcastLayout layout{};
    layout.ndim = predictions.dim();
    for (int d = 0; d < layout.ndim; ++d) {
        layout.sizes[d] = predictions.size(d);
        layout.prediction_strides[d] = predictions.stride(d);
        layout.target_strides[d] = targets.stride(d);
    }

    const int64_t count = predictions.numel();
    const int blocks = static_cast<int>(
        std::min<int64_t>(1024, std::max<int64_t>(1, (count + 4095) / 4096))
    );

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

    const bool pc = predictions.is_contiguous();
    const bool tc = targets.is_contiguous();

    #define LAUNCH(P, T)                                                   \
        hinge_partial_kernel<P, T><<<blocks, 256, 0, stream.stream()>>>(   \
            predictions.data_ptr<float>(), targets.data_ptr<float>(),    \
            partials.data_ptr<float>(), count, layout, blocks == 1)

    if (pc && tc) {
        LAUNCH(true, true);
    } else if (pc) {
        LAUNCH(true, false);
    } else if (tc) {
        LAUNCH(false, true);
    } else {
        LAUNCH(false, false);
    }

    #undef LAUNCH
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    if (blocks > 1) {
        hinge_finish_kernel<<<1, 256, 0, stream.stream()>>>(
            partials.data_ptr<float>(), output.data_ptr<float>(),
            blocks, count
        );
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }

    return output;
}
"""

_hinge_cuda = 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):
        predictions, targets = torch.broadcast_tensors(predictions, targets)
        return _hinge_cuda.hinge_loss_cuda(predictions, targets)
```