```python
# Copyright (c) Meta Platforms, Inc. and affiliates.
# All rights reserved.
#
# This source code is licensed under the license found in the
# LICENSE file in the root directory of this source tree.

import torch
import torch.nn as nn
import torch.nn.functional as F
from torch.utils.cpp_extension import load_inline

# Fused CUDA kernel: MaxPool1d(kernel_size=2) + sum over features + scale.
# The matmul (which dominates runtime) is left to cuBLAS via F.linear;
# everything after it is fused into a single kernel that reads the
# (batch, out_features) matrix once and produces a (batch,) result.
fused_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>

__inline__ __device__ float warpReduceSum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_down_sync(0xffffffff, val, offset);
    return val;
}

__inline__ __device__ float blockReduceSum(float val) {
    __shared__ float shared[32];
    int lane = threadIdx.x & 31;
    int wid  = threadIdx.x >> 5;
    val = warpReduceSum(val);
    if (lane == 0) shared[wid] = val;
    __syncthreads();
    int nwarp = (blockDim.x + 31) >> 5;
    if (wid == 0) {
        val = (lane < nwarp) ? shared[lane] : 0.0f;
        val = warpReduceSum(val);
    }
    return val;
}

// Vectorized path: one float4 load covers two pooling pairs.
__global__ void maxpool2_sum_scale_vec4_kernel(const float4* __restrict__ x,
                                               float* __restrict__ out,
                                               int n_vec,
                                               float scale) {
    const float4* row = x + (size_t)blockIdx.x * (size_t)n_vec;
    float acc = 0.0f;
    for (int i = threadIdx.x; i < n_vec; i += blockDim.x) {
        float4 v = row[i];
        acc += fmaxf(v.x, v.y) + fmaxf(v.z, v.w);
    }
    acc = blockReduceSum(acc);
    if (threadIdx.x == 0) out[blockIdx.x] = acc * scale;
}

// Generic path for feature counts not divisible by 4.
__global__ void maxpool2_sum_scale_scalar_kernel(const float* __restrict__ x,
                                                 float* __restrict__ out,
                                                 int n_pairs,
                                                 float scale) {
    const float* row = x + (size_t)blockIdx.x * ((size_t)n_pairs * 2);
    float acc = 0.0f;
    for (int i = threadIdx.x; i < n_pairs; i += blockDim.x) {
        float a = row[2 * i];
        float b = row[2 * i + 1];
        acc += fmaxf(a, b);
    }
    acc = blockReduceSum(acc);
    if (threadIdx.x == 0) out[blockIdx.x] = acc * scale;
}

torch::Tensor maxpool_sum_scale(torch::Tensor x, double scale) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "x must be float32");
    TORCH_CHECK(x.dim() == 2, "x must be 2D (batch, features)");

    auto xc = x.contiguous();
    const int batch = (int)xc.size(0);
    const int n = (int)xc.size(1);
    auto out = torch::empty({batch}, xc.options());
    if (batch == 0 || n == 0) return out;

    const int threads = 256;
    const float s = (float)scale;
    if (n % 4 == 0) {
        maxpool2_sum_scale_vec4_kernel<<<batch, threads>>>(
            reinterpret_cast<const float4*>(xc.data_ptr<float>()),
            out.data_ptr<float>(), n / 4, s);
    } else {
        maxpool2_sum_scale_scalar_kernel<<<batch, threads>>>(
            xc.data_ptr<float>(), out.data_ptr<float>(), n / 2, s);
    }
    return out;
}
"""

fused_cpp_source = "torch::Tensor maxpool_sum_scale(torch::Tensor x, double scale);"

fused_ops = load_inline(
    name="fused_maxpool_sum_scale",
    cpp_sources=fused_cpp_source,
    cuda_sources=fused_cuda_source,
    functions=["maxpool_sum_scale"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model: keeps the cuBLAS-backed linear layer and fuses
    max pooling (kernel_size=2), feature sum, and scaling into a single
    custom CUDA kernel.
    """
    def __init__(self, in_features, out_features, kernel_size, scale_factor):
        super(ModelNew, self).__init__()
        self.matmul = nn.Linear(in_features, out_features)
        self.kernel_size = kernel_size
        self.scale_factor = scale_factor

    def forward(self, x):
        x = F.linear(x, self.matmul.weight, self.matmul.bias)
        if (self.kernel_size == 2 and x.is_cuda
                and x.dtype == torch.float32 and x.dim() == 2):
            return fused_ops.maxpool_sum_scale(x, float(self.scale_factor))
        # Fallback to the original sequence for any other configuration.
        x = F.max_pool1d(x.unsqueeze(1), self.kernel_size).squeeze(1)
        x = torch.sum(x, dim=1)
        return x * self.scale_factor
```