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


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

torch::Tensor pooled_conv_sum_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias);
"""

_cuda_source = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <c10/cuda/CUDAGuard.h>
#include <c10/cuda/CUDAException.h>
#include <cuda_runtime.h>
#include <mma.h>

namespace {

constexpr int CHANNELS = 8;
constexpr int OUTPUTS = 64;
constexpr int K = CHANNELS * 9;
constexpr int TILE = 64;

__global__ void lower_pool(
    const float* __restrict__ x,
    float* __restrict__ lowered,
    int height,
    int width,
    int out_width,
    int positions,
    int pitch)
{
    const int p = blockIdx.x * blockDim.x + threadIdx.x;
    if (p >= pitch)
        return;

    const int channel = blockIdx.y;
    const int batch = blockIdx.z;

    float* dst = lowered +
        (static_cast<int64_t>(batch) * CHANNELS + channel) * 9 * pitch
        + p;

    if (p >= positions) {
        #pragma unroll
        for (int k = 0; k < 9; ++k)
            dst[static_cast<int64_t>(k) * pitch] = 0.0f;
        return;
    }

    const int iy = (p / out_width) * 4;
    const int ix = (p % out_width) * 4;
    const float* src = x +
        ((static_cast<int64_t>(batch) * CHANNELS + channel) * height + iy)
        * width + ix;

    float horizontal[6][3];

    #pragma unroll
    for (int r = 0; r < 6; ++r) {
        const float* row = src + r * width;
        float a, b, c, d, e, f;

        if ((width & 3) == 0) {
            const float4 first =
                *reinterpret_cast<const float4*>(row);
            const float2 last =
                *reinterpret_cast<const float2*>(row + 4);
            a = first.x;
            b = first.y;
            c = first.z;
            d = first.w;
            e = last.x;
            f = last.y;
        } else {
            a = row[0];
            b = row[1];
            c = row[2];
            d = row[3];
            e = row[4];
            f = row[5];
        }

        horizontal[r][0] = ((a + b) + c) + d;
        horizontal[r][1] = ((b + c) + d) + e;
        horizontal[r][2] = ((c + d) + e) + f;
    }

    #pragma unroll
    for (int r = 0; r < 3; ++r) {
        #pragma unroll
        for (int c = 0; c < 3; ++c) {
            const float value =
                ((horizontal[r][c] + horizontal[r + 1][c]) +
                 horizontal[r + 2][c]) + horizontal[r + 3][c];

            dst[static_cast<int64_t>(r * 3 + c) * pitch] =
                value * 0.0625f;
        }
    }
}

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

// Each block computes all 64 output channels for 64 pooled positions,
// then immediately applies bias, sigmoid, and reduction.
__global__ void fused_matmul_sigmoid_sum(
    const float* __restrict__ lowered,
    const float* __restrict__ weight,
    const float* __restrict__ bias,
    float* __restrict__ partial,
    int positions,
    int pitch,
    int tiles)
{
    const int tid = threadIdx.x;
    const int warp = tid >> 5;
    const int lane = tid & 31;
    const int tile = blockIdx.x;
    const int batch = blockIdx.y;
    const int first_position = tile * TILE;

    const float* input =
        lowered + static_cast<int64_t>(batch) * K * pitch;

    float total = 0.0f;

#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
    namespace wmma = nvcuda::wmma;

    // Reused for the matrix product after both operand tiles are consumed.
    __shared__ __align__(32) float storage[OUTPUTS * K + K * TILE];
    float* shared_weight = storage;
    float* shared_input = storage + OUTPUTS * K;

    for (int i = tid; i < OUTPUTS * K; i += 128)
        shared_weight[i] = __float_to_tf32(weight[i]);

    for (int i = tid; i < K * TILE; i += 128) {
        const int k = i / TILE;
        const int p = i % TILE;
        shared_input[i] = __float_to_tf32(
            input[static_cast<int64_t>(k) * pitch + first_position + p]);
    }
    __syncthreads();

    wmma::fragment<
        wmma::accumulator, 16, 16, 8, float
    > accum[4];

    #pragma unroll
    for (int m = 0; m < 4; ++m)
        wmma::fill_fragment(accum[m], 0.0f);

    #pragma unroll
    for (int k = 0; k < K; k += 8) {
        wmma::fragment<
            wmma::matrix_b, 16, 16, 8,
            wmma::precision::tf32, wmma::row_major
        > b;

        wmma::load_matrix_sync(
            b, shared_input + k * TILE + warp * 16, TILE);

        #pragma unroll
        for (int m = 0; m < 4; ++m) {
            wmma::fragment<
                wmma::matrix_a, 16, 16, 8,
                wmma::precision::tf32, wmma::row_major
            > a;

            wmma::load_matrix_sync(
                a, shared_weight + m * 16 * K + k, K);
            wmma::mma_sync(accum[m], a, b, accum[m]);
        }
    }

    __syncthreads();

    #pragma unroll
    for (int m = 0; m < 4; ++m) {
        wmma::store_matrix_sync(
            storage + m * 16 * TILE + warp * 16,
            accum[m], TILE, wmma::mem_row_major);
    }
    __syncthreads();

    #pragma unroll
    for (int j = 0; j < OUTPUTS * TILE / 128; ++j) {
        const int i = tid + j * 128;
        const int channel = i / TILE;
        const int p = i % TILE;

        if (first_position + p < positions) {
            const float value = storage[i] + bias[channel];
            total += 1.0f / (1.0f + __expf(-value));
        }
    }
#else
    // Functional fallback for GPUs without TF32 tensor cores.
    for (int i = tid; i < OUTPUTS * TILE; i += 128) {
        const int channel = i / TILE;
        const int p = first_position + i % TILE;
        if (p < positions) {
            float value = 0.0f;
            #pragma unroll
            for (int k = 0; k < K; ++k) {
                value = fmaf(
                    weight[channel * K + k],
                    input[static_cast<int64_t>(k) * pitch + p],
                    value);
            }
            value += bias[channel];
            total += 1.0f / (1.0f + __expf(-value));
        }
    }
#endif

    __shared__ float warp_totals[4];
    total = warp_sum(total);
    if (lane == 0)
        warp_totals[warp] = total;
    __syncthreads();

    if (warp == 0) {
        total = lane < 4 ? warp_totals[lane] : 0.0f;
        total = warp_sum(total);
        if (lane == 0)
            partial[static_cast<int64_t>(batch) * tiles + tile] = total;
    }
}

__global__ void finish_sum(
    const float* __restrict__ partial,
    float* __restrict__ output,
    int tiles)
{
    const int batch = blockIdx.x;
    const int lane = threadIdx.x & 31;
    const int warp = threadIdx.x >> 5;

    float total = 0.0f;
    for (int i = threadIdx.x; i < tiles; i += blockDim.x)
        total += partial[static_cast<int64_t>(batch) * tiles + i];

    __shared__ float warp_totals[8];
    total = warp_sum(total);
    if (lane == 0)
        warp_totals[warp] = total;
    __syncthreads();

    if (warp == 0) {
        total = lane < 8 ? warp_totals[lane] : 0.0f;
        total = warp_sum(total);
        if (lane == 0)
            output[batch] = total;
    }
}

} // namespace

torch::Tensor pooled_conv_sum_cuda(
    torch::Tensor x,
    torch::Tensor weight,
    torch::Tensor bias)
{
    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 tensors");
    TORCH_CHECK(x.is_contiguous() && weight.is_contiguous() &&
                bias.is_contiguous(), "Expected contiguous tensors");
    TORCH_CHECK(x.dim() == 4 && x.size(1) == CHANNELS,
                "Expected NCHW input with eight channels");
    TORCH_CHECK(weight.dim() == 4 &&
                weight.size(0) == OUTPUTS &&
                weight.size(1) == CHANNELS &&
                weight.size(2) == 3 && weight.size(3) == 3 &&
                bias.numel() == OUTPUTS,
                "Expected 64x8x3x3 weights and 64 biases");
    TORCH_CHECK(x.size(2) >= 6 && x.size(3) >= 6,
                "Input is too small");

    const c10::cuda::CUDAGuard guard(x.device());

    const int batch = x.size(0);
    const int height = x.size(2);
    const int width = x.size(3);
    const int out_height = (height - 2) / 4;
    const int out_width = (width - 2) / 4;
    const int positions = out_height * out_width;
    const int tiles = (positions + TILE - 1) / TILE;
    const int pitch = tiles * TILE;

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

    auto lowered = torch::empty({batch, K, pitch}, x.options());
    auto partial = torch::empty({batch, tiles}, x.options());
    const auto stream = at::cuda::getCurrentCUDAStream();

    lower_pool<<<dim3((pitch + 255) / 256, CHANNELS, batch),
                 256, 0, stream>>>(
        x.data_ptr<float>(),
        lowered.data_ptr<float>(),
        height, width, out_width, positions, pitch);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    fused_matmul_sigmoid_sum<<<dim3(tiles, batch), 128, 0, stream>>>(
        lowered.data_ptr<float>(),
        weight.data_ptr<float>(),
        bias.data_ptr<float>(),
        partial.data_ptr<float>(),
        positions, pitch, tiles);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    finish_sum<<<batch, 256, 0, stream>>>(
        partial.data_ptr<float>(), result.data_ptr<float>(), tiles);
    C10_CUDA_KERNEL_LAUNCH_CHECK();

    return result;
}
"""

_ops = load_inline(
    name="pooled_conv_fused_sigmoid_sum_v3",
    cpp_sources=_cpp_source,
    cuda_sources=_cuda_source,
    functions=["pooled_conv_sum_cuda"],
    extra_cflags=["-O3"],
    extra_cuda_cflags=["-O3"],
    verbose=False,
)


class ModelNew(nn.Module):
    def __init__(self, in_channels, out_channels, kernel_size, pool_kernel_size):
        super().__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.avg_pool = nn.AvgPool2d(pool_kernel_size)

        pool_size = (
            (pool_kernel_size, pool_kernel_size)
            if isinstance(pool_kernel_size, int)
            else tuple(pool_kernel_size)
        )
        self._specialized = (
            in_channels == 8
            and out_channels == 64
            and self.conv.kernel_size == (3, 3)
            and pool_size == (4, 4)
        )

    def forward(self, x):
        if (
            self._specialized
            and x.is_cuda
            and x.ndim == 4
            and x.dtype == torch.float32
            and self.conv.weight.dtype == torch.float32
            and self.conv.bias.dtype == torch.float32
            and not torch.is_autocast_enabled()
        ):
            return _ops.pooled_conv_sum_cuda(
                x.contiguous(),
                self.conv.weight.contiguous(),
                self.conv.bias.contiguous(),
            )

        return torch.sigmoid(self.avg_pool(self.conv(x))).sum(dim=(1, 2, 3))
```