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

_matvec_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>
#include <cstdint>
#include <limits>

namespace {
constexpr int kBlockSize = 256;

__global__ void __launch_bounds__(kBlockSize) matvec_kernel(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float* __restrict__ C,
    int M,
    int K) {
  const int row = blockIdx.x;
  const float* a = A + static_cast<int64_t>(row) * K;
  const float4* a4 = reinterpret_cast<const float4*>(a);
  const float4* b4 = reinterpret_cast<const float4*>(B);

  const int n4 = K >> 2;
  float acc0 = 0.0f;
  float acc1 = 0.0f;
  float acc2 = 0.0f;
  float acc3 = 0.0f;

  int i = threadIdx.x;
  for (; i + 3 * kBlockSize < n4; i += 4 * kBlockSize) {
#pragma unroll
    for (int u = 0; u < 4; ++u) {
      const int idx = i + u * kBlockSize;
      const float4 av = __ldg(a4 + idx);
      const float4 bv = __ldg(b4 + idx);
      acc0 = fmaf(av.x, bv.x, acc0);
      acc1 = fmaf(av.y, bv.y, acc1);
      acc2 = fmaf(av.z, bv.z, acc2);
      acc3 = fmaf(av.w, bv.w, acc3);
    }
  }
  for (; i < n4; i += kBlockSize) {
    const float4 av = __ldg(a4 + i);
    const float4 bv = __ldg(b4 + i);
    acc0 = fmaf(av.x, bv.x, acc0);
    acc1 = fmaf(av.y, bv.y, acc1);
    acc2 = fmaf(av.z, bv.z, acc2);
    acc3 = fmaf(av.w, bv.w, acc3);
  }

  float sum = (acc0 + acc1) + (acc2 + acc3);

#pragma unroll
  for (int offset = 16; offset > 0; offset >>= 1) {
    sum += __shfl_down_sync(0xffffffffu, sum, offset);
  }

  __shared__ float warp_sums[kBlockSize / 32];
  const int lane = threadIdx.x & 31;
  const int warp = threadIdx.x >> 5;
  if (lane == 0) {
    warp_sums[warp] = sum;
  }
  __syncthreads();

  if (warp == 0) {
    constexpr int kWarps = kBlockSize / 32;
    sum = (lane < kWarps) ? warp_sums[lane] : 0.0f;
#pragma unroll
    for (int offset = kWarps / 2; offset > 0; offset >>= 1) {
      sum += __shfl_down_sync(0xffffffffu, sum, offset);
    }
    if (lane == 0) {
      C[row] = sum;
    }
  }
}
}  // namespace

torch::Tensor matvec_cuda(torch::Tensor A, torch::Tensor B) {
  TORCH_CHECK(A.is_cuda() && B.is_cuda(), "A and B must be CUDA tensors");
  TORCH_CHECK(A.scalar_type() == at::kFloat && B.scalar_type() == at::kFloat,
              "A and B must be float32");
  TORCH_CHECK(A.dim() == 2 && B.dim() == 2 && B.size(1) == 1,
              "Expected A with shape (M, K) and B with shape (K, 1)");
  TORCH_CHECK(A.size(1) == B.size(0), "K dimension mismatch");
  TORCH_CHECK(A.is_contiguous() && B.is_contiguous(), "A and B must be contiguous");

  const int64_t M64 = A.size(0);
  const int64_t K64 = A.size(1);
  TORCH_CHECK(M64 <= std::numeric_limits<int>::max() &&
                  K64 <= std::numeric_limits<int>::max(),
              "Matrix is too large for this kernel");
  TORCH_CHECK((K64 & 3) == 0, "K must be divisible by 4");
  TORCH_CHECK((reinterpret_cast<std::uintptr_t>(A.data_ptr<float>()) & 15) == 0 &&
                  (reinterpret_cast<std::uintptr_t>(B.data_ptr<float>()) & 15) == 0,
              "A and B must be 16-byte aligned");

  auto C = torch::empty({M64, 1}, A.options());
  if (M64 == 0) {
    return C;
  }
  if (K64 == 0) {
    C.zero_();
    return C;
  }

  const int M = static_cast<int>(M64);
  const int K = static_cast<int>(K64);
  cudaStream_t stream = at::cuda::getCurrentCUDAStream();
  matvec_kernel<<<M, kBlockSize, 0, stream>>>(
      A.data_ptr<float>(), B.data_ptr<float>(), C.data_ptr<float>(), M, K);
  cudaError_t err = cudaGetLastError();
  TORCH_CHECK(err == cudaSuccess, "matvec kernel launch failed: ",
              cudaGetErrorString(err));
  return C;
}
"""

_matvec_cpp_src = "torch::Tensor matvec_cuda(torch::Tensor A, torch::Tensor B);"

_matvec_ext = load_inline(
    name="custom_matvec_gemv",
    cpp_sources=_matvec_cpp_src,
    cuda_sources=_matvec_cuda_src,
    functions=["matvec_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


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

    def forward(self, A: torch.Tensor, B: torch.Tensor) -> torch.Tensor:
        if (
            A.is_cuda
            and B.is_cuda
            and A.dtype == torch.float32
            and B.dtype == torch.float32
            and A.dim() == 2
            and B.dim() == 2
            and B.size(1) == 1
            and A.size(1) == B.size(0)
            and A.is_contiguous()
            and B.is_contiguous()
            and (A.size(1) & 3) == 0
            and (A.data_ptr() & 15) == 0
            and (B.data_ptr() & 15) == 0
        ):
            return self._matvec.matvec_cuda(A, B)
        return torch.matmul(A, B)
```