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

__device__ __forceinline__ float4 streaming_load(const float* p) {
    float4 v;
    asm volatile(
        "ld.global.cg.v4.f32 {%0, %1, %2, %3}, [%4];"
        : "=f"(v.x), "=f"(v.y), "=f"(v.z), "=f"(v.w)
        : "l"(p));
    return v;
}

__device__ __forceinline__ void accumulate(
    float4 p, float4 t, float& a, float& b, float& c, float& d
) {
    a += hinge(p.x, t.x);
    b += hinge(p.y, t.y);
    c += hinge(p.z, t.z);
    d += hinge(p.w, t.w);
}

__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 write_partial(
    float value, float* output, int64_t count
) {
    value = block_sum(value);
    if (threadIdx.x == 0) {
        const int block = blockIdx.y * gridDim.x + blockIdx.x;
        output[block] = gridDim.x * gridDim.y == 1
            ? (count ? value / static_cast<float>(count) : nanf(""))
            : value;
    }
}

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

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 prs,
    int64_t trs,
    int64_t count
) {
    float a = 0.0f, b = 0.0f, c = 0.0f, d = 0.0f;
    constexpr int tile = Vectorized ? 4096 : 1024;
    const int64_t step = static_cast<int64_t>(gridDim.x) * tile;

    for (int64_t row = blockIdx.y; row < rows; row += gridDim.y) {
        const float* p = predictions + row * prs;
        const float* t = targets + row * trs;
        const float pv = PScalar ? p[0] : 0.0f;
        const float tv = TScalar ? t[0] : 0.0f;

        if (Vectorized) {
            for (int64_t start = static_cast<int64_t>(blockIdx.x) * tile;
                 start < columns; start += step) {
                const int64_t base = start + threadIdx.x * 4;
                #pragma unroll
                for (int j = 0; j < 4; ++j) {
                    const int64_t col = base + j * 1024;
                    if (col < columns)
                        accumulate(
                            read4<PScalar>(p, col, pv),
                            read4<TScalar>(t, col, tv), a, b, c, d);
                }
            }
        } else {
            for (int64_t base = static_cast<int64_t>(blockIdx.x) * tile
                                + threadIdx.x;
                 base < columns; base += step) {
                a += hinge(PScalar ? pv : p[base],
                           TScalar ? tv : t[base]);
                if (base + 256 < columns)
                    b += hinge(PScalar ? pv : p[base + 256],
                               TScalar ? tv : t[base + 256]);
                if (base + 512 < columns)
                    c += hinge(PScalar ? pv : p[base + 512],
                               TScalar ? tv : t[base + 512]);
                if (base + 768 < columns)
                    d += hinge(PScalar ? pv : p[base + 768],
                               TScalar ? tv : t[base + 768]);
            }
        }
    }
    write_partial((a + b) + (c + d), output, count);
}

// Wider contiguous matrix tiles, with four target vectors retained per thread.
template<bool FullColumns>
__global__ __launch_bounds__(256, 2) void reused_row_kernel(
    const float* __restrict__ matrix,
    const float* __restrict__ vector,
    float* __restrict__ output,
    int64_t rows,
    int64_t columns,
    int64_t row_stride
) {
    float a = 0.0f, b = 0.0f, c = 0.0f, d = 0.0f;
    const int64_t column_step = static_cast<int64_t>(gridDim.x) * 4096;
    const int64_t row_step = static_cast<int64_t>(gridDim.y) * 4;

    for (int64_t tile = static_cast<int64_t>(blockIdx.x) * 4096;
         tile < columns; tile += column_step) {
        const int64_t base = tile + threadIdx.x * 4;
        float4 tv[4];

        #pragma unroll
        for (int j = 0; j < 4; ++j) {
            const int64_t col = base + j * 1024;
            if (FullColumns || col < columns)
                tv[j] = __ldg(
                    reinterpret_cast<const float4*>(vector + col));
            else
                tv[j] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
        }

        for (int64_t rb = static_cast<int64_t>(blockIdx.y) * 4;
             rb < rows; rb += row_step) {
            if (rb + 4 <= rows) {
                #pragma unroll
                for (int r = 0; r < 4; ++r) {
                    const float* p = matrix + (rb + r) * row_stride + base;
                    #pragma unroll
                    for (int j = 0; j < 4; ++j) {
                        if (FullColumns || base + j * 1024 < columns)
                            accumulate(
                                streaming_load(p + j * 1024),
                                tv[j], a, b, c, d);
                    }
                }
            } else {
                #pragma unroll
                for (int r = 0; r < 4; ++r) {
                    if (rb + r < rows) {
                        const float* p =
                            matrix + (rb + r) * row_stride + base;
                        #pragma unroll
                        for (int j = 0; j < 4; ++j) {
                            if (FullColumns || base + j * 1024 < columns)
                                accumulate(
                                    streaming_load(p + j * 1024),
                                    tv[j], a, b, c, d);
                        }
                    }
                }
            }
        }
    }
    write_partial((a + b) + (c + d), output, rows * columns);
}

__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, 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 a = 0.0f, b = 0.0f, c = 0.0f, d = 0.0f;
    for (int i = threadIdx.x; i < blocks; i += 1024) {
        a += partials[i];
        if (i + 256 < blocks) b += partials[i + 256];
        if (i + 512 < blocks) c += partials[i + 512];
        if (i + 768 < blocks) d += partials[i + 768];
    }
    const float sum = block_sum((a + b) + (c + d));
    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 ps, bool ts, dim3 grid, cudaStream_t stream
) {
    if (ps && ts)
        rectangular_kernel<true, true, Vectorized>
            <<<grid, 256, 0, stream>>>(
                p, t, output, rows, columns, prs, trs, count);
    else if (ps)
        rectangular_kernel<true, false, Vectorized>
            <<<grid, 256, 0, stream>>>(
                p, t, output, rows, columns, prs, trs, count);
    else if (ts)
        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());
    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    const bool ps = predictions.numel() == 1;
    const bool ts = targets.numel() == 1;
    const bool direct_linear =
        predictions.is_contiguous() && targets.is_contiguous() &&
        (predictions.sizes() == targets.sizes() || ps || ts);

    torch::Tensor p = predictions;
    torch::Tensor t = targets;
    BroadcastInfo info{};
    int64_t count;

    if (direct_linear) {
        count = ps ? targets.numel() : predictions.numel();
    } else {
        auto expanded = at::broadcast_tensors({predictions, targets});
        p = expanded[0];
        t = expanded[1];
        count = p.numel();
        TORCH_CHECK(p.dim() <= 64, "At most 64 dimensions are supported");

        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 linear = direct_linear ||
        ((p.is_contiguous() || ps) && (t.is_contiguous() || ts));

    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, columns = count, prs = 0, trs = 0;
    bool pcs = ps, tcs = ts;
    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));

    const bool reuse_t = vectorized && rows > 1 &&
        !pcs && !tcs && trs == 0 && prs != 0;
    const bool reuse_p = vectorized && rows > 1 &&
        !pcs && !tcs && prs == 0 && trs != 0;
    const bool reuse = reuse_t || reuse_p;

    constexpr int64_t max_blocks = 32768;
    dim3 grid(1, 1);
    if (reuse) {
        constexpr int64_t reuse_blocks = 8192;
        grid.x = static_cast<unsigned int>(
            std::min<int64_t>(reuse_blocks, (columns + 4095) / 4096));
        grid.y = static_cast<unsigned int>(std::min<int64_t>(
            (rows + 3) / 4,
            std::max<int64_t>(1, reuse_blocks / grid.x)));
    } else if (outer) {
        grid.x = static_cast<unsigned int>(
            std::min<int64_t>(max_blocks, (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, max_blocks / grid.x)));
    } else if (rectangular) {
        const int tile = vectorized ? 4096 : 1024;
        grid.x = static_cast<unsigned int>(
            std::min<int64_t>(max_blocks, (columns + tile - 1) / tile));
        grid.y = static_cast<unsigned int>(std::min<int64_t>(
            rows, std::max<int64_t>(1, max_blocks / 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>();

    if (reuse) {
        const float* matrix = reuse_t ? pp : tp;
        const float* vector = reuse_t ? tp : pp;
        const int64_t stride = reuse_t ? prs : trs;
        if (columns % 4096 == 0)
            reused_row_kernel<true><<<grid, 256, 0, stream>>>(
                matrix, vector, destination, rows, columns, stride);
        else
            reused_row_kernel<false><<<grid, 256, 0, stream>>>(
                matrix, vector, destination, rows, columns, stride);
    } else 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_wide_reuse_v6",
    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)
```