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


_cpp_source = r"""
#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_runtime.h>
#include <algorithm>
#include <cstdint>
#include <limits>

struct BroadcastInfo {
    int ndim;
    int64_t sizes[64];
    int64_t ps[64];
    int64_t ts[64];
};

__device__ __forceinline__ float hinge(float p, float t) {
    const float v = __fsub_rn(1.0f, __fmul_rn(p, t));
    return v < 0.0f ? 0.0f : v;
}

__device__ __forceinline__ float block_sum(float v) {
    __shared__ float scratch[8];
    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;

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

    if (lane == 0)
        scratch[warp] = v;
    __syncthreads();

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

__device__ __forceinline__ void store_sum(
    float sum, float* out, int64_t count
) {
    sum = block_sum(sum);
    if (threadIdx.x == 0) {
        out[blockIdx.x] = gridDim.x == 1
            ? (count ? sum / static_cast<float>(count) : nanf(""))
            : sum;
    }
}

template<bool Scalar>
__device__ __forceinline__ float4 load_four(const float* ptr, int64_t col) {
    if (Scalar) {
        const float x = ptr[0];
        return make_float4(x, x, x, x);
    } else {
        return *reinterpret_cast<const float4*>(ptr + col);
    }
}

// Each block processes a contiguous segment of a row. Broadcast offsets
// are calculated once per tile, and aligned inputs use 128-bit loads.
template<bool PScalar, bool TScalar, bool Vectorized>
__global__ void rectangular_hinge_kernel(
    const float* __restrict__ p,
    const float* __restrict__ t,
    float* __restrict__ out,
    int64_t cols,
    int64_t p_row_stride,
    int64_t t_row_stride,
    int64_t column_tiles,
    int64_t tiles,
    int64_t count
) {
    float s0 = 0.0f;
    float s1 = 0.0f;
    float s2 = 0.0f;
    float s3 = 0.0f;

    for (int64_t tile = blockIdx.x; tile < tiles; tile += gridDim.x) {
        const int64_t row = tile / column_tiles;
        const int64_t tile_col = tile - row * column_tiles;
        const float* pr = p + row * p_row_stride;
        const float* tr = t + row * t_row_stride;

        if (Vectorized) {
            const int64_t col = tile_col * 1024 + threadIdx.x * 4;
            if (col < cols) {
                const float4 pv = load_four<PScalar>(pr, col);
                const float4 tv = load_four<TScalar>(tr, col);
                s0 += hinge(pv.x, tv.x);
                s1 += hinge(pv.y, tv.y);
                s2 += hinge(pv.z, tv.z);
                s3 += hinge(pv.w, tv.w);
            }
        } else {
            const int64_t col = tile_col * 1024 + threadIdx.x;
            if (col < cols)
                s0 += hinge(pr[PScalar ? 0 : col],
                            tr[TScalar ? 0 : col]);
            if (col + 256 < cols)
                s1 += hinge(pr[PScalar ? 0 : col + 256],
                            tr[TScalar ? 0 : col + 256]);
            if (col + 512 < cols)
                s2 += hinge(pr[PScalar ? 0 : col + 512],
                            tr[TScalar ? 0 : col + 512]);
            if (col + 768 < cols)
                s3 += hinge(pr[PScalar ? 0 : col + 768],
                            tr[TScalar ? 0 : col + 768]);
        }
    }
    store_sum((s0 + s1) + (s2 + s3), out, count);
}

__global__ void outer_hinge_kernel(
    const float* __restrict__ rows,
    const float* __restrict__ cols,
    float* __restrict__ out,
    int64_t m,
    int64_t n,
    int64_t row_stride,
    int64_t col_stride,
    int64_t column_tiles,
    int64_t tile_count
) {
    float sums[4] = {0.0f, 0.0f, 0.0f, 0.0f};
    for (int64_t tile = blockIdx.x; tile < tile_count; tile += gridDim.x) {
        const int64_t row_base = (tile / column_tiles) * 8;
        const int64_t col_base =
            (tile % column_tiles) * 1024 + threadIdx.x;

        float tv[4];
        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            const int64_t col = col_base + j * 256;
            tv[j] = col < n ? cols[col * col_stride] : 0.0f;
        }

        #pragma unroll
        for (int r = 0; r < 8; ++r) {
            if (row_base + r < m) {
                const float pv = rows[(row_base + r) * row_stride];
                #pragma unroll
                for (int j = 0; j < 4; ++j) {
                    if (col_base + j * 256 < n)
                        sums[j] += hinge(pv, tv[j]);
                }
            }
        }
    }
    store_sum((sums[0] + sums[1]) + (sums[2] + sums[3]), out, m * n);
}

template<typename index_t>
__global__ void broadcast_hinge_kernel(
    const float* __restrict__ p,
    const float* __restrict__ t,
    float* __restrict__ out,
    int64_t count,
    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 < count; i += step) {
        index_t remaining = static_cast<index_t>(i);
        int64_t po = 0;
        int64_t to = 0;

        for (int d = info.ndim - 1; d >= 0; --d) {
            const index_t size = static_cast<index_t>(info.sizes[d]);
            const index_t coord = remaining % size;
            remaining /= size;
            po += static_cast<int64_t>(coord) * info.ps[d];
            to += static_cast<int64_t>(coord) * info.ts[d];
        }
        sum += hinge(p[po], t[to]);
    }
    store_sum(sum, out, count);
}

__global__ void finish_kernel(
    const float* __restrict__ partials,
    float* __restrict__ out,
    int blocks,
    int64_t count
) {
    float sum = 0.0f;
    for (int i = threadIdx.x; i < blocks; i += 256)
        sum += partials[i];
    sum = block_sum(sum);
    if (threadIdx.x == 0)
        out[0] = sum / static_cast<float>(count);
}

template<bool Vectorized>
void launch_rectangular(
    const float* p, const float* t, float* out,
    int64_t cols, int64_t prs, int64_t trs,
    int64_t column_tiles, int64_t tiles, int64_t count,
    bool pscalar, bool tscalar, int blocks, cudaStream_t stream
) {
    if (pscalar && tscalar) {
        rectangular_hinge_kernel<true, true, Vectorized>
            <<<blocks, 256, 0, stream>>>(
                p, t, out, cols, prs, trs, column_tiles, tiles, count);
    } else if (pscalar) {
        rectangular_hinge_kernel<true, false, Vectorized>
            <<<blocks, 256, 0, stream>>>(
                p, t, out, cols, prs, trs, column_tiles, tiles, count);
    } else if (tscalar) {
        rectangular_hinge_kernel<false, true, Vectorized>
            <<<blocks, 256, 0, stream>>>(
                p, t, out, cols, prs, trs, column_tiles, tiles, count);
    } else {
        rectangular_hinge_kernel<false, false, Vectorized>
            <<<blocks, 256, 0, stream>>>(
                p, t, out, cols, prs, trs, column_tiles, tiles, 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");

    const c10::cuda::CUDAGuard 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");

    const int64_t count = p.numel();
    BroadcastInfo info{};
    for (int d = 0; d < p.dim(); ++d) {
        if (p.size(d) == 1)
            continue;
        const int k = info.ndim - 1;
        if (k >= 0 &&
            info.ps[k] == p.stride(d) * p.size(d) &&
            info.ts[k] == t.stride(d) * p.size(d)) {
            info.sizes[k] *= p.size(d);
            info.ps[k] = p.stride(d);
            info.ts[k] = t.stride(d);
        } else {
            const int j = info.ndim++;
            info.sizes[j] = p.size(d);
            info.ps[j] = p.stride(d);
            info.ts[j] = t.stride(d);
        }
    }

    const bool pscalar = predictions.numel() == 1;
    const bool tscalar = targets.numel() == 1;
    const bool linear = (p.is_contiguous() || pscalar) &&
                        (t.is_contiguous() || tscalar);
    const bool outer_pt = info.ndim == 2 &&
                          info.ps[1] == 0 && info.ts[0] == 0;
    const bool outer_tp = info.ndim == 2 &&
                          info.ts[1] == 0 && info.ps[0] == 0;
    const bool outer = count > 0 && (outer_pt || outer_tp);

    const bool rect2d = count > 0 && info.ndim == 2 &&
                       info.sizes[1] >= 128 &&
                       (info.ps[1] == 0 || info.ps[1] == 1) &&
                       (info.ts[1] == 0 || info.ts[1] == 1);
    const bool rectangular = count > 0 && (linear || rect2d);

    int64_t rows = 1;
    int64_t cols = count;
    int64_t prs = 0, trs = 0;
    bool pcs = pscalar, tcs = tscalar;
    if (!linear && rect2d) {
        rows = info.sizes[0];
        cols = info.sizes[1];
        prs = info.ps[0];
        trs = info.ts[0];
        pcs = info.ps[1] == 0;
        tcs = info.ts[1] == 0;
    }

    int64_t column_tiles = 0;
    int64_t tiles = 0;
    if (outer) {
        column_tiles = (info.sizes[1] + 1023) / 1024;
        tiles = ((info.sizes[0] + 7) / 8) * column_tiles;
    } else if (rectangular) {
        column_tiles = (cols + 1023) / 1024;
        tiles = rows * column_tiles;
    } else {
        tiles = (count + 1023) / 1024;
    }
    const int blocks = static_cast<int>(
        std::max<int64_t>(1, std::min<int64_t>(4096, tiles)));

    auto output = torch::empty({}, predictions.options());
    auto partials = blocks == 1
        ? output
        : torch::empty({blocks}, predictions.options());
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const float* pp = p.data_ptr<float>();
    const float* tp = t.data_ptr<float>();
    float* dst = partials.data_ptr<float>();

    if (outer) {
        outer_hinge_kernel<<<blocks, 256, 0, stream>>>(
            outer_pt ? pp : tp, outer_pt ? tp : pp, dst,
            info.sizes[0], info.sizes[1],
            outer_pt ? info.ps[0] : info.ts[0],
            outer_pt ? info.ts[1] : info.ps[1],
            column_tiles, tiles);
    } else if (rectangular) {
        const bool vectorized = (cols % 4 == 0) &&
            (pcs || ((reinterpret_cast<uintptr_t>(pp) % 16 == 0) &&
                     (prs % 4 == 0))) &&
            (tcs || ((reinterpret_cast<uintptr_t>(tp) % 16 == 0) &&
                     (trs % 4 == 0)));
        if (vectorized) {
            launch_rectangular<true>(
                pp, tp, dst, cols, prs, trs, column_tiles, tiles, count,
                pcs, tcs, blocks, stream);
        } else {
            launch_rectangular<false>(
                pp, tp, dst, cols, prs, trs, column_tiles, tiles, count,
                pcs, tcs, blocks, stream);
        }
    } else if (count <= std::numeric_limits<int32_t>::max()) {
        broadcast_hinge_kernel<int32_t><<<blocks, 256, 0, stream>>>(
            pp, tp, dst, count, info);
    } else {
        broadcast_hinge_kernel<int64_t><<<blocks, 256, 0, stream>>>(
            pp, tp, dst, count, info);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();

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

_hinge_extension = load_inline(
    name="vectorized_broadcast_hinge_loss_v3",
    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)
```