```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 implementing:
#     MaxPool1d(kernel_size=2) -> sum(dim=1) -> multiply by scale
# i.e. out[b] = scale * sum_j max(h[b, 2*j], h[b, 2*j + 1])
# This collapses three PyTorch operators (and their intermediate tensors
# and kernel launches) into a single lightweight kernel.
fused_maxpool_sum_scale_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

__device__ __forceinline__ double blockReduceSum(double val) {
    __shared__ double warp_sums[32];
    const int lane = threadIdx.x & 31;
    const int wid  = threadIdx.x >> 5;

    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffffu, val, offset);
    }
    if (lane == 0) warp_sums[wid] = val;
    __syncthreads();

    const int num_warps = (blockDim.x + 31) >> 5;
    val = (threadIdx.x < num_warps) ? warp_sums[threadIdx.x] : 0.0;
    if (wid == 0) {
        #pragma unroll
        for (int offset = 16; offset > 0; offset >>= 1) {
            val += __shfl_down_sync(0xffffffffu, val, offset);
        }
    }
    return val;  // result is valid in thread 0
}

__global__ void maxpool2_sum_scale_kernel(const float* __restrict__ x,
                                          float* __restrict__ out,
                                          const long n,
                                          const float scale) {
    const float* row = x + (long)blockIdx.x * n;
    double acc = 0.0;

    if ((n & 3L) == 0L) {
        // Vectorized path: every float4 contains exactly two pool windows.
        // (Rows stay 16-byte aligned because n is a multiple of 4.)
        const float4* row4 = reinterpret_cast<const float4*>(row);
        const long nvec = n >> 2;
        for (long i = threadIdx.x; i < nvec; i += blockDim.x) {
            const float4 v = row4[i];
            acc += (double)fmaxf(v.x, v.y);
            acc += (double)fmaxf(v.z, v.w);
        }
    } else {
        // Generic path for any feature count (odd tail is dropped, matching
        // MaxPool1d with kernel_size=stride=2, padding=0).
        const long pairs = n >> 1;
        for (long i = threadIdx.x; i < pairs; i += blockDim.x) {
            acc += (double)fmaxf(row[2 * i], row[2 * i + 1]);
        }
    }

    const double total = blockReduceSum(acc);
    if (threadIdx.x == 0) {
        out[blockIdx.x] = (float)(total * (double)scale);
    }
}

torch::Tensor maxpool2_sum_scale_cuda(torch::Tensor x, 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 must be 2D (batch, features)");

    auto xc = x.contiguous();
    const long batch = xc.size(0);
    const long n = xc.size(1);

    auto out = torch::empty({batch}, xc.options());
    if (batch == 0) return out;
    if (n < 2) {
        out.zero_();
        return out;
    }

    const int threads = 256;
    const dim3 blocks((unsigned int)batch);
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();
    maxpool2_sum_scale_kernel<<<blocks, threads, 0, stream>>>(
        xc.data_ptr<float>(), out.data_ptr<float>(), n, (float)scale);
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

fused_maxpool_sum_scale_cpp_src = (
    "torch::Tensor maxpool2_sum_scale_cuda(torch::Tensor x, double scale);"
)

fused_maxpool_sum_scale = load_inline(
    name="fused_maxpool2_sum_scale",
    cpp_sources=fused_maxpool_sum_scale_cpp_src,
    cuda_sources=fused_maxpool_sum_scale_cuda_src,
    functions=["maxpool2_sum_scale_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized version of Model.

    The nn.Linear GEMM dominates the runtime and is left to the highly
    tuned cuBLAS implementation. Everything after it
    (MaxPool1d(kernel_size=2) -> sum over features -> scaling) is replaced
    by a single fused custom CUDA kernel that computes
        out[b] = scale_factor * sum_j max(h[b, 2j], h[b, 2j + 1])
    in one pass, eliminating the intermediate pooled tensor and the extra
    kernel launches while accumulating in double precision for accuracy.
    """
    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

    def forward(self, x):
        x = self.matmul(x)
        if self.kernel_size == 2 and x.is_cuda and x.dtype == torch.float32:
            return fused_maxpool_sum_scale.maxpool2_sum_scale_cuda(
                x, float(self.scale_factor))
        # Generic fallback preserving the original semantics.
        x = self.max_pool(x.unsqueeze(1)).squeeze(1)
        x = torch.sum(x, dim=1)
        x = x * self.scale_factor
        return x
```