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

_fused_cuda_source = r"""
#include <torch/extension.h>
#include <cuda_runtime.h>
#include <mma.h>

using namespace nvcuda;

constexpr int IN_C = 8;
constexpr int OUT_C = 64;
constexpr int K = 3;
constexpr int KTOT = IN_C * K * K;      // 72
constexpr int POOL = 4;
constexpr int POS = POOL * POOL;        // 16
constexpr int WARPS = 6;
constexpr int THREADS = WARPS * 32;
constexpr int BUF_FLOATS = KTOT * POS;  // 1152, also enough for OUT_C x POS = 1024

__global__ void fused_conv3_pool4_sigmoid_sum_kernel(
    const float* __restrict__ x,
    const float* __restrict__ w,
    const float* __restrict__ b,
    float* __restrict__ out,
    int H, int W,
    int pool_h, int pool_w,
    int spatial_blocks)
{
    __shared__ float s_w[OUT_C][KTOT];
    __shared__ float s_buf[WARPS][BUF_FLOATS];

    const int tid = threadIdx.x;
    for (int idx = tid; idx < OUT_C * KTOT; idx += THREADS) {
        reinterpret_cast<float*>(s_w)[idx] = w[idx];
    }
    __syncthreads();

    const int n = blockIdx.y;
    const int warp = tid >> 5;
    const int lane = tid & 31;

    const int r0 = lane;
    const int r1 = lane + 32;
    const float bias0 = b[r0];
    const float bias1 = b[r1];

    const float* x_n = x + (long long)n * IN_C * H * W;
    float* warp_buf = s_buf[warp];

    const int cells = pool_h * pool_w;
    const int stride = spatial_blocks * WARPS;

    float acc = 0.0f;

    for (int cell = blockIdx.x * WARPS + warp; cell < cells; cell += stride) {
        const int ph = cell / pool_w;
        const int pw = cell - ph * pool_w;
        const int base_y = ph * POOL;
        const int base_x = pw * POOL;

        // Gather B as column-major (K x POS): buf[p*KTOT + k] = B[k][p].
        for (int idx = lane; idx < KTOT * POS; idx += 32) {
            const int p = idx / KTOT;
            const int k = idx - p * KTOT;

            const int i = p >> 2;
            const int j = p & 3;
            const int y = base_y + i;
            const int xx = base_x + j;

            const int ic = k / 9;
            const int r = k - ic * 9;
            const int ky = r / 3;
            const int kx = r - ky * 3;

            warp_buf[idx] = x_n[((long long)ic * H + (y + ky)) * W + (xx + kx)];
        }
        __syncwarp();

        wmma::fragment<wmma::accumulator, 16, 16, 8, float> c_frag[4];
        #pragma unroll
        for (int m = 0; m < 4; ++m) {
            wmma::fill_fragment(c_frag[m], 0.0f);
        }

        #pragma unroll
        for (int k0 = 0; k0 < KTOT; k0 += 8) {
            wmma::fragment<wmma::matrix_b, 16, 16, 8, wmma::precision::tf32, wmma::col_major> b_frag;
            wmma::load_matrix_sync(b_frag, warp_buf + k0, KTOT);

            #pragma unroll
            for (int m = 0; m < 4; ++m) {
                wmma::fragment<wmma::matrix_a, 16, 16, 8, wmma::precision::tf32, wmma::row_major> a_frag;
                wmma::load_matrix_sync(a_frag, &s_w[m * 16][k0], KTOT);
                wmma::mma_sync(c_frag[m], a_frag, b_frag, c_frag[m]);
            }
        }
        __syncwarp();

        // Reuse the same warp buffer for C, stored as row-major OUT_C x POS.
        float* s_c = warp_buf;
        #pragma unroll
        for (int m = 0; m < 4; ++m) {
            wmma::store_matrix_sync(s_c + (m * 16) * POS, c_frag[m], POS, wmma::mem_row_major);
        }
        __syncwarp();

        float sum0 = 0.0f;
        float sum1 = 0.0f;
        #pragma unroll
        for (int p = 0; p < POS; ++p) {
            sum0 += s_c[r0 * POS + p];
            sum1 += s_c[r1 * POS + p];
        }

        const float inv_pool = 1.0f / (float)POS;
        acc += 1.0f / (1.0f + expf(-(bias0 + sum0 * inv_pool)));
        acc += 1.0f / (1.0f + expf(-(bias1 + sum1 * inv_pool)));
        __syncwarp();
    }

    #pragma unroll
    for (int off = 16; off > 0; off >>= 1) {
        acc += __shfl_down_sync(0xffffffffu, acc, off);
    }
    if (lane == 0) {
        atomicAdd(out + n, acc);
    }
}

torch::Tensor fused_conv3_pool4_sigmoid_sum_cuda(torch::Tensor x, torch::Tensor w, torch::Tensor b) {
    TORCH_CHECK(x.is_cuda(), "x must be CUDA");
    TORCH_CHECK(w.is_cuda(), "w must be CUDA");
    TORCH_CHECK(b.is_cuda(), "b must be CUDA");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "x must be float32");
    TORCH_CHECK(w.scalar_type() == at::kFloat, "w must be float32");
    TORCH_CHECK(b.scalar_type() == at::kFloat, "b must be float32");
    TORCH_CHECK(x.is_contiguous(), "x must be contiguous");
    TORCH_CHECK(w.is_contiguous(), "w must be contiguous");
    TORCH_CHECK(b.is_contiguous(), "b must be contiguous");

    const int N = x.size(0);
    const int H = x.size(2);
    const int W = x.size(3);
    const int conv_h = H - 2;
    const int conv_w = W - 2;
    const int pool_h = conv_h / 4;
    const int pool_w = conv_w / 4;

    auto out = torch::zeros({N}, x.options());

    int sm_count = 0;
    cudaDeviceGetAttribute(&sm_count, cudaDevAttrMultiProcessorCount, x.get_device());
    int spatial_blocks = (sm_count * 16 + N - 1) / N;
    if (spatial_blocks < 1) spatial_blocks = 1;
    const int max_spatial_blocks = (pool_h * pool_w + WARPS - 1) / WARPS;
    if (spatial_blocks > max_spatial_blocks) spatial_blocks = max_spatial_blocks;

    dim3 grid(spatial_blocks, N);
    dim3 block(THREADS);

    fused_conv3_pool4_sigmoid_sum_kernel<<<grid, block>>>(
        x.data_ptr<float>(),
        w.data_ptr<float>(),
        b.data_ptr<float>(),
        out.data_ptr<float>(),
        H, W,
        pool_h, pool_w,
        spatial_blocks);

    return out;
}
"""

_fused_cpp_source = "torch::Tensor fused_conv3_pool4_sigmoid_sum_cuda(torch::Tensor x, torch::Tensor w, torch::Tensor b);"

_fused = load_inline(
    name="fused_conv3_pool4_sigmoid_sum",
    cpp_sources=_fused_cpp_source,
    cuda_sources=_fused_cuda_source,
    functions=["fused_conv3_pool4_sigmoid_sum_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3", "--use_fast_math"],
)


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

    def forward(self, x):
        if (
            x.is_cuda
            and x.dtype == torch.float32
            and self.in_channels == 8
            and self.out_channels == 64
            and self.kernel_size == 3
            and self.pool_kernel_size == 4
            and x.dim() == 4
            and x.size(1) == 8
            and x.size(2) >= 6
            and x.size(3) >= 6
        ):
            return _fused.fused_conv3_pool4_sigmoid_sum_cuda(
                x.contiguous(),
                self.conv.weight.contiguous(),
                self.conv.bias.contiguous(),
            )

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