```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 pool_sum_scale_cuda(
    torch::Tensor x, int64_t kernel_size, double scale);
"""

_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 <math_constants.h>

__device__ __forceinline__ float maximum_nan(float a, float b) {
    return isnan(b) || a < b ? b : a;
}

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

template <bool VECTOR_PAIRS>
__global__ void pool_sum_scale_kernel(
    const float* __restrict__ input,
    float* __restrict__ output,
    int64_t columns,
    int64_t kernel_size,
    float scale
) {
    const int tid = threadIdx.x;
    const int64_t row = blockIdx.x;
    const float* row_input = input + row * columns;

    float sum0 = 0.0f;
    float sum1 = 0.0f;

    if constexpr (VECTOR_PAIRS) {
        const float4* vectors =
            reinterpret_cast<const float4*>(row_input);
        const int64_t count = columns / 4;

        for (int64_t i = tid; i < count; i += blockDim.x) {
            const float4 values = vectors[i];
            sum0 += maximum_nan(values.x, values.y);
            sum1 += maximum_nan(values.z, values.w);
        }
    } else {
        const int64_t pooled_columns = columns / kernel_size;
        for (int64_t i = tid; i < pooled_columns; i += blockDim.x) {
            const float* window = row_input + i * kernel_size;
            float maximum = -CUDART_INF_F;
            for (int64_t j = 0; j < kernel_size; ++j) {
                maximum = maximum_nan(maximum, window[j]);
            }
            sum0 += maximum;
        }
    }

    float sum = warp_sum(sum0 + sum1);

    __shared__ float warp_sums[8];
    const int lane = tid & 31;
    const int warp = tid >> 5;

    if (lane == 0) {
        warp_sums[warp] = sum;
    }
    __syncthreads();

    if (warp == 0) {
        sum = lane < 8 ? warp_sums[lane] : 0.0f;
        sum = warp_sum(sum);
        if (lane == 0) {
            output[row] = sum * scale;
        }
    }
}

torch::Tensor pool_sum_scale_cuda(
    torch::Tensor x, int64_t kernel_size, double scale
) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(x.dim() == 2 && x.is_contiguous(),
                "x must be a contiguous matrix");
    TORCH_CHECK(kernel_size > 0 && x.size(1) >= kernel_size,
                "invalid pooling kernel size");

    const c10::cuda::CUDAGuard device_guard(x.device());
    auto output = torch::empty({x.size(0)}, x.options());

    if (x.size(0) == 0) {
        return output;
    }

    const auto stream = at::cuda::getCurrentCUDAStream();
    const dim3 blocks(static_cast<unsigned int>(x.size(0)));
    constexpr int threads = 256;

    if (kernel_size == 2 && x.size(1) % 4 == 0 &&
        reinterpret_cast<uintptr_t>(x.data_ptr<float>()) % 16 == 0) {
        pool_sum_scale_kernel<true><<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), output.data_ptr<float>(),
            x.size(1), kernel_size, static_cast<float>(scale));
    } else {
        pool_sum_scale_kernel<false><<<blocks, threads, 0, stream>>>(
            x.data_ptr<float>(), output.data_ptr<float>(),
            x.size(1), kernel_size, static_cast<float>(scale));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return output;
}
"""

_fused_ops = load_inline(
    name="fused_maxpool_sum_scale_cuda_v2",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["pool_sum_scale_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class _PoolSumScale(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, kernel_size, scale):
        ctx.kernel_size = kernel_size
        ctx.scale = scale
        if ctx.needs_input_grad[0]:
            ctx.save_for_backward(x)
        return _fused_ops.pool_sum_scale_cuda(x, kernel_size, scale)

    @staticmethod
    def backward(ctx, grad_output):
        (x,) = ctx.saved_tensors
        _, indices = torch.nn.functional.max_pool1d(
            x.unsqueeze(1),
            kernel_size=ctx.kernel_size,
            stride=ctx.kernel_size,
            return_indices=True,
        )
        indices = indices.squeeze(1)
        values = (grad_output * ctx.scale).unsqueeze(1).expand_as(indices)
        grad_input = torch.zeros_like(x)
        grad_input.scatter_(1, indices, values)
        return grad_input, None, None


class ModelNew(nn.Module):
    def __init__(self, in_features, out_features, kernel_size, scale_factor):
        super().__init__()
        self.matmul = nn.Linear(in_features, out_features)
        self.max_pool = nn.MaxPool1d(kernel_size)
        self.kernel_size = (
            kernel_size[0]
            if isinstance(kernel_size, (tuple, list))
            else kernel_size
        )
        self.scale_factor = scale_factor

    def forward(self, x):
        x = self.matmul(x)
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 2
            and x.is_contiguous()
        ):
            if torch.is_grad_enabled() and x.requires_grad:
                return _PoolSumScale.apply(
                    x, self.kernel_size, self.scale_factor
                )
            return _fused_ops.pool_sum_scale_cuda(
                x, self.kernel_size, self.scale_factor
            )

        x = self.max_pool(x.unsqueeze(1)).squeeze(1)
        return x.sum(dim=1) * self.scale_factor
```