```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 kernel: max-pool (kernel_size=2, non-overlapping) + sum over features + scale.
# This turns three separate ops (each materializing intermediate tensors) into a
# single memory-bound pass over the matmul output that writes only `batch_size` floats.
# The GEMM itself stays on cuBLAS (F.linear), which is already optimal for this shape.
_fused_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>

#define THREADS 512

__device__ __forceinline__ double blockReduceSum(double val) {
    __shared__ double shared[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) shared[wid] = val;
    __syncthreads();

    const int nwarp = (blockDim.x + 31) >> 5;
    if (wid == 0) {
        val = (lane < nwarp) ? shared[lane] : 0.0;
        #pragma unroll
        for (int offset = 16; offset > 0; offset >>= 1)
            val += __shfl_down_sync(0xffffffffu, val, offset);
    }
    return val;  // only thread 0 holds the block total
}

// Fast path: feature count divisible by 4 -> aligned float4 loads.
__global__ void maxpool2_sum_scale_vec4_kernel(
    const float* __restrict__ x,
    float* __restrict__ out,
    const int row_stride,
    const int num_pairs,
    const float scale)
{
    const int row = blockIdx.x;
    const float4* __restrict__ p =
        reinterpret_cast<const float4*>(x + (long long)row * row_stride);
    const int n4 = num_pairs >> 1;

    double acc = 0.0;
    for (int i = threadIdx.x; i < n4; i += blockDim.x) {
        float4 v = p[i];
        acc += (double)(fmaxf(v.x, v.y) + fmaxf(v.z, v.w));
    }
    // Tail pair when num_pairs is odd.
    if ((num_pairs & 1) && threadIdx.x == 0) {
        const float* sp = x + (long long)row * row_stride
                          + (long long)(num_pairs - 1) * 2;
        acc += (double)fmaxf(sp[0], sp[1]);
    }

    acc = blockReduceSum(acc);
    if (threadIdx.x == 0) out[row] = (float)(acc * (double)scale);
}

// Generic fallback for any feature count.
__global__ void maxpool2_sum_scale_scalar_kernel(
    const float* __restrict__ x,
    float* __restrict__ out,
    const int row_stride,
    const int num_pairs,
    const float scale)
{
    const int row = blockIdx.x;
    const float* __restrict__ p = x + (long long)row * row_stride;

    double acc = 0.0;
    for (int i = threadIdx.x; i < num_pairs; i += blockDim.x) {
        acc += (double)fmaxf(p[2 * i], p[2 * i + 1]);
    }

    acc = blockReduceSum(acc);
    if (threadIdx.x == 0) out[row] = (float)(acc * (double)scale);
}

torch::Tensor maxpool_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() == torch::kFloat32, "x must be float32");
    TORCH_CHECK(x.dim() == 2, "x must be 2D (batch, features)");
    TORCH_CHECK(kernel_size == 2, "only kernel_size == 2 is supported");

    auto xc = x.contiguous();
    const int rows = xc.size(0);
    const int cols = xc.size(1);
    const int num_pairs = cols / 2;

    auto out = torch::empty({rows}, xc.options());
    if (rows == 0) return out;
    if (num_pairs == 0) { out.zero_(); return out; }

    const dim3 grid(rows);
    const dim3 block(THREADS);
    const float s = (float)scale;
    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if ((cols & 3) == 0) {
        maxpool2_sum_scale_vec4_kernel<<<grid, block, 0, stream>>>(
            xc.data_ptr<float>(), out.data_ptr<float>(), cols, num_pairs, s);
    } else {
        maxpool2_sum_scale_scalar_kernel<<<grid, block, 0, stream>>>(
            xc.data_ptr<float>(), out.data_ptr<float>(), cols, num_pairs, s);
    }
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

_fused_cpp_source = (
    "torch::Tensor maxpool_sum_scale_cuda(torch::Tensor x, int64_t kernel_size, double scale);"
)

_fused_op = load_inline(
    name="maxpool_sum_scale_fused",
    cpp_sources=_fused_cpp_source,
    cuda_sources=_fused_cuda_source,
    functions=["maxpool_sum_scale_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model:
      - GEMM + bias stays on cuBLAS via F.linear (optimal for this shape).
      - max-pool(k=2) + sum + scale are fused into a single custom CUDA kernel
        that reads the GEMM output once and writes only `batch_size` floats,
        eliminating two intermediate tensor materializations.
    """
    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 x.is_cuda and x.dtype == torch.float32 and self.kernel_size == 2:
            return _fused_op.maxpool_sum_scale_cuda(
                x, self.kernel_size, float(self.scale_factor)
            )
        # Fallback (CPU / other dtypes / other kernel sizes): original ops.
        x = F.max_pool1d(x.unsqueeze(1), self.kernel_size).squeeze(1)
        x = torch.sum(x, dim=1)
        return x * self.scale_factor
```