Optimize `GemPooling` using a "Warp-per-Channel" strategy.

**Performance Analysis:**
The previous "Block-per-Channel" approach was slower than PyTorch (0.77x). The spatial size `H*W` (e.g., 784) is too small to justify dedicating an entire Thread Block (e.g., 256 threads). This leads to low computational intensity per thread and excessive overhead from block scheduling and intra-block synchronization barriers.

**Optimization Strategy: Warp-per-Channel with Vectorized Loads**

The goal is to maximize "work per thread" and eliminate synchronization overhead.

1.  **Warp-Parallelism**: Assign one **Warp (32 threads)** to process one spatial channel `(H, W)`.
    *   A single CUDA Block (e.g., 128 threads) can now handle 4 channels simultaneously (4 warps).
    *   This reduces the grid size significantly and removes the need for `__syncthreads()` and Shared Memory, as Warp primitives (`__shfl_down_sync`) allow instant register-to-register communication.

2.  **Vectorized I/O (Float4)**: Use `float4` types to load 4 float values per instruction.
    *   With 32 threads, a single Warp iteration processes `32 * 4 = 128` elements.
    *   For `H*W=784`, the loop runs only ~6 times. This results in extremely high instruction throughput.

3.  **Fused Arithmetic**: Perform `pow(clamp(x, eps), p)` immediately in registers after loading.

4.  **Warp Reduction**: Perform the summation using efficient Warp Shuffle instructions. The first lane (lane 0) of the warp computes the final root mean and writes the result to global memory.

This approach matches the parallelism granularity to the problem size (Small Spatial, Large Batch/Channel), minimizing framework overhead and memory latency.
  
Here's an example to show you the syntax of inline embedding custom CUDA operators in torch: The example given architecture is:   
  
```python
import torch
import torch.nn as nn
import torch.nn.functional as F

# --- 用于基准测试的配置 ---
BATCH_SIZE = 256
CHANNELS = 1024
HEIGHT = 32
WIDTH = 32
SHAPE = (BATCH_SIZE, CHANNELS, HEIGHT, WIDTH)

# GeM Pooling 的超参数
P_VALUE = 3.0
EPS_VALUE = 1e-6

class GemPooling(nn.Module):
    """
    待优化的算子类：Generalized Mean Pooling (GeM)
    公式: f(X) = (mean(clamp(X, min=eps)^p))^(1/p)
    """
    def __init__(self, p=3.0, eps=1e-6):
        super(GemPooling, self).__init__()
        self.p = p
        self.eps = eps

    def forward(self, x: torch.Tensor) -> torch.Tensor:
        x = x.clamp(min=self.eps)
        x = x.pow(self.p)
        # 输入形状 (N, C, H, W) -> 输出形状 (N, C)
        x = x.mean(dim=[-2, -1])
        x = x.pow(1.0 / self.p)
        
        return x

class Model(nn.Module):
    def __init__(self, p=3.0, eps=1e-6):
        super(Model, self).__init__()
        self.gem_pooling = GemPooling(p=p, eps=eps)
    
    def forward(self, x: torch.Tensor) -> torch.Tensor:
        return self.gem_pooling(x)

def get_inputs():
    input_tensor = torch.randn(SHAPE, dtype=torch.float32)
    return [input_tensor.contiguous()]

def get_init_inputs():
    return [P_VALUE, EPS_VALUE]