```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 that replaces:
#   MaxPool1d(kernel_size=2) -> sum(dim=1) -> * scale_factor
# For each row of the (batch, features) matmul output it computes:
#   out[b] = scale * sum_i max(x[b, 2*i], x[b, 2*i + 1])
# This avoids materializing the pooled tensor and collapses 3 kernel
# launches + extra global memory traffic into a single lightweight pass.
fused_pool_sum_scale_cuda_source = r"""
#include <torch/extension.h>
#include <c10/cuda/CUDAStream.h>
#include <cuda_runtime.h>

__global__ void fused_pool_sum_scale_kernel(const float* __restrict__ x,
                                            float* __restrict__ out,
                                            const int N,
                                            const float scale) {
    const int row = blockIdx.x;
    const float* __restrict__ row_ptr = x + (long)row * N;
    const int stride = blockDim.x;

    float acc0 = 0.0f, acc1 = 0.0f, acc2 = 0.0f, acc3 = 0.0f;

    if ((N & 3) == 0) {
        // Vectorized path: each float4 covers two pool pairs.
        const float4* __restrict__ p4 = reinterpret_cast<const float4*>(row_ptr);
        const int nv4 = N >> 2;
        int i = threadIdx.x;
        for (; i + 3 * stride < nv4; i += 4 * stride) {
            float4 a = p4[i];
            float4 b = p4[i + stride];
            float4 c = p4[i + 2 * stride];
            float4 d = p4[i + 3 * stride];
            acc0 += fmaxf(a.x, a.y) + fmaxf(a.z, a.w);
            acc1 += fmaxf(b.x, b.y) + fmaxf(b.z, b.w);
            acc2 += fmaxf(c.x, c.y) + fmaxf(c.z, c.w);
            acc3 += fmaxf(d.x, d.y) + fmaxf(d.z, d.w);
        }
        for (; i < nv4; i += stride) {
            float4 a = p4[i];
            acc0 += fmaxf(a.x, a.y) + fmaxf(a.z, a.w);
        }
    } else {
        // Scalar fallback for feature counts not divisible by 4.
        const int npairs = N >> 1;
        for (int i = threadIdx.x; i < npairs; i += stride) {
            acc0 += fmaxf(row_ptr[2 * i], row_ptr[2 * i + 1]);
        }
    }

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

    // Warp-level reduction.
    for (int offset = 16; offset > 0; offset >>= 1) {
        acc += __shfl_down_sync(0xffffffffu, acc, offset);
    }
    __shared__ float warp_sums[32];
    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;
    if (lane == 0) {
        warp_sums[warp] = acc;
    }
    __syncthreads();
    if (warp == 0) {
        const int nwarps = (blockDim.x + 31) >> 5;
        float v = (lane < nwarps) ? warp_sums[lane] : 0.0f;
        for (int offset = 16; offset > 0; offset >>= 1) {
            v += __shfl_down_sync(0xffffffffu, v, offset);
        }
        if (lane == 0) {
            out[row] = v * scale;
        }
    }
}

torch::Tensor fused_pool_sum_scale_cuda(torch::Tensor x, double scale) {
    TORCH_CHECK(x.is_cuda(), "x must be a CUDA tensor");
    TORCH_CHECK(x.dim() == 2, "x must be 2D (batch, features)");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    auto xc = x.contiguous();
    const int64_t B = xc.size(0);
    const int64_t N = xc.size(1);
    TORCH_CHECK((N % 2) == 0, "feature count must be even");

    auto out = torch::empty({B}, xc.options());
    if (B == 0 || N == 0) {
        out.zero_();
        return out;
    }

    const int threads = 256;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    fused_pool_sum_scale_kernel<<<(int)B, threads, 0, stream>>>(
        xc.data_ptr<float>(), out.data_ptr<float>(), (int)N, (float)scale);
    return out;
}
"""

fused_pool_sum_scale_cpp_source = (
    "torch::Tensor fused_pool_sum_scale_cuda(torch::Tensor x, double scale);"
)

fused_pool_sum_scale = load_inline(
    name="fused_pool_sum_scale_ext",
    cpp_sources=fused_pool_sum_scale_cpp_source,
    cuda_sources=fused_pool_sum_scale_cuda_source,
    functions=["fused_pool_sum_scale_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model:
      - The (batch, in) x (in, out) linear layer stays on the highly-tuned
        cuBLAS path via F.linear.
      - MaxPool1d(kernel_size=2) + sum + scaling are fused into a single
        custom CUDA kernel (one pass over the GEMM output, no intermediate
        pooled tensor, one kernel launch instead of three).
    """
    def __init__(self, in_features, out_features, kernel_size, scale_factor):
        super(ModelNew, self).__init__()
        self.matmul = nn.Linear(in_features, out_features)
        self.max_pool = nn.MaxPool1d(kernel_size)
        self.kernel_size = kernel_size
        self.scale_factor = scale_factor
        self.fused_op = fused_pool_sum_scale

    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):
            return self.fused_op.fused_pool_sum_scale_cuda(
                x, float(self.scale_factor))
        # Generic fallback for other configurations.
        x = self.max_pool(x.unsqueeze(1)).squeeze(1)
        x = torch.sum(x, dim=1)
        x = x * self.scale_factor
        return x
```