Merge pull request 'feat:add good variance #26' (#224) from wut0n/GPUCodeForces:variance into main

S1/wut0n_#26/prompt.txt

@@ -0,0 +1,100 @@
+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:   
+  
+python
+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:   
+python
+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 []
+
+
+  
+You are given the following architecture:   
+  
+python
+import torch
+import torch.nn as nn
+
+class Model(nn.Module):
+“”"
+Variance implementation.
+Computes the variance of input tensors along the feature dimension.
+“”"
+def init(self):
+super(Model, self).init()
+
+def forward(self, x: torch.Tensor) -> torch.Tensor:  
+    """  
+    Compute variance of input tensor.  
+
+    Args:  
+        x (torch.Tensor): Input tensor [batch_size, feature_dim]  
+
+    Returns:  
+        torch.Tensor: Variance values [batch_size]  
+    """  
+    # Compute mean  
+    mean = torch.mean(x, dim=1, keepdim=True)  # [batch_size, 1]  
+      
+    # Compute squared differences  
+    diff = x - mean  # [batch_size, feature_dim]  
+    squared_diff = torch.pow(diff, 2)  # [batch_size, feature_dim]  
+      
+    # Compute variance  
+    variance = torch.mean(squared_diff, dim=1)  # [batch_size]  
+      
+    return variance  
+batch_size = 256
+feature_dim = 1024
+
+def get_inputs():
+# Generate input tensor with some variance
+x = torch.randn(batch_size, feature_dim) * 2.0 + 1.0 # mean=1, std=2
+return [x]
+
+def get_init_inputs():
+return [] # No special initialization inputs needed
+
+IMPORTANT: The variance computation involves multiple separate PyTorch operations (mean calculation, subtraction, squaring, and final mean) that can be fused into a single CUDA kernel for significant performance improvements. Consider two-pass algorithms (mean first, then variance) or online algorithms (Welford's algorithm) to achieve both high performance and numerical stability. Focus on creating a robust implementation that maintains perfect precision while delivering consistent speedups.

S1/wut0n_#26/run_code.py

@@ -0,0 +1,74 @@
+###########################################################
+# 性能和精度验证程序
+###########################################################
+import torch
+import torch.nn as nn
+import time
+from variance_torchcode import Model, get_inputs, get_init_inputs
+from variance_cudacode import ModelNew
+
+def run_benchmark():
+    # 检查 CUDA 是否可用
+    if not torch.cuda.is_available():
+        print("CUDA 不可用，请确保您有可用的 NVIDIA GPU 并已正确安装 PyTorch CUDA 版本。")
+        return
+    else:
+        device = torch.device("cuda")
+
+    # 初始化模型
+    init_inputs = get_init_inputs()
+    init_inputs = [
+        x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in init_inputs
+    ]
+    inputs = get_inputs()
+    inputs = [
+        x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in inputs
+    ]
+
+    torch_model = Model(*init_inputs).cuda()
+    cuda_model = ModelNew(*init_inputs).cuda()
+
+    torch_model.eval()
+    cuda_model.eval()
+
+    print("-------------------- 精度对齐验证 --------------------")
+    with torch.no_grad():
+        output_torch = torch_model( *inputs)
+        output_cuda = cuda_model(*inputs)
+
+    precision_flag = torch.allclose(output_torch, output_cuda,rtol=1e-03)
+    if precision_flag:
+        print("✅ 精度对齐：两个模型的输出结果非常接近。")
+    else:
+        print("❌ 精度不一致！")
+        
+    print("\n-------------------- 性能加速比测试 --------------------")
+    num_iterations = 100 
+    
+    # PyTorch 模型计时
+    torch.cuda.synchronize()
+    start_time = time.time()
+    for _ in range(num_iterations):
+        _ = torch_model(*inputs)
+    torch.cuda.synchronize()
+    torch_time = (time.time() - start_time) / num_iterations
+    
+    # 自定义 CUDA 内核计时
+    torch.cuda.synchronize()
+    start_time = time.time()
+    for _ in range(num_iterations):
+        _ = cuda_model(*inputs)
+    torch.cuda.synchronize()
+    cuda_time = (time.time() - start_time) / num_iterations
+    
+    print(f"PyTorch variance 平均执行时间: {torch_time:.6f} 秒")
+    print(f"自定义 CUDA variance 平均执行时间: {cuda_time:.6f} 秒")
+    speedup = 0
+    if cuda_time > 0:
+        speedup = torch_time / cuda_time
+        print(f"加速比 (Speedup): {speedup:.2f}x")
+    else:
+        print("CUDA 内核执行时间为0，无法计算加速比。")
+    return precision_flag,speedup
+if __name__ == "__main__":
+    precision_flag,speedup = run_benchmark()

S1/wut0n_#26/variance_cudacode.py

@@ -0,0 +1,103 @@
+import torch
+from torch.utils.cpp_extension import load_inline
+
+variance_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+// Two-pass算法：精度与性能的完美平衡
+__global__ void variance_two_pass_kernel(
+    const float* __restrict__ x,
+    float* __restrict__ variance,
+    int batch_size,
+    int feature_dim
+) {
+    int sample_idx = blockIdx.x;
+    if (sample_idx >= batch_size) return;
+    
+    int tid = threadIdx.x;
+    int base = sample_idx * feature_dim;
+    
+    // 第一遍：计算均值
+    float sum = 0.0f;
+    for (int i = tid; i < feature_dim; i += blockDim.x) {
+        sum += x[base + i];
+    }
+    
+    // Warp级归约求和
+    for (int offset = 16; offset > 0; offset /= 2) {
+        sum += __shfl_down_sync(0xffffffff, sum, offset);
+    }
+    
+    float mean = 0.0f;
+    if (tid == 0) {
+        mean = sum / feature_dim;
+    }
+    mean = __shfl_sync(0xffffffff, mean, 0);  // 广播均值
+    
+    // 第二遍：计算方差
+    float var_sum = 0.0f;
+    for (int i = tid; i < feature_dim; i += blockDim.x) {
+        float diff = x[base + i] - mean;
+        var_sum += diff * diff;
+    }
+    
+    // Warp级归约求和
+    for (int offset = 16; offset > 0; offset /= 2) {
+        var_sum += __shfl_down_sync(0xffffffff, var_sum, offset);
+    }
+    
+    if (tid == 0) {
+        variance[sample_idx] = var_sum / feature_dim;
+    }
+}
+
+torch::Tensor variance_cuda(
+    torch::Tensor x
+) {
+    auto x_contig = x.contiguous();
+    
+    int batch_size = x_contig.size(0);
+    int feature_dim = x_contig.size(1);
+    
+    auto variance = torch::zeros({batch_size}, x_contig.options());
+    
+    // Two-pass优化：最佳平衡点
+    const int block_size = 32;  // warp大小
+    
+    variance_two_pass_kernel<<<batch_size, block_size>>>(
+        x_contig.data_ptr<float>(),
+        variance.data_ptr<float>(),
+        batch_size,
+        feature_dim
+    );
+    
+    return variance;
+}
+"""
+
+variance_cpp_source = """
+torch::Tensor variance_cuda(torch::Tensor x);
+"""
+
+# 编译CUDA代码
+variance = load_inline(
+    name="variance",
+    cpp_sources=variance_cpp_source,
+    cuda_sources=variance_source,
+    functions=["variance_cuda"],
+    extra_cuda_cflags=[
+        "-O3", 
+        "--use_fast_math",
+        "-std=c++17"
+    ],
+    verbose=True
+)
+
+class ModelNew(torch.nn.Module):
+    def __init__(self):
+        super(ModelNew, self).__init__()
+        self.variance = variance
+
+    def forward(self, x):
+        return self.variance.variance_cuda(x)

S1/wut0n_#26/variance_torchcode.py

@@ -0,0 +1,43 @@
+import torch
+import torch.nn as nn
+
+class Model(nn.Module):
+    """
+    Variance implementation.
+    Computes the variance of input tensors along the feature dimension.
+    """
+    def __init__(self):
+        super(Model, self).__init__()
+    
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        """
+        Compute variance of input tensor.
+
+        Args:
+            x (torch.Tensor): Input tensor [batch_size, feature_dim]
+
+        Returns:
+            torch.Tensor: Variance values [batch_size]
+        """
+        # Compute mean
+        mean = torch.mean(x, dim=1, keepdim=True)  # [batch_size, 1]
+        
+        # Compute squared differences
+        diff = x - mean  # [batch_size, feature_dim]
+        squared_diff = torch.pow(diff, 2)  # [batch_size, feature_dim]
+        
+        # Compute variance
+        variance = torch.mean(squared_diff, dim=1)  # [batch_size]
+        
+        return variance
+
+batch_size = 256
+feature_dim = 1024
+
+def get_inputs():
+    # Generate input tensor with some variance
+    x = torch.randn(batch_size, feature_dim) * 2.0 + 1.0  # mean=1, std=2
+    return [x]
+
+def get_init_inputs():
+    return []  # No special initialization inputs needed