```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 value = __fsub_rn(1.0f, __fmul_rn(p, t));
    return value < 0.0f ? 0.0f : value;
}

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

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

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

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

__device__ __forceinline__ void write_partial(
    float value, float* output, int64_t count
) {
    value = block_sum(value);
    if (threadIdx.x == 0) {
        const int index = blockIdx.y * gridDim.x + blockIdx.x;
        output[index] = gridDim.x * gridDim.y == 1
            ? (count ? value / static_cast<float>(count) : nanf(""))
            : value;
    }
}

template<bool Scalar>
__device__ __forceinline__ float4 read4(
    const float* pointer, int64_t offset, float scalar
) {
    if (Scalar)
        return make_float4(scalar, scalar, scalar, scalar);
    return *reinterpret_cast<const float4*>(pointer + offset);
}

// The two-dimensional grid removes integer division from the input loop.
// Each iteration handles 4096 elements with independent accumulators.
template<bool PScalar, bool TScalar, bool Vectorized>
__global__ void rectangular_kernel(
    const float* __restrict__ predictions,
    const float* __restrict__ targets,
    float* __restrict__ output,
    int64_t rows,
    int64_t columns,
    int64_t prediction_row_stride,
    int64_t target_row_stride,
    int64_t count
) {
    float s0 = 0.0f;
    float s1 = 0.0f;
    float s2 = 0.0f;
    float s3 = 0.0f;

    constexpr int tile_size = Vectorized ? 4096 : 1024;
    const int64_t step = static_cast<int64_t>(gridDim.x) * tile_size;

    for (int64_t row = blockIdx.y; row < rows; row += gridDim.y) {
        const float* p = predictions + row * prediction_row_stride;
        const float* t = targets + row * target_row_stride;
        const float pv_scalar = PScalar ? p[0] : 0.0f;
        const float tv_scalar = TScalar ? t[0] : 0.0f;

        if (Vectorized) {
            for (int64_t base = static_cast<int64_t>(blockIdx.x) * 4096
                                + threadIdx.x * 4;
                 base < columns; base += step) {
                #pragma unroll
                for (int j = 0; j < 4; ++j) {
                    const int64_t col = base + j * 1024;
                    if (col < columns) {
                        const float4 p4 = read4<PScalar>(p, col, pv_scalar);
                        const float4 t4 = read4<TScalar>(t, col, tv_scalar);
                        s0 += hinge(p4.x, t4.x);
                        s1 += hinge(p4.y, t4.y);
                        s2 += hinge(p4.z, t4.z);
                        s3 += hinge(p4.w, t4.w);
                    }
                }
            }
        } else {
            for (int64_t base = static_cast<int64_t>(blockIdx.x) * 1024
                                + threadIdx.x;
                 base < columns; base += step) {
                if (base < columns)
                    s0 += hinge(PScalar ? pv_scalar : p[base],
                                TScalar ? tv_scalar : t[base]);
                if (base + 256 < columns)
                    s1 += hinge(PScalar ? pv_scalar : p[base + 256],
                                TScalar ? tv_scalar : t[base + 256]);
                if (base + 512 < columns)
                    s2 += hinge(PScalar ? pv_scalar : p[base + 512],
                                TScalar ? tv_scalar : t[base + 512]);
                if (base + 768 < columns)
                    s3 += hinge(PScalar ? pv_scalar : p[base + 768],
                                TScalar ? tv_scalar : t[base + 768]);
            }
        }
    }

    write_partial((s0 + s1) + (s2 + s3), output, count);
}

__global__ void outer_kernel(
    const float* __restrict__ rows,
    const float* __restrict__ columns,
    float* __restrict__ output,
    int64_t m,
    int64_t n,
    int64_t row_stride,
    int64_t column_stride
) {
    float sums[4] = {0.0f, 0.0f, 0.0f, 0.0f};

    for (int64_t rb = static_cast<int64_t>(blockIdx.y) * 8;
         rb < m; rb += static_cast<int64_t>(gridDim.y) * 8) {
        for (int64_t cb = static_cast<int64_t>(blockIdx.x) * 1024
                          + threadIdx.x;
             cb < n; cb += static_cast<int64_t>(gridDim.x) * 1024) {
            float tv[4];
            #pragma unroll
            for (int j = 0; j < 4; ++j) {
                const int64_t col = cb + j * 256;
                tv[j] = col < n ? columns[col * column_stride] : 0.0f;
            }

            #pragma unroll
            for (int r = 0; r < 8; ++r) {
                if (rb + r < m) {
                    const float pv = rows[(rb + r) * row_stride];
                    #pragma unroll
                    for (int j = 0; j < 4; ++j) {
                        if (cb + j * 256 < n)
                            sums[j] += hinge(pv, tv[j]);
                    }
                }
            }
        }
    }

    write_partial(
        (sums[0] + sums[1]) + (sums[2] + sums[3]), output, m * n);
}

template<typename Index>
__global__ void broadcast_kernel(
    const float* __restrict__ predictions,
    const float* __restrict__ targets,
    float* __restrict__ output,
    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 remaining = static_cast<Index>(i);
        int64_t po = 0;
        int64_t to = 0;
        for (int d = info.ndim - 1; d >= 0; --d) {
            const Index size = static_cast<Index>(info.sizes[d]);
            const Index coordinate = remaining % size;
            remaining /= size;
            po += static_cast<int64_t>(coordinate) * info.ps[d];
            to += static_cast<int64_t>(coordinate) * info.ts[d];
        }
        sum += hinge(predictions[po], targets[to]);
    }

    write_partial(sum, output, count);
}

__global__ void finish_kernel(
    const float* __restrict__ partials,
    float* __restrict__ output,
    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)
        output[0] = sum / static_cast<float>(count);
}

template<bool Vectorized>
void launch_rectangular(
    const float* p,
    const float* t,
    float* output,
    int64_t rows,
    int64_t columns,
    int64_t prs,
    int64_t trs,
    int64_t count,
    bool pscalar,
    bool tscalar,
    dim3 grid,
    cudaStream_t stream
) {
    if (pscalar && tscalar) {
        rectangular_kernel<true, true, Vectorized>
            <<<grid, 256, 0, stream>>>(
                p, t, output, rows, columns, prs, trs, count);
    } else if (pscalar) {
        rectangular_kernel<true, false, Vectorized>
            <<<grid, 256, 0, stream>>>(
                p, t, output, rows, columns, prs, trs, count);
    } else if (tscalar) {
        rectangular_kernel<false, true, Vectorized>
            <<<grid, 256, 0, stream>>>(
                p, t, output, rows, columns, prs, trs, count);
    } else {
        rectangular_kernel<false, false, Vectorized>
            <<<grid, 256, 0, stream>>>(
                p, t, output, rows, columns, prs, trs, 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 previous = info.ndim - 1;
        if (previous >= 0 &&
            info.ps[previous] == p.stride(d) * p.size(d) &&
            info.ts[previous] == t.stride(d) * p.size(d)) {
            info.sizes[previous] *= p.size(d);
            info.ps[previous] = p.stride(d);
            info.ts[previous] = t.stride(d);
        } else {
            const int index = info.ndim++;
            info.sizes[index] = p.size(d);
            info.ps[index] = p.stride(d);
            info.ts[index] = 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 columns = count;
    int64_t prs = 0;
    int64_t trs = 0;
    bool pcs = pscalar;
    bool tcs = tscalar;

    if (!linear && rect2d) {
        rows = info.sizes[0];
        columns = info.sizes[1];
        prs = info.ps[0];
        trs = info.ts[0];
        pcs = info.ps[1] == 0;
        tcs = info.ts[1] == 0;
    }

    const float* pp = p.data_ptr<float>();
    const float* tp = t.data_ptr<float>();
    const bool vectorized = rectangular && columns % 4 == 0 &&
        (pcs || (reinterpret_cast<uintptr_t>(pp) % 16 == 0 &&
                 prs % 4 == 0)) &&
        (tcs || (reinterpret_cast<uintptr_t>(tp) % 16 == 0 &&
                 trs % 4 == 0));

    dim3 grid(1, 1);
    if (outer) {
        grid.x = static_cast<unsigned int>(std::min<int64_t>(
            2048, (info.sizes[1] + 1023) / 1024));
        grid.y = static_cast<unsigned int>(std::min<int64_t>(
            (info.sizes[0] + 7) / 8,
            std::max<int64_t>(1, 2048 / grid.x)));
    } else if (rectangular) {
        const int tile = vectorized ? 4096 : 1024;
        grid.x = static_cast<unsigned int>(std::min<int64_t>(
            2048, (columns + tile - 1) / tile));
        grid.y = static_cast<unsigned int>(std::min<int64_t>(
            rows, std::max<int64_t>(1, 2048 / grid.x)));
    } else {
        grid.x = static_cast<unsigned int>(std::max<int64_t>(
            1, std::min<int64_t>(2048, (count + 1023) / 1024)));
    }

    const int blocks = grid.x * grid.y;
    auto output = torch::empty({}, predictions.options());
    auto partials = blocks == 1
        ? output : torch::empty({blocks}, predictions.options());
    float* destination = partials.data_ptr<float>();
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (outer) {
        outer_kernel<<<grid, 256, 0, stream>>>(
            outer_pt ? pp : tp,
            outer_pt ? tp : pp,
            destination,
            info.sizes[0],
            info.sizes[1],
            outer_pt ? info.ps[0] : info.ts[0],
            outer_pt ? info.ts[1] : info.ps[1]);
    } else if (rectangular) {
        if (vectorized) {
            launch_rectangular<true>(
                pp, tp, destination, rows, columns, prs, trs,
                count, pcs, tcs, grid, stream);
        } else {
            launch_rectangular<false>(
                pp, tp, destination, rows, columns, prs, trs,
                count, pcs, tcs, grid, stream);
        }
    } else if (count <= std::numeric_limits<int32_t>::max()) {
        broadcast_kernel<int32_t><<<grid, 256, 0, stream>>>(
            pp, tp, destination, count, info);
    } else {
        broadcast_kernel<int64_t><<<grid, 256, 0, stream>>>(
            pp, tp, destination, count, info);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();

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

    return output;
}
"""

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