Merge pull request 'feat:add high performance weighted_sum #25' (#223) from wut0n/GPUCodeForces:weighted_sum into main

S1/wut0n_#25/prompt.txt

@@ -0,0 +1,98 @@
+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):
+“”"
+Weighted Sum implementation.
+Computes the weighted sum of values using corresponding weights.
+“”"
+def init(self):
+super(Model, self).init()
+
+def forward(self, values: torch.Tensor, weights: torch.Tensor) -> torch.Tensor:  
+    """  
+    Compute weighted sum of values.  
+
+    Args:  
+        values (torch.Tensor): Input values [batch_size, feature_dim]  
+        weights (torch.Tensor): Corresponding weights [batch_size, feature_dim]  
+
+    Returns:  
+        torch.Tensor: Weighted sums [batch_size]  
+    """  
+    # Element-wise multiplication  
+    elementwise_product = values * weights  
+      
+    # Sum along feature dimension  
+    result = torch.sum(elementwise_product, dim=1)  
+      
+    return result  
+batch_size = 256
+feature_dim = 512
+
+def get_inputs():
+# Generate values and corresponding weights
+values = torch.randn(batch_size, feature_dim)
+weights = torch.rand(batch_size, feature_dim) # Random weights between 0 and 1
+return [values, weights]
+
+def get_init_inputs():
+return [] # No special initialization inputs needed
+
+IMPORTANT: The weighted sum computation involves two separate PyTorch operations (element-wise multiplication and reduction) that can be fused into a single CUDA kernel for significant performance improvements. Consider warp-level optimizations and efficient reduction techniques to achieve both high performance and accuracy. Focus on creating a robust implementation that maintains perfect precision while delivering consistent speedups.

S1/wut0n_#25/run_code.py

@@ -0,0 +1,74 @@
+###########################################################
+# 性能和精度验证程序
+###########################################################
+import torch
+import torch.nn as nn
+import time
+from weighted_sum_torchcode import Model, get_inputs, get_init_inputs
+from weighted_sum_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 weighted_sum 平均执行时间: {torch_time:.6f} 秒")
+    print(f"自定义 CUDA weighted_sum 平均执行时间: {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_#25/weighted_sum_cudacode.py

@@ -0,0 +1,90 @@
+import torch
+from torch.utils.cpp_extension import load_inline
+
+weighted_sum_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+// Warp优化版本：精度与性能的完美平衡
+__global__ void weighted_sum_warp_kernel(
+    const float* __restrict__ values,
+    const float* __restrict__ weights,
+    float* __restrict__ results,
+    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;
+    
+    // 高效的warp级处理
+    for (int i = tid; i < feature_dim; i += 32) {  // warp大小为32
+        sum += values[base + i] * weights[base + i];
+    }
+    
+    // Warp级归约
+    for (int offset = 16; offset > 0; offset /= 2) {
+        sum += __shfl_down_sync(0xffffffff, sum, offset);
+    }
+    
+    if (tid == 0) {
+        results[sample_idx] = sum;
+    }
+}
+
+torch::Tensor weighted_sum_cuda(
+    torch::Tensor values,
+    torch::Tensor weights
+) {
+    auto values_contig = values.contiguous();
+    auto weights_contig = weights.contiguous();
+    
+    int batch_size = values_contig.size(0);
+    int feature_dim = values_contig.size(1);
+    
+    auto results = torch::zeros({batch_size}, values.options());
+    
+    // Warp优化：最佳平衡点
+    const int block_size = 32;  // warp大小
+    
+    weighted_sum_warp_kernel<<<batch_size, block_size>>>(
+        values_contig.data_ptr<float>(),
+        weights_contig.data_ptr<float>(),
+        results.data_ptr<float>(),
+        batch_size,
+        feature_dim
+    );
+    
+    return results;
+}
+"""
+
+weighted_sum_cpp_source = """
+torch::Tensor weighted_sum_cuda(torch::Tensor values, torch::Tensor weights);
+"""
+
+# 编译CUDA代码
+weighted_sum = load_inline(
+    name="weighted_sum",
+    cpp_sources=weighted_sum_cpp_source,
+    cuda_sources=weighted_sum_source,
+    functions=["weighted_sum_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.weighted_sum = weighted_sum
+
+    def forward(self, values, weights):
+        return self.weighted_sum.weighted_sum_cuda(values, weights)

S1/wut0n_#25/weighted_sum_torchcode.py

@@ -0,0 +1,41 @@
+import torch
+import torch.nn as nn
+
+class Model(nn.Module):
+    """
+    Weighted Sum implementation.
+    Computes the weighted sum of values using corresponding weights.
+    """
+    def __init__(self):
+        super(Model, self).__init__()
+    
+    def forward(self, values: torch.Tensor, weights: torch.Tensor) -> torch.Tensor:
+        """
+        Compute weighted sum of values.
+
+        Args:
+            values (torch.Tensor): Input values [batch_size, feature_dim]
+            weights (torch.Tensor): Corresponding weights [batch_size, feature_dim]
+
+        Returns:
+            torch.Tensor: Weighted sums [batch_size]
+        """
+        # 逐元素乘法
+        elementwise_product = values * weights
+        
+        # 求和
+        result = torch.sum(elementwise_product, dim=1)
+        
+        return result
+
+batch_size = 1024
+feature_dim = 512
+
+def get_inputs():
+    # Generate values and corresponding weights
+    values = torch.randn(batch_size, feature_dim)
+    weights = torch.rand(batch_size, feature_dim)  # Random weights between 0 and 1
+    return [values, weights]
+
+def get_init_inputs():
+    return []  # No special initialization inputs needed