```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 streaming_gemv_cuda(torch::Tensor A, torch::Tensor B);
"""

_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>

constexpr int THREADS = 256;
constexpr int ROWS = 4;
constexpr int CHUNK = 32768;

__device__ __forceinline__ float warp_sum(float value) {
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1)
        value += __shfl_down_sync(0xffffffffu, value, offset);
    return value;
}

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

template<bool FULL_ROWS, bool FULL_CHUNKS>
__global__ __launch_bounds__(THREADS)
void streaming_gemv_partial(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ partial,
    int M,
    int K,
    int splits
) {
    // Adjacent blocks traverse adjacent column tiles, avoiding the
    // power-of-two row-stride access pattern across neighboring blocks.
    const int split = blockIdx.x % splits;
    const int row = (blockIdx.x / splits) * ROWS;
    const int begin = split * CHUNK;
    const int remaining = K - begin;
    const int count = FULL_CHUNKS ? CHUNK
                                : (remaining < CHUNK ? remaining : CHUNK);

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

    const float* matrix = A + static_cast<int64_t>(row) * K + begin;
    const float* vector = B + begin;

    float4 acc[ROWS];
    #pragma unroll
    for (int r = 0; r < ROWS; ++r)
        acc[r] = make_float4(0.f, 0.f, 0.f, 0.f);

    #pragma unroll 4
    for (int offset = tid * 4; offset < count; offset += THREADS * 4) {
        const float4 b = __ldg(
            reinterpret_cast<const float4*>(vector + offset)
        );

        #pragma unroll
        for (int r = 0; r < ROWS; ++r) {
            if (FULL_ROWS || row + r < M) {
                const float4 a = streaming_load(
                    matrix + static_cast<int64_t>(r) * K + offset
                );
                acc[r].x = fmaf(a.x, b.x, acc[r].x);
                acc[r].y = fmaf(a.y, b.y, acc[r].y);
                acc[r].z = fmaf(a.z, b.z, acc[r].z);
                acc[r].w = fmaf(a.w, b.w, acc[r].w);
            }
        }
    }

    __shared__ float warp_sums[ROWS][THREADS / 32];

    #pragma unroll
    for (int r = 0; r < ROWS; ++r) {
        float sum = (acc[r].x + acc[r].y) + (acc[r].z + acc[r].w);
        sum = warp_sum(sum);
        if (lane == 0)
            warp_sums[r][warp] = sum;
    }
    __syncthreads();

    if (warp < ROWS) {
        float sum = lane < THREADS / 32 ? warp_sums[warp][lane] : 0.f;
        sum = warp_sum(sum);
        if (lane == 0 && (FULL_ROWS || row + warp < M))
            partial[static_cast<int64_t>(row + warp) * splits + split] = sum;
    }
}

__global__ void streaming_gemv_finish(
    const float* __restrict__ partial,
    float* __restrict__ output,
    int M,
    int splits
) {
    const int lane = threadIdx.x & 31;
    const int row = blockIdx.x * (THREADS / 32) + (threadIdx.x >> 5);
    if (row >= M)
        return;

    const float* values = partial + static_cast<int64_t>(row) * splits;
    float sum = 0.f;
    for (int s = lane; s < splits; s += 32)
        sum += values[s];

    sum = warp_sum(sum);
    if (lane == 0)
        output[row] = sum;
}

torch::Tensor streaming_gemv_cuda(torch::Tensor A, torch::Tensor B) {
    const c10::cuda::CUDAGuard guard(A.device());
    const int M = static_cast<int>(A.size(0));
    const int K = static_cast<int>(A.size(1));
    const int splits = (K + CHUNK - 1) / CHUNK;
    const int64_t blocks = ((static_cast<int64_t>(M) + ROWS - 1) / ROWS)
                         * splits;

    auto output = torch::empty({M, 1}, A.options());
    auto partial = splits == 1
        ? output
        : torch::empty({M, splits}, A.options());

    const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    const float* a = A.data_ptr<float>();
    const float* b = B.data_ptr<float>();
    float* p = partial.data_ptr<float>();

    if (M % ROWS == 0) {
        if (K % CHUNK == 0) {
            streaming_gemv_partial<true, true>
                <<<static_cast<unsigned int>(blocks), THREADS, 0, stream>>>(
                    a, b, p, M, K, splits);
        } else {
            streaming_gemv_partial<true, false>
                <<<static_cast<unsigned int>(blocks), THREADS, 0, stream>>>(
                    a, b, p, M, K, splits);
        }
    } else {
        if (K % CHUNK == 0) {
            streaming_gemv_partial<false, true>
                <<<static_cast<unsigned int>(blocks), THREADS, 0, stream>>>(
                    a, b, p, M, K, splits);
        } else {
            streaming_gemv_partial<false, false>
                <<<static_cast<unsigned int>(blocks), THREADS, 0, stream>>>(
                    a, b, p, M, K, splits);
        }
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    if (splits > 1) {
        streaming_gemv_finish<<<
            (static_cast<int64_t>(M) + THREADS / 32 - 1) / (THREADS / 32),
            THREADS, 0, stream
        >>>(p, output.data_ptr<float>(), M, splits);
        C10_CUDA_KERNEL_LAUNCH_CHECK();
    }
    return output;
}
"""

_gemv = load_inline(
    name="streaming_gemv_contiguous_tiles_v3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["streaming_gemv_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self):
        super().__init__()

    def forward(self, A: torch.Tensor, B: torch.Tensor) -> torch.Tensor:
        if (
            A.is_cuda
            and B.is_cuda
            and A.device == B.device
            and A.dtype == torch.float32
            and B.dtype == torch.float32
            and A.ndim == 2
            and B.ndim == 2
            and B.shape == (A.shape[1], 1)
            and A.is_contiguous()
            and B.is_contiguous()
            and 0 < A.shape[0] < 2**31
            and 0 < A.shape[1] <= 2**30
            and A.shape[1] % 4 == 0
            and ((A.shape[0] + 3) // 4)
                * ((A.shape[1] + 32767) // 32768) < 2**31
            and A.data_ptr() % 16 == 0
            and B.data_ptr() % 16 == 0
            and not (
                torch.is_grad_enabled() and (A.requires_grad or B.requires_grad)
            )
        ):
            return _gemv.streaming_gemv_cuda(A, B)
        return torch.matmul(A, B)
```