```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: max-pool (kernel_size=2, stride=2) + full reduction over
# the feature dimension + scaling, all in a single kernel launch. This removes
# the intermediate tensors and extra kernel launches produced by the separate
# MaxPool1d / sum / mul operators. The large GEMM is left to cuBLAS (F.linear),
# which is already near-optimal for this shape.
fused_cuda_source = """
#include <torch/extension.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <cstdint>

__global__ void maxpool_sum_scale_kernel(
    const float* __restrict__ y,
    float* __restrict__ out,
    const long out_features,
    const float scale)
{
    const long row = blockIdx.x;
    const float* __restrict__ y_row = y + row * out_features;
    // MaxPool1d(kernel=2, stride=2, padding=0) keeps floor(L/2) pairs and
    // drops a trailing odd element.
    const long num_pairs = out_features >> 1;

    float local = 0.0f;

    if ((reinterpret_cast<std::uintptr_t>(y_row) & 15) == 0) {
        // 16-byte aligned: vectorized path, one float4 covers two pairs.
        const float4* __restrict__ y4 = reinterpret_cast<const float4*>(y_row);
        const long num_quads = num_pairs >> 1;
        for (long i = threadIdx.x; i < num_quads; i += blockDim.x) {
            const float4 v = y4[i];
            local += fmaxf(v.x, v.y);
            local += fmaxf(v.z, v.w);
        }
        if ((num_pairs & 1) && threadIdx.x == 0) {
            const long p = num_pairs - 1;
            local += fmaxf(y_row[2 * p], y_row[2 * p + 1]);
        }
    } else {
        for (long i = threadIdx.x; i < num_pairs; i += blockDim.x) {
            local += fmaxf(y_row[2 * i], y_row[2 * i + 1]);
        }
    }

    // Warp-level reduction.
    for (int offset = 16; offset > 0; offset >>= 1) {
        local += __shfl_down_sync(0xffffffffu, local, offset);
    }

    __shared__ float smem[32];
    const int lane = threadIdx.x & 31;
    const int wid  = threadIdx.x >> 5;
    if (lane == 0) {
        smem[wid] = local;
    }
    __syncthreads();

    if (wid == 0) {
        const int num_warps = (blockDim.x + 31) >> 5;
        float v = (lane < num_warps) ? smem[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_maxpool_sum_scale(torch::Tensor y, double scale) {
    TORCH_CHECK(y.is_cuda(), "input must be a CUDA tensor");
    TORCH_CHECK(y.scalar_type() == at::kFloat, "input must be float32");
    TORCH_CHECK(y.is_contiguous(), "input must be contiguous");
    TORCH_CHECK(y.dim() == 2, "input must be 2D (batch, features)");

    const long batch = y.size(0);
    const long out_features = y.size(1);

    auto out = torch::empty({batch}, y.options());
    if (batch == 0 || out_features == 0) {
        out.zero_();
        return out;
    }

    const int threads = 256;
    maxpool_sum_scale_kernel<<<(int)batch, threads>>>(
        y.data_ptr<float>(),
        out.data_ptr<float>(),
        out_features,
        static_cast<float>(scale));
    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return out;
}
"""

fused_cpp_source = (
    "torch::Tensor fused_maxpool_sum_scale(torch::Tensor y, double scale);"
)

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


class ModelNew(nn.Module):
    """
    Optimized model: GEMM via cuBLAS, with max-pool + sum + scaling fused 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.max_pool = nn.MaxPool1d(kernel_size)
        self.kernel_size = kernel_size
        self.scale_factor = scale_factor

    def forward(self, x):
        if x.is_cuda and x.dtype == torch.float32 and self.kernel_size == 2:
            y = F.linear(x, self.matmul.weight, self.matmul.bias)
            return fused_ops.fused_maxpool_sum_scale(y, self.scale_factor)
        # Generic fallback path (non-CUDA / other kernel sizes).
        x = self.matmul(x)
        x = self.max_pool(x.unsqueeze(1)).squeeze(1)
        x = torch.sum(x, dim=1)
        x = x * self.scale_factor
        return x
```