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

_cpp_source = r"""
#include <torch/extension.h>
#include <vector>

std::vector<torch::Tensor> linear_pool_sum_cuda(
    torch::Tensor x, torch::Tensor weight, torch::Tensor bias,
    int64_t kernel_size, double scale);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cublasLt.h>
#include <cuda_runtime.h>
#include <math_constants.h>
#include <algorithm>
#include <climits>
#include <cstdint>
#include <limits>
#include <map>
#include <mutex>
#include <tuple>
#include <vector>

static void check_blas(cublasStatus_t status) {
    TORCH_CHECK(status == CUBLAS_STATUS_SUCCESS,
                "cuBLASLt error: ", static_cast<int>(status));
}

struct LtHandle {
    cublasLtHandle_t handle;
    LtHandle() { check_blas(cublasLtCreate(&handle)); }
    ~LtHandle() { cublasLtDestroy(handle); }
    LtHandle(const LtHandle&) = delete;
    LtHandle& operator=(const LtHandle&) = delete;
};

struct Descriptors {
    cublasLtMatmulDesc_t operation;
    cublasLtMatrixLayout_t a, b, c;

    Descriptors(int batch, int inputs, int outputs) {
        check_blas(cublasLtMatmulDescCreate(
            &operation, CUBLAS_COMPUTE_32F_FAST_TF32, CUDA_R_32F));

        cublasOperation_t trans_a = CUBLAS_OP_T;
        cublasOperation_t trans_b = CUBLAS_OP_N;
        check_blas(cublasLtMatmulDescSetAttribute(
            operation, CUBLASLT_MATMUL_DESC_TRANSA,
            &trans_a, sizeof(trans_a)));
        check_blas(cublasLtMatmulDescSetAttribute(
            operation, CUBLASLT_MATMUL_DESC_TRANSB,
            &trans_b, sizeof(trans_b)));

        check_blas(cublasLtMatrixLayoutCreate(
            &a, CUDA_R_32F, inputs, outputs, inputs));
        check_blas(cublasLtMatrixLayoutCreate(
            &b, CUDA_R_32F, inputs, batch, inputs));
        check_blas(cublasLtMatrixLayoutCreate(
            &c, CUDA_R_32F, outputs, batch, outputs));
    }

    ~Descriptors() {
        cublasLtMatrixLayoutDestroy(a);
        cublasLtMatrixLayoutDestroy(b);
        cublasLtMatrixLayoutDestroy(c);
        cublasLtMatmulDescDestroy(operation);
    }
};

struct Plan {
    cublasLtMatmulAlgo_t algorithm;
    size_t workspace;
};

static std::mutex plan_mutex;
static std::map<int, LtHandle> handles;
static std::map<std::tuple<int, int, int, int, int, int>, Plan> plans;

static cublasStatus_t launch_gemm(
    cublasLtHandle_t handle,
    const Descriptors& desc,
    const Plan& plan,
    const float* weight,
    const float* x,
    float* output,
    void* workspace,
    cudaStream_t stream
) {
    const float alpha = 1.0f;
    const float beta = 0.0f;
    return cublasLtMatmul(
        handle, desc.operation,
        &alpha, weight, desc.a, x, desc.b,
        &beta, output, desc.c, output, desc.c,
        &plan.algorithm, workspace, plan.workspace, stream);
}

static int pointer_alignment(const void* pointer) {
    const uintptr_t address = reinterpret_cast<uintptr_t>(pointer);
    int alignment = 1;
    while (alignment < 256 && address % (alignment * 2) == 0)
        alignment *= 2;
    return alignment;
}

static Plan select_plan(
    cublasLtHandle_t handle,
    const Descriptors& desc,
    const torch::Tensor& x,
    const torch::Tensor& weight,
    torch::Tensor& output,
    cudaStream_t stream,
    bool tune
) {
    constexpr size_t workspace_limit = 128ULL * 1024 * 1024;

    cublasLtMatmulPreference_t preference;
    check_blas(cublasLtMatmulPreferenceCreate(&preference));
    check_blas(cublasLtMatmulPreferenceSetAttribute(
        preference, CUBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES,
        &workspace_limit, sizeof(workspace_limit)));

    const uint32_t alignment_a = pointer_alignment(weight.data_ptr());
    const uint32_t alignment_b = pointer_alignment(x.data_ptr());
    check_blas(cublasLtMatmulPreferenceSetAttribute(
        preference, CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_A_BYTES,
        &alignment_a, sizeof(alignment_a)));
    check_blas(cublasLtMatmulPreferenceSetAttribute(
        preference, CUBLASLT_MATMUL_PREF_MIN_ALIGNMENT_B_BYTES,
        &alignment_b, sizeof(alignment_b)));

    cublasLtMatmulHeuristicResult_t heuristics[16];
    int count = 0;
    const auto status = cublasLtMatmulAlgoGetHeuristic(
        handle, desc.operation, desc.a, desc.b, desc.c, desc.c,
        preference, 16, heuristics, &count);
    cublasLtMatmulPreferenceDestroy(preference);
    check_blas(status);
    TORCH_CHECK(count > 0, "No supported cuBLASLt algorithm");

    std::vector<Plan> candidates;
    auto add_candidate = [&](const cublasLtMatmulAlgo_t& algorithm) {
        cublasLtMatmulHeuristicResult_t checked;
        const auto s = cublasLtMatmulAlgoCheck(
            handle, desc.operation, desc.a, desc.b, desc.c, desc.c,
            &algorithm, &checked);
        if (s == CUBLAS_STATUS_SUCCESS &&
            checked.state == CUBLAS_STATUS_SUCCESS &&
            checked.workspaceSize <= workspace_limit) {
            candidates.push_back({algorithm, checked.workspaceSize});
        }
    };

    for (int i = 0; i < count; ++i) {
        if (heuristics[i].state == CUBLAS_STATUS_SUCCESS)
            add_candidate(heuristics[i].algo);
    }

    TORCH_CHECK(!candidates.empty(), "No valid cuBLASLt algorithm");
    if (!tune)
        return candidates.front();

    for (int h = 0; h < std::min(count, 4); ++h) {
        if (heuristics[h].state != CUBLAS_STATUS_SUCCESS)
            continue;
        for (uint32_t split : {2u, 4u, 8u}) {
            auto algorithm = heuristics[h].algo;
            const uint32_t reduction = CUBLASLT_REDUCTION_SCHEME_INPLACE;
            if (cublasLtMatmulAlgoConfigSetAttribute(
                    &algorithm, CUBLASLT_ALGO_CONFIG_SPLITK_NUM,
                    &split, sizeof(split)) != CUBLAS_STATUS_SUCCESS)
                continue;
            if (cublasLtMatmulAlgoConfigSetAttribute(
                    &algorithm, CUBLASLT_ALGO_CONFIG_REDUCTION_SCHEME,
                    &reduction, sizeof(reduction)) != CUBLAS_STATUS_SUCCESS)
                continue;
            add_candidate(algorithm);
        }
    }

    const uint32_t tiles[] = {
        CUBLASLT_MATMUL_TILE_64x64,
        CUBLASLT_MATMUL_TILE_64x128,
        CUBLASLT_MATMUL_TILE_128x64,
        CUBLASLT_MATMUL_TILE_128x128,
        CUBLASLT_MATMUL_TILE_256x64,
        CUBLASLT_MATMUL_TILE_256x128
    };

    for (int h = 0; h < std::min(count, 2); ++h) {
        if (heuristics[h].state != CUBLAS_STATUS_SUCCESS)
            continue;
        for (uint32_t tile : tiles) {
            for (uint32_t split : {1u, 4u}) {
                auto algorithm = heuristics[h].algo;
                const uint32_t reduction =
                    split == 1 ? CUBLASLT_REDUCTION_SCHEME_NONE
                               : CUBLASLT_REDUCTION_SCHEME_INPLACE;
                if (cublasLtMatmulAlgoConfigSetAttribute(
                        &algorithm, CUBLASLT_ALGO_CONFIG_TILE_ID,
                        &tile, sizeof(tile)) != CUBLAS_STATUS_SUCCESS)
                    continue;
                if (cublasLtMatmulAlgoConfigSetAttribute(
                        &algorithm, CUBLASLT_ALGO_CONFIG_SPLITK_NUM,
                        &split, sizeof(split)) != CUBLAS_STATUS_SUCCESS)
                    continue;
                if (cublasLtMatmulAlgoConfigSetAttribute(
                        &algorithm, CUBLASLT_ALGO_CONFIG_REDUCTION_SCHEME,
                        &reduction, sizeof(reduction)) != CUBLAS_STATUS_SUCCESS)
                    continue;
                add_candidate(algorithm);
            }
        }
    }

    size_t scratch_bytes = 0;
    for (const auto& candidate : candidates)
        scratch_bytes = std::max(scratch_bytes, candidate.workspace);

    auto scratch = torch::empty(
        {static_cast<int64_t>(scratch_bytes)},
        x.options().dtype(at::kByte));
    void* scratch_pointer = scratch_bytes ? scratch.data_ptr() : nullptr;

    cudaEvent_t start, stop;
    C10_CUDA_CHECK(cudaEventCreate(&start));
    C10_CUDA_CHECK(cudaEventCreate(&stop));

    Plan best = candidates.front();
    float best_time = std::numeric_limits<float>::infinity();

    for (const auto& candidate : candidates) {
        if (launch_gemm(
                handle, desc, candidate,
                weight.data_ptr<float>(), x.data_ptr<float>(),
                output.data_ptr<float>(), scratch_pointer, stream)
            != CUBLAS_STATUS_SUCCESS)
            continue;

        C10_CUDA_CHECK(cudaEventRecord(start, stream));
        const auto s = launch_gemm(
            handle, desc, candidate,
            weight.data_ptr<float>(), x.data_ptr<float>(),
            output.data_ptr<float>(), scratch_pointer, stream);
        C10_CUDA_CHECK(cudaEventRecord(stop, stream));
        C10_CUDA_CHECK(cudaEventSynchronize(stop));

        if (s != CUBLAS_STATUS_SUCCESS)
            continue;

        float elapsed = 0.0f;
        C10_CUDA_CHECK(cudaEventElapsedTime(&elapsed, start, stop));
        if (elapsed < best_time) {
            best_time = elapsed;
            best = candidate;
        }
    }

    C10_CUDA_CHECK(cudaEventDestroy(start));
    C10_CUDA_CHECK(cudaEventDestroy(stop));
    TORCH_CHECK(best_time < std::numeric_limits<float>::infinity(),
                "All cuBLASLt configurations failed");
    return best;
}

__device__ __forceinline__ float maximum_nan(float a, float b) {
    return isnan(b) || a < b ? b : a;
}

__device__ __forceinline__ float warp_sum(float value) {
    #pragma unroll
    for (int offset = 16; offset > 0; offset >>= 1)
        value += __shfl_down_sync(0xffffffff, value, offset);
    return value;
}

template<bool PAIRS>
__global__ void bias_pool_sum_kernel(
    const float* __restrict__ input,
    const float* __restrict__ bias,
    float* __restrict__ output,
    int columns,
    int kernel_size,
    float scale
) {
    const int tid = threadIdx.x;
    const float* row = input + static_cast<int64_t>(blockIdx.x) * columns;
    float sum0 = 0.0f;
    float sum1 = 0.0f;

    if constexpr (PAIRS) {
        const float4* values = reinterpret_cast<const float4*>(row);
        const float4* biases = reinterpret_cast<const float4*>(bias);
        for (int i = tid; i < columns / 4; i += blockDim.x) {
            const float4 v = values[i];
            const float4 b = biases[i];
            sum0 += maximum_nan(v.x + b.x, v.y + b.y);
            sum1 += maximum_nan(v.z + b.z, v.w + b.w);
        }
    } else {
        for (int i = tid; i < columns / kernel_size; i += blockDim.x) {
            const int start = i * kernel_size;
            float maximum = -CUDART_INF_F;
            for (int j = 0; j < kernel_size; ++j)
                maximum = maximum_nan(
                    maximum, row[start + j] + bias[start + j]);
            sum0 += maximum;
        }
    }

    float sum = warp_sum(sum0 + sum1);
    __shared__ float partial[8];
    const int lane = tid & 31;
    const int warp = tid >> 5;

    if (lane == 0)
        partial[warp] = sum;
    __syncthreads();

    if (warp == 0) {
        sum = warp_sum(lane < 8 ? partial[lane] : 0.0f);
        if (lane == 0)
            output[blockIdx.x] = sum * scale;
    }
}

std::vector<torch::Tensor> linear_pool_sum_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias,
    int64_t kernel_size,
    double scale
) {
    TORCH_CHECK(x.is_cuda() && weight.is_cuda() && bias.is_cuda(),
                "Expected CUDA tensors");
    TORCH_CHECK(x.device() == weight.device() &&
                x.device() == bias.device(), "Device mismatch");
    TORCH_CHECK(x.scalar_type() == at::kFloat &&
                weight.scalar_type() == at::kFloat &&
                bias.scalar_type() == at::kFloat, "Expected float32");
    TORCH_CHECK(x.dim() == 2 && weight.dim() == 2 && bias.dim() == 1,
                "Invalid dimensions");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous() &&
                bias.is_contiguous(), "Expected contiguous tensors");
    TORCH_CHECK(x.size(1) == weight.size(1) &&
                weight.size(0) == bias.size(0), "Shape mismatch");
    TORCH_CHECK(kernel_size > 0 && kernel_size <= weight.size(0),
                "Invalid pooling kernel");
    TORCH_CHECK(x.size(0) <= INT_MAX && x.size(1) <= INT_MAX &&
                weight.size(0) <= INT_MAX, "Dimensions too large");

    const c10::cuda::CUDAGuard guard(x.device());
    const int batch = static_cast<int>(x.size(0));
    const int inputs = static_cast<int>(x.size(1));
    const int outputs = static_cast<int>(weight.size(0));

    auto linear = torch::empty({batch, outputs}, x.options());
    auto result = torch::empty({batch}, x.options());
    if (batch == 0)
        return {result, linear};

    cudaStream_t stream = at::cuda::getCurrentCUDAStream();

    if (inputs == 0) {
        C10_CUDA_CHECK(cudaMemsetAsync(
            linear.data_ptr<float>(), 0,
            linear.numel() * sizeof(float), stream));
    } else {
        std::lock_guard<std::mutex> lock(plan_mutex);
        const int device = x.get_device();
        const auto handle = handles.try_emplace(device).first->second.handle;
        Descriptors desc(batch, inputs, outputs);

        const auto key = std::make_tuple(
            device, batch, inputs, outputs,
            pointer_alignment(weight.data_ptr()),
            pointer_alignment(x.data_ptr()));
        auto it = plans.find(key);
        Plan plan;

        if (it != plans.end()) {
            plan = it->second;
        } else {
            cudaStreamCaptureStatus capture;
            C10_CUDA_CHECK(cudaStreamIsCapturing(stream, &capture));
            const bool tune = capture == cudaStreamCaptureStatusNone;
            plan = select_plan(handle, desc, x, weight, linear, stream, tune);
            if (tune)
                plans.emplace(key, plan);
        }

        auto workspace = torch::empty(
            {static_cast<int64_t>(plan.workspace)},
            x.options().dtype(at::kByte));
        check_blas(launch_gemm(
            handle, desc, plan,
            weight.data_ptr<float>(), x.data_ptr<float>(),
            linear.data_ptr<float>(),
            plan.workspace ? workspace.data_ptr() : nullptr, stream));
    }

    if (kernel_size == 2 && outputs % 4 == 0 &&
        pointer_alignment(bias.data_ptr()) >= 16) {
        bias_pool_sum_kernel<true><<<batch, 256, 0, stream>>>(
            linear.data_ptr<float>(), bias.data_ptr<float>(),
            result.data_ptr<float>(), outputs, 2,
            static_cast<float>(scale));
    } else {
        bias_pool_sum_kernel<false><<<batch, 256, 0, stream>>>(
            linear.data_ptr<float>(), bias.data_ptr<float>(),
            result.data_ptr<float>(), outputs,
            static_cast<int>(kernel_size), static_cast<float>(scale));
    }

    C10_CUDA_KERNEL_LAUNCH_CHECK();
    return {result, linear};
}
"""

_ops = load_inline(
    name="lt_tuned_linear_fused_pool_sum_v5",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["linear_pool_sum_cuda"],
    # Let PyTorch select the C++ standard required by its installed headers.
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    extra_ldflags=["-lcublasLt"],
    verbose=False,
)


class _LinearPoolSum(torch.autograd.Function):
    @staticmethod
    def forward(ctx, x, weight, bias, kernel_size, scale):
        result, linear = _ops.linear_pool_sum_cuda(
            x, weight, bias, kernel_size, scale
        )
        ctx.kernel_size = kernel_size
        ctx.scale = scale
        if any(ctx.needs_input_grad[:3]):
            ctx.save_for_backward(x, weight, bias, linear)
        return result

    @staticmethod
    def backward(ctx, grad_output):
        x, weight, bias, linear = ctx.saved_tensors
        _, indices = F.max_pool1d(
            (linear + bias).unsqueeze(1),
            kernel_size=ctx.kernel_size,
            stride=ctx.kernel_size,
            return_indices=True,
        )
        indices = indices.squeeze(1)
        grad_linear = torch.zeros_like(linear)
        values = (grad_output * ctx.scale).unsqueeze(1).expand_as(indices)
        grad_linear.scatter_(1, indices, values)

        grad_x = grad_linear.matmul(weight) if ctx.needs_input_grad[0] else None
        grad_weight = (
            grad_linear.t().matmul(x) if ctx.needs_input_grad[1] else None
        )
        grad_bias = grad_linear.sum(dim=0) if ctx.needs_input_grad[2] else None
        return grad_x, grad_weight, grad_bias, None, None


class ModelNew(nn.Module):
    def __init__(self, in_features, out_features, kernel_size, scale_factor):
        super().__init__()
        self.matmul = nn.Linear(in_features, out_features)
        self.max_pool = nn.MaxPool1d(kernel_size)
        self.kernel_size = (
            kernel_size[0]
            if isinstance(kernel_size, (tuple, list))
            else kernel_size
        )
        self.scale_factor = scale_factor

    def forward(self, x):
        weight = self.matmul.weight
        bias = self.matmul.bias
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and x.ndim == 2
            and weight.dtype == torch.float32
            and bias is not None
            and bias.dtype == torch.float32
            and not torch.is_autocast_enabled()
        ):
            return _LinearPoolSum.apply(
                x.contiguous(),
                weight.contiguous(),
                bias.contiguous(),
                self.kernel_size,
                self.scale_factor,
            )

        x = self.matmul(x)
        x = self.max_pool(x.unsqueeze(1)).squeeze(1)
        return x.sum(dim=1) * self.scale_factor
```