You write custom CUDA kernels to replace the pytorch operators in the given architecture to get speedups. 

    You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.


        Here's an example to show you the syntax of inline embedding custom CUDA operators in torch: The example given architecture is: 

        ``` 

        # 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


class Model(nn.Module):
    def __init__(self) -> None:
        super().__init__()

    def forward(self, a, b):
        return a + b


def get_inputs():
    # randomly generate input tensors based on the model architecture
    a = torch.randn(1, 128).cuda()
    b = torch.randn(1, 128).cuda()
    return [a, b]


def get_init_inputs():
    # randomly generate tensors required for initialization based on the model architecture
    return []

        ``` 

        The example new arch with custom CUDA kernels looks like this:
        ```
        # 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

# Define the custom CUDA kernel for element-wise addition
elementwise_add_source = """
#include <torch/extension.h>
#include <cuda_runtime.h>

__global__ void elementwise_add_kernel(const float* a, const float* b, float* out, int size) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        out[idx] = a[idx] + b[idx];
    }
}

torch::Tensor elementwise_add_cuda(torch::Tensor a, torch::Tensor b) {
    auto size = a.numel();
    auto out = torch::zeros_like(a);

    const int block_size = 256;
    const int num_blocks = (size + block_size - 1) / block_size;

    elementwise_add_kernel<<<num_blocks, block_size>>>(a.data_ptr<float>(), b.data_ptr<float>(), out.data_ptr<float>(), size);

    return out;
}
"""

elementwise_add_cpp_source = (
    "torch::Tensor elementwise_add_cuda(torch::Tensor a, torch::Tensor b);"
)

# Compile the inline CUDA code for element-wise addition
elementwise_add = load_inline(
    name="elementwise_add",
    cpp_sources=elementwise_add_cpp_source,
    cuda_sources=elementwise_add_source,
    functions=["elementwise_add_cuda"],
    verbose=True,
    extra_cflags=[""],
    extra_ldflags=[""],
)


class ModelNew(nn.Module):
    def __init__(self) -> None:
        super().__init__()
        self.elementwise_add = elementwise_add

    def forward(self, a, b):
        return self.elementwise_add.elementwise_add_cuda(a, b)

        ``` 

        
    You are given the following architecture: 

    ```
    # 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

class Model(nn.Module):
    """
    This model performs a convolution, average pooling, applies sigmoid, and sums the result.
    """
    def __init__(self, in_channels, out_channels, kernel_size, pool_kernel_size):
        super(Model, self).__init__()
        self.conv = nn.Conv2d(in_channels, out_channels, kernel_size)
        self.avg_pool = nn.AvgPool2d(pool_kernel_size)

    def forward(self, x):
        x = self.conv(x)
        x = self.avg_pool(x)
        x = torch.sigmoid(x)
        x = torch.sum(x, dim=[1,2,3]) # Sum over all spatial dimensions
        return x

batch_size = 128
in_channels = 8
out_channels = 64
height, width = 384, 384
kernel_size = 3
pool_kernel_size = 4

def get_inputs():
    return [torch.rand(batch_size, in_channels, height, width)]

def get_init_inputs():
    return [in_channels, out_channels, kernel_size, pool_kernel_size]
    ```
    
Optimize the architecture named Model with custom CUDA operators! Name your optimized output architecture ModelNew. Output the new code in codeblocks. Please generate real code, NOT pseudocode, make sure the code compiles and is fully functional. Just output the new model code, no other text, and NO testing code! 



This is turn 6 of 6 for this task. Your previous solution (turn 5) was:
```python
(no code was extracted)
```
Result of evaluating it:
Your reply contained no usable ```python code block.

The fastest correct solution so far (2.061x speedup) is:
```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 kernels: AvgPool + sigmoid + per-batch global sum, for both
# channels_last (NHWC) and contiguous (NCHW) conv outputs. Reading the conv
# output in its native channels_last layout lets us run the cuDNN conv fully
# on its fastest NHWC tensor-core path and skip every layout transform.
_fused_cuda_src = r"""
#include <torch/extension.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda_runtime.h>

__device__ __forceinline__ float sigmoid_f(float v) {
    return 1.0f / (1.0f + expf(-v));
}

__device__ __forceinline__ void block_sum_reduce(float val, float* out) {
    __shared__ float smem[32];
    const int lane = threadIdx.x & 31;
    const int wid  = threadIdx.x >> 5;
    #pragma unroll
    for (int off = 16; off > 0; off >>= 1)
        val += __shfl_down_sync(0xffffffffu, val, off);
    if (lane == 0) smem[wid] = val;
    __syncthreads();
    if (wid == 0) {
        const int nw = (blockDim.x + 31) >> 5;
        float v = (lane < nw) ? smem[lane] : 0.0f;
        #pragma unroll
        for (int off = 16; off > 0; off >>= 1)
            v += __shfl_down_sync(0xffffffffu, v, off);
        if (lane == 0) atomicAdd(out, v);
    }
}

// ------------------------- channels_last (NHWC) -------------------------
// grid = (OH, N). Each block reduces one (b, oi) band: all channels x all oj.
// Consecutive threads map to consecutive channels -> fully coalesced loads.
__global__ void nhwc_pool4_sigmoid_sum_kernel(
    const float* __restrict__ x, float* __restrict__ out,
    const int C, const int H, const int W, const int OH, const int OW)
{
    const int b  = blockIdx.y;
    const int oi = blockIdx.x;
    const long WC = (long)W * C;
    const float* xrow = x + ((long)b * H + (long)oi * 4) * WC;

    const int step = blockDim.x;
    const int dc  = step % C;
    const int doj = step / C;
    int c  = (int)(threadIdx.x) % C;
    int oj = (int)(threadIdx.x) / C;

    float acc = 0.0f;
    while (oj < OW) {
        const float* p = xrow + (long)oj * 4 * C + c;
        float s0 = 0.0f, s1 = 0.0f;
        #pragma unroll
        for (int r = 0; r < 4; ++r) {
            const float* pr = p + r * WC;
            s0 += pr[0] + pr[2 * C];
            s1 += pr[C] + pr[3 * C];
        }
        acc += sigmoid_f((s0 + s1) * 0.0625f);
        c += dc;
        if (c >= C) { c -= C; ++oj; }
        oj += doj;
    }
    block_sum_reduce(acc, out + b);
}

__global__ void nhwc_poolP_sigmoid_sum_kernel(
    const float* __restrict__ x, float* __restrict__ out,
    const int C, const int H, const int W, const int OH, const int OW,
    const int P, const float scale)
{
    const int b  = blockIdx.y;
    const int oi = blockIdx.x;
    const long WC = (long)W * C;
    const float* xrow = x + ((long)b * H + (long)oi * P) * WC;

    const int step = blockDim.x;
    const int dc  = step % C;
    const int doj = step / C;
    int c  = (int)(threadIdx.x) % C;
    int oj = (int)(threadIdx.x) / C;

    float acc = 0.0f;
    while (oj < OW) {
        const float* p = xrow + (long)oj * P * C + c;
        float s = 0.0f;
        for (int r = 0; r < P; ++r) {
            const float* pr = p + r * WC;
            for (int q = 0; q < P; ++q) s += pr[q * C];
        }
        acc += sigmoid_f(s * scale);
        c += dc;
        if (c >= C) { c -= C; ++oj; }
        oj += doj;
    }
    block_sum_reduce(acc, out + b);
}

// --------------------------- contiguous (NCHW) ---------------------------
// grid = (channel chunks, N). Incremental (oi, oj) updates: no div in loop.
__global__ void nchw_pool4_sigmoid_sum_kernel(
    const float* __restrict__ x, float* __restrict__ out,
    const int C, const int H, const int W, const int OH, const int OW,
    const int c_per_block)
{
    const int b  = blockIdx.y;
    const int c0 = blockIdx.x * c_per_block;
    const int c1 = min(c0 + c_per_block, C);

    const float* xb = x + (long)b * C * H * W;
    const long cstride = (long)H * W;
    const int step = blockDim.x;
    const int doi = step / OW;
    const int doj = step - doi * OW;

    float acc = 0.0f;
    for (int c = c0; c < c1; ++c) {
        const float* xc = xb + c * cstride;
        int idx = threadIdx.x;
        int oi = idx / OW;
        int oj = idx - oi * OW;
        while (oi < OH) {
            const float* p = xc + (long)oi * 4 * W + oj * 4;
            float s = 0.0f;
            #pragma unroll
            for (int r = 0; r < 4; ++r) {
                const float* pr = p + r * W;
                s += (pr[0] + pr[1]) + (pr[2] + pr[3]);
            }
            acc += sigmoid_f(s * 0.0625f);
            oj += doj;
            if (oj >= OW) { oj -= OW; ++oi; }
            oi += doi;
        }
    }
    block_sum_reduce(acc, out + b);
}

__global__ void nchw_poolP_sigmoid_sum_kernel(
    const float* __restrict__ x, float* __restrict__ out,
    const int C, const int H, const int W, const int OH, const int OW,
    const int P, const float scale, const int c_per_block)
{
    const int b  = blockIdx.y;
    const int c0 = blockIdx.x * c_per_block;
    const int c1 = min(c0 + c_per_block, C);

    const float* xb = x + (long)b * C * H * W;
    const long cstride = (long)H * W;
    const int nwin = OH * OW;

    float acc = 0.0f;
    for (int c = c0; c < c1; ++c) {
        const float* xc = xb + c * cstride;
        for (int idx = threadIdx.x; idx < nwin; idx += blockDim.x) {
            const int oi = idx / OW;
            const int oj = idx - oi * OW;
            const float* p = xc + (long)oi * P * W + oj * P;
            float s = 0.0f;
            for (int r = 0; r < P; ++r) {
                const float* pr = p + r * W;
                for (int q = 0; q < P; ++q) s += pr[q];
            }
            acc += sigmoid_f(s * scale);
        }
    }
    block_sum_reduce(acc, out + b);
}

torch::Tensor pool_sigmoid_sum_cuda(torch::Tensor x, int64_t pool) {
    TORCH_CHECK(x.is_cuda(), "input must be a CUDA tensor");
    TORCH_CHECK(x.scalar_type() == at::kFloat, "input must be float32");
    TORCH_CHECK(x.dim() == 4, "input must be NCHW/NHWC 4D");

    const int N = (int)x.size(0);
    const int C = (int)x.size(1);
    const int H = (int)x.size(2);
    const int W = (int)x.size(3);
    const int P = (int)pool;
    const int OH = H / P;
    const int OW = W / P;

    auto out = at::zeros({N}, x.options());
    if (N == 0 || C == 0 || OH <= 0 || OW <= 0) return out;

    const int block = 256;
    auto stream = at::cuda::getCurrentCUDAStream();

    at::Tensor xc = x;
    bool channels_last = x.is_contiguous(at::MemoryFormat::ChannelsLast);
    if (!channels_last && !x.is_contiguous()) {
        xc = x.contiguous(at::MemoryFormat::ChannelsLast);
        channels_last = true;
    }

    if (channels_last) {
        const dim3 grid(OH, N);
        if (P == 4) {
            nhwc_pool4_sigmoid_sum_kernel<<<grid, block, 0, stream>>>(
                xc.data_ptr<float>(), out.data_ptr<float>(), C, H, W, OH, OW);
        } else {
            nhwc_poolP_sigmoid_sum_kernel<<<grid, block, 0, stream>>>(
                xc.data_ptr<float>(), out.data_ptr<float>(), C, H, W, OH, OW,
                P, 1.0f / (float)(P * P));
        }
    } else {
        int chunks = (1024 + N - 1) / N;
        if (chunks < 1) chunks = 1;
        if (chunks > C) chunks = C;
        const int c_per_block = (C + chunks - 1) / chunks;
        chunks = (C + c_per_block - 1) / c_per_block;
        const dim3 grid(chunks, N);
        if (P == 4) {
            nchw_pool4_sigmoid_sum_kernel<<<grid, block, 0, stream>>>(
                xc.data_ptr<float>(), out.data_ptr<float>(),
                C, H, W, OH, OW, c_per_block);
        } else {
            nchw_poolP_sigmoid_sum_kernel<<<grid, block, 0, stream>>>(
                xc.data_ptr<float>(), out.data_ptr<float>(),
                C, H, W, OH, OW, P, 1.0f / (float)(P * P), c_per_block);
        }
    }
    return out;
}
"""

_fused_cpp_src = "torch::Tensor pool_sigmoid_sum_cuda(torch::Tensor x, int64_t pool);"

_fused_pool_sigmoid_sum = load_inline(
    name="fused_pool_sigmoid_sum_v2",
    cpp_sources=_fused_cpp_src,
    cuda_sources=_fused_cuda_src,
    functions=["pool_sigmoid_sum_cuda"],
    verbose=False,
    extra_cuda_cflags=["-O3"],
)


class ModelNew(nn.Module):
    """
    Optimized model:
      * Conv2d runs on cuDNN's autotuned TF32 tensor-core path in native
        channels_last (NHWC) layout, avoiding the per-call NCHW<->NHWC layout
        transforms of both the activations and the filter.
      * AvgPool + sigmoid + global per-batch sum are fused into a single
        custom CUDA kernel that reads the conv output exactly once (in its
        native layout) and writes just one float per batch element.
    """
    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.pool_kernel_size = pool_kernel_size
        torch.backends.cudnn.benchmark = True
        torch.backends.cudnn.benchmark_limit = 0  # exhaustive engine search
        torch.backends.cudnn.allow_tf32 = True

    def forward(self, x):
        if x.is_cuda and x.dtype == torch.float32 and x.dim() == 4:
            xcl = x.contiguous(memory_format=torch.channels_last)
            wcl = self.conv.weight.contiguous(memory_format=torch.channels_last)
            y = F.conv2d(xcl, wcl, self.conv.bias, self.conv.stride,
                         self.conv.padding, self.conv.dilation, self.conv.groups)
            return _fused_pool_sigmoid_sum.pool_sigmoid_sum_cuda(
                y, self.pool_kernel_size)
        # Fallback (CPU / other dtypes): original op sequence.
        x = self.conv(x)
        x = F.avg_pool2d(x, self.pool_kernel_size)
        x = torch.sigmoid(x)
        return torch.sum(x, dim=[1, 2, 3])
```

Write an improved complete ModelNew: fix any error, and if the solution is already correct, make it faster. Output the new code in one ```python codeblock. Just the model code, no other text, and NO testing code.


Precision rule: the timed forward must compute in FP32 (TF32 tensor-core math is allowed, as in the PyTorch reference). Do not cast inputs, weights or intermediate results to float16 or bfloat16 and do not call half-precision kernels in forward().
