feat:add mahattan_gelu #63

S1/wut0n_#63/manhattan_gelu_cudacode.py

@@ -0,0 +1,332 @@
+import torch
+from torch.utils.cpp_extension import load_inline
+
+manhattan_gelu_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+// 纯CUDA实现：Manhattan + GELU融合
+__global__ void manhattan_gelu_kernel(
+    const float* __restrict__ x,
+    const float* __restrict__ y,
+    float* __restrict__ distances,
+    float* __restrict__ gelu_output,
+    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;
+    
+    // 使用共享内存进行归约
+    extern __shared__ float shared_sum[];
+    shared_sum[tid] = 0.0f;
+    
+    // 每个线程处理4个元素（float4向量化）
+    int stride = blockDim.x * 4;
+    for (int dim = tid * 4; dim < feature_dim; dim += stride) {
+        // 确保不越界
+        if (dim + 3 < feature_dim) {
+            float4 x_val = *reinterpret_cast<const float4*>(&x[base + dim]);
+            float4 y_val = *reinterpret_cast<const float4*>(&y[base + dim]);
+            
+            // 应用GELU激活
+            float4 x_activated;
+            // GELU(x) = 0.5x * (1 + erf(x/√2))
+            const float sqrt_2_over_pi = 0.7978845608028654f; // sqrt(2/π)
+            const float coeff = 0.044715f;
+            
+            // 对每个组件应用GELU
+            float x_temp = x_val.x;
+            float tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+            x_activated.x = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+            
+            x_temp = x_val.y;
+            tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+            x_activated.y = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+            
+            x_temp = x_val.z;
+            tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+            x_activated.z = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+            
+            x_temp = x_val.w;
+            tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+            x_activated.w = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+            
+            // 存储激活后的输出
+            *reinterpret_cast<float4*>(&gelu_output[base + dim]) = x_activated;
+            
+            // 计算Manhattan距离
+            shared_sum[tid] += fabsf(x_activated.x - y_val.x) + fabsf(x_activated.y - y_val.y) + 
+                              fabsf(x_activated.z - y_val.z) + fabsf(x_activated.w - y_val.w);
+        } else {
+            // 处理剩余元素
+            for (int i = dim; i < feature_dim; i++) {
+                float x_val = x[base + i];
+                float y_val = y[base + i];
+                
+                // 应用GELU激活
+                const float sqrt_2_over_pi = 0.7978845608028654f;
+                const float coeff = 0.044715f;
+                float tanh_arg = sqrt_2_over_pi * (x_val + coeff * x_val * x_val * x_val);
+                float x_activated = 0.5f * x_val * (1.0f + tanhf(tanh_arg));
+                gelu_output[base + i] = x_activated;
+                
+                // 计算Manhattan距离
+                float diff = x_activated - y_val;
+                shared_sum[tid] += fabsf(diff);
+            }
+        }
+    }
+    
+    __syncthreads();
+    
+    // 块内归约求和
+    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
+        if (tid < stride) {
+            shared_sum[tid] += shared_sum[tid + stride];
+        }
+        __syncthreads();
+    }
+    
+    // 第一个线程写入结果
+    if (tid == 0) {
+        distances[sample_idx] = shared_sum[0];
+    }
+}
+
+// Warp级优化版本 - Manhattan + GELU融合
+__global__ void manhattan_gelu_kernel_warp(
+    const float* __restrict__ x,
+    const float* __restrict__ y,
+    float* __restrict__ distances,
+    float* __restrict__ gelu_output,
+    int batch_size,
+    int feature_dim
+) {
+    int sample_idx = blockIdx.x;
+    if (sample_idx >= batch_size) return;
+    
+    int tid = threadIdx.x;
+    int warp_id = tid / 32;
+    int lane_id = tid % 32;
+    
+    int base = sample_idx * feature_dim;
+    
+    // Warp级计算累加和
+    float warp_sum = 0.0f;
+    
+    // 每个warp处理一部分特征
+    int elements_per_warp = (feature_dim + 8 - 1) / 8;
+    int start_dim = warp_id * elements_per_warp;
+    int end_dim = min(start_dim + elements_per_warp, feature_dim);
+    
+    for (int dim = start_dim + lane_id; dim < end_dim; dim += 32) {
+        float x_val = x[base + dim];
+        float y_val = y[base + dim];
+        
+        // 应用GELU激活
+        const float sqrt_2_over_pi = 0.7978845608028654f;
+        const float coeff = 0.044715f;
+        float tanh_arg = sqrt_2_over_pi * (x_val + coeff * x_val * x_val * x_val);
+        float x_activated = 0.5f * x_val * (1.0f + tanhf(tanh_arg));
+        gelu_output[base + dim] = x_activated;
+        
+        // 计算Manhattan距离
+        float diff = x_activated - y_val;
+        warp_sum += fabsf(diff);
+    }
+    
+    // Warp级归约求和
+    for (int offset = 16; offset > 0; offset /= 2) {
+        warp_sum += __shfl_down_sync(0xffffffff, warp_sum, offset);
+    }
+    
+    // 使用共享内存进行跨warp归约
+    extern __shared__ float shared_data[];
+    if (lane_id == 0) {
+        shared_data[warp_id] = warp_sum;
+    }
+    __syncthreads();
+    
+    // 第一个线程找到全局总和
+    if (tid == 0) {
+        float total_sum = 0.0f;
+        int num_warps = blockDim.x / 32;
+        for (int i = 0; i < num_warps; i++) {
+            total_sum += shared_data[i];
+        }
+        distances[sample_idx] = total_sum;
+    }
+}
+
+// 向量化Warp级优化版本
+__global__ void manhattan_gelu_kernel_vectorized_warp(
+    const float* __restrict__ x,
+    const float* __restrict__ y,
+    float* __restrict__ distances,
+    float* __restrict__ gelu_output,
+    int batch_size,
+    int feature_dim
+) {
+    int sample_idx = blockIdx.x;
+    if (sample_idx >= batch_size) return;
+    
+    int tid = threadIdx.x;
+    int warp_id = tid / 32;
+    int lane_id = tid % 32;
+    
+    int base = sample_idx * feature_dim;
+    
+    // Warp级计算累加和
+    float warp_sum = 0.0f;
+    
+    // 使用float4向量化
+    const float4* x_vec = reinterpret_cast<const float4*>(x + base);
+    const float4* y_vec = reinterpret_cast<const float4*>(y + base);
+    float4* gelu_vec = reinterpret_cast<float4*>(gelu_output + base);
+    
+    int feature_dim_vec = feature_dim / 4;
+    int elements_per_warp_vec = (feature_dim_vec + 8 - 1) / 8;
+    int start_vec = warp_id * elements_per_warp_vec;
+    int end_vec = min(start_vec + elements_per_warp_vec, feature_dim_vec);
+    
+    // GELU常量
+    const float sqrt_2_over_pi = 0.7978845608028654f;
+    const float coeff = 0.044715f;
+    
+    for (int vec_idx = start_vec + lane_id; vec_idx < end_vec; vec_idx += 32) {
+        float4 x_val = x_vec[vec_idx];
+        float4 y_val = y_vec[vec_idx];
+        
+        // 应用GELU激活
+        float4 x_activated;
+        
+        // 对每个组件应用GELU
+        float x_temp = x_val.x;
+        float tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+        x_activated.x = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+        
+        x_temp = x_val.y;
+        tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+        x_activated.y = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+        
+        x_temp = x_val.z;
+        tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+        x_activated.z = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+        
+        x_temp = x_val.w;
+        tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+        x_activated.w = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+        
+        // 存储激活后的输出
+        gelu_vec[vec_idx] = x_activated;
+        
+        // 计算Manhattan距离
+        warp_sum += fabsf(x_activated.x - y_val.x) + fabsf(x_activated.y - y_val.y) + 
+                   fabsf(x_activated.z - y_val.z) + fabsf(x_activated.w - y_val.w);
+    }
+    
+    // 处理剩余元素
+    int remaining_start = feature_dim_vec * 4;
+    for (int dim = remaining_start + warp_id * 32 + lane_id; dim < feature_dim; dim += 256) {
+        float x_val = x[base + dim];
+        float y_val = y[base + dim];
+        
+        // 应用GELU激活
+        float tanh_arg = sqrt_2_over_pi * (x_val + coeff * x_val * x_val * x_val);
+        float x_activated = 0.5f * x_val * (1.0f + tanhf(tanh_arg));
+        gelu_output[base + dim] = x_activated;
+        
+        // 计算Manhattan距离
+        float diff = x_activated - y_val;
+        warp_sum += fabsf(diff);
+    }
+    
+    // Warp级归约求和
+    for (int offset = 16; offset > 0; offset /= 2) {
+        warp_sum += __shfl_down_sync(0xffffffff, warp_sum, offset);
+    }
+    
+    // 使用共享内存进行跨warp归约
+    extern __shared__ float shared_data[];
+    if (lane_id == 0) {
+        shared_data[warp_id] = warp_sum;
+    }
+    __syncthreads();
+    
+    // 第一个线程找到全局总和
+    if (tid == 0) {
+        float total_sum = 0.0f;
+        int num_warps = blockDim.x / 32;
+        for (int i = 0; i < num_warps; i++) {
+            total_sum += shared_data[i];
+        }
+        distances[sample_idx] = total_sum;
+    }
+}
+
+// 主函数 - 纯CUDA实现
+torch::Tensor manhattan_gelu_cuda(
+    torch::Tensor x,
+    torch::Tensor y
+) {
+    // 输入验证
+    TORCH_CHECK(x.scalar_type() == torch::kFloat32, "X must be float32");
+    TORCH_CHECK(y.scalar_type() == torch::kFloat32, "Y must be float32");
+    TORCH_CHECK(x.sizes() == y.sizes(), "X and Y must have same shape");
+    
+    auto x_contig = x.contiguous();
+    auto y_contig = y.contiguous();
+    
+    int batch_size = x_contig.size(0);
+    int feature_dim = x_contig.size(1);
+    
+    // 创建输出张量
+    auto distances = torch::zeros({batch_size}, x.options());
+    auto gelu_output = torch::empty_like(x_contig);
+    
+    const int block_size = 256; // 8个warps
+    size_t shared_mem = 8 * sizeof(float); // 8个warp的结果
+    
+    // 使用向量化Warp级优化版本
+    manhattan_gelu_kernel_vectorized_warp<<<batch_size, block_size, shared_mem>>>(
+        x_contig.data_ptr<float>(),
+        y_contig.data_ptr<float>(),
+        distances.data_ptr<float>(),
+        gelu_output.data_ptr<float>(),
+        batch_size,
+        feature_dim
+    );
+    
+    return distances;
+}
+"""
+
+manhattan_gelu_cpp_source = """
+torch::Tensor manhattan_gelu_cuda(torch::Tensor x, torch::Tensor y);
+"""
+
+# 编译CUDA代码
+manhattan_gelu = load_inline(
+    name="manhattan_gelu",
+    cpp_sources=manhattan_gelu_cpp_source,
+    cuda_sources=manhattan_gelu_source,
+    functions=["manhattan_gelu_cuda"],
+    extra_cuda_cflags=[
+        "-O3", 
+        "--use_fast_math",
+        "-gencode=arch=compute_80,code=sm_80"
+    ],
+    verbose=True
+)
+
+class ModelNew(torch.nn.Module):
+    def __init__(self):
+        super(ModelNew, self).__init__()
+        self.manhattan_gelu = manhattan_gelu
+
+    def forward(self, x, y):
+        return self.manhattan_gelu.manhattan_gelu_cuda(x, y)

S1/wut0n_#63/manhattan_gelu_torchcode.py

@@ -0,0 +1,49 @@
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+
+class Model(nn.Module):
+    """
+    Manhattan + GELU融合实现。
+    先对输入应用GELU激活，然后计算Manhattan距离。
+    """
+    def __init__(self):
+        super(Model, self).__init__()
+    
+    def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+        """
+        Compute Manhattan + GELU fusion.
+
+        Args:
+            x (torch.Tensor): First set of vectors [batch_size, feature_dim]
+            y (torch.Tensor): Second set of vectors [batch_size, feature_dim]
+
+        Returns:
+            torch.Tensor: Manhattan distances after GELU activation [batch_size]
+        """
+        # Input validation
+        if x.shape != y.shape:
+            raise ValueError(f"Input tensors must have the same shape, got {x.shape} and {y.shape}")
+        
+        if x.dim() != 2:
+            raise ValueError(f"Input tensors must be 2D, got {x.dim()}D")
+        
+        # Apply GELU activation to x
+        x_activated = F.gelu(x)
+        
+        # Compute Manhattan distance: Σ|x_activated - y|
+        manhattan_dist = torch.sum(torch.abs(x_activated - y), dim=1)
+        
+        return manhattan_dist
+
+batch_size = 256
+feature_dim = 512
+
+def get_inputs():
+    # Generate two sets of vectors
+    x = torch.randn(batch_size, feature_dim)
+    y = torch.randn(batch_size, feature_dim)
+    return [x, y]
+
+def get_init_inputs():
+    return []  # No special initialization inputs needed

S1/wut0n_#63/prompt.txt

@@ -0,0 +1,245 @@
+You write custom CUDA kernels to replace pytorch operators in given architecture to get speedups. You have complete freedom to choose 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.
+
+**SPECIAL INSTRUCTIONS FOR MANHATTAN + GELU FUSION:**
+
+When implementing Manhattan Distance + GELU fusion, you MUST implement the following optimized strategy:
+
+1. **FUSION ARCHITECTURE**: Combine GELU activation and Manhattan distance computation in a single kernel:
+   - Apply GELU activation to input tensor x first
+   - Compute Manhattan distance between activated x and y
+   - Eliminate intermediate tensor storage for maximum efficiency
+   - Store both activated output and distance results
+
+2. **FLOAT4 VECTORIZATION**: Use float4 vectorization for maximum memory bandwidth utilization:
+   - Process 4 elements simultaneously using float4 loads/stores
+   - Apply GELU activation to all 4 components in parallel
+   - Compute Manhattan distance for all 4 components together
+   - Handle remaining elements with scalar processing
+
+3. **WARP-LEVEL OPTIMIZATION**: Use warp-level processing for maximum performance:
+   - Each block processes one sample from the batch
+   - Use 8 warps per block (256 threads) for optimal GPU utilization
+   - Use __shfl_down_sync for efficient warp-level reduction of sums
+   - Divide feature dimensions among warps for parallel processing
+
+4. **MEMORY COALESCING**: Ensure efficient memory access patterns:
+   - Use float4 vectorized loads for coalesced memory access
+   - Store GELU results using float4 vectorized stores
+   - Each thread processes multiple elements with stride pattern
+   - Minimize global memory accesses through fusion
+
+5. **EFFICIENT SUM REDUCTION**: Implement optimized sum reduction for Manhattan distance:
+cpp
+// Warp-level sum reduction
+for (int offset = 16; offset > 0; offset /= 2) {
+    warp_sum += __shfl_down_sync(0xffffffff, warp_sum, offset);
+}
+
+// Cross-warp reduction using shared memory
+extern __shared__ float shared_data[];
+if (lane_id == 0) {
+    shared_data[warp_id] = warp_sum;
+}
+__syncthreads();
+
+
+
+6. **GELU FUSION**: Integrate GELU activation seamlessly with vectorization:
+cpp
+// GELU constants
+const float sqrt_2_over_pi = 0.7978845608028654f; // sqrt(2/π)
+const float coeff = 0.044715f;
+
+// Apply GELU activation to float4 vector
+float4 x_activated;
+float x_temp = x_val.x;
+float tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+x_activated.x = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+
+x_temp = x_val.y;
+tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+x_activated.y = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+
+x_temp = x_val.z;
+tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+x_activated.z = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+
+x_temp = x_val.w;
+tanh_arg = sqrt_2_over_pi * (x_temp + coeff * x_temp * x_temp * x_temp);
+x_activated.w = 0.5f * x_temp * (1.0f + tanhf(tanh_arg));
+
+// Store activated output
+gelu_vec[vec_idx] = x_activated;
+
+// Compute Manhattan distance for all 4 components
+warp_sum += fabsf(x_activated.x - y_val.x) + fabsf(x_activated.y - y_val.y) + 
+           fabsf(x_activated.z - y_val.z) + fabsf(x_activated.w - y_val.w);
+
+
+
+7. **SHARED MEMORY PATTERN**: Use efficient shared memory organization:
+cpp
+// For sum reduction across warps
+extern __shared__ float shared_data[];
+if (lane_id == 0) {
+    shared_data[warp_id] = warp_sum;
+}
+__syncthreads();
+
+// Final sum calculation
+if (tid == 0) {
+    float total_sum = 0.0f;
+    int num_warps = blockDim.x / 32;
+    for (int i = 0; i < num_warps; i++) {
+        total_sum += shared_data[i];
+    }
+    distances[sample_idx] = total_sum;
+}
+
+
+
+8. **BLOCK CONFIGURATION**: Use optimal settings for vectorized processing:
+   - Block size: 256 threads (8 warps)
+   - Shared memory: 8 * sizeof(float) for warp reduction results
+   - One block per sample for maximum parallelism
+   - Elements per warp: (feature_dim + 8 - 1) / 8
+
+9. **PRECISION REQUIREMENTS**: Ensure exact mathematical alignment:
+   - GELU Activation: gelu(x) = 0.5x * (1 + erf(x/√2))
+   - Approximate GELU: gelu(x) = 0.5x * (1 + tanh(√(2/π) * (x + 0.044715x³)))
+   - Manhattan Distance: Σ|gelu(x) - y|
+   - Use fabsf for absolute value computation
+   - Use tanhf for GELU approximation
+   - Verify with torch.allclose(rtol=1e-03, atol=1e-6)
+
+10. **FUNCTION SIGNATURE**: The main CUDA function must accept all parameters:
+cpp
+torch::Tensor manhattan_gelu_cuda(
+    torch::Tensor x,
+    torch::Tensor y
+)
+
+
+
+11. **MATHEMATICAL FORMULAS**: Implement exact mathematical operations:
+    - GELU Activation: gelu(x) = 0.5x * (1 + erf(x/√2))
+    - Approximate GELU: gelu(x) = 0.5x * (1 + tanh(√(2/π) * (x + 0.044715x³)))
+    - Absolute Difference: abs_diff = |gelu(x) - y|
+    - Manhattan Distance: manhattan_dist = Σabs_diff
+
+12. **PYTHON CALLING CONVENTION**: The ModelNew forward method must pass parameters correctly:
+python
+def forward(self, x, y):
+    return self.manhattan_gelu.manhattan_gelu_cuda(x, y)
+
+
+
+13. **OUTPUT REQUIREMENTS**: Generate both distances and activated outputs:
+    - Primary output: Manhattan distances after GELU activation [batch_size]
+    - Secondary output: GELU activated tensor [batch_size, feature_dim]
+    - Both outputs must match PyTorch reference implementation exactly
+
+14. **PERFORMANCE OPTIMIZATIONS**: Include advanced optimizations:
+    - Use fast math optimizations (--use_fast_math)
+    - Optimize for compute capability 8.0+ (sm_80)
+    - Use -O3 optimization level
+    - Avoid bank conflicts in shared memory access
+    - Use efficient memory access patterns
+
+15. **ALGORITHM CHOICE**: Prioritize the vectorized fused approach:
+    - float4 vectorization is mandatory for this implementation
+    - Do NOT implement scalar-only versions
+    - The fusion must happen at the CUDA kernel level, not Python level
+    - Eliminate all intermediate tensor storage
+
+16. **BOUNDARY HANDLING**: Properly handle non-multiple-of-4 feature dimensions:
+    - Use float4 for vectorized processing of main portion
+    - Handle remaining elements with scalar processing
+    - Ensure no memory access violations
+    - Maintain mathematical correctness for all dimensions
+
+17. **GELU APPROXIMATION**: Use the standard tanh approximation for efficiency:
+    - GELU(x) ≈ 0.5x * (1 + tanh(√(2/π) * (x + 0.044715x³)))
+    - This approximation is widely used in practice and provides good accuracy
+    - Precompute constants: sqrt(2/π) ≈ 0.7978845608028654
+    - Coefficient: 0.044715
+
+Here's the target architecture to optimize:
+
+python
+import torch
+import torch.nn as nn
+
+class Model(nn.Module):
+"""
+Manhattan Distance implementation.
+Computes the Manhattan distance (L1 distance) between two sets of vectors.
+"""
+def init(self):
+super(Model, self).init()
+
+def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
+    """
+    Compute Manhattan distance between x and y.
+
+    Args:
+        x (torch.Tensor): First set of vectors [batch_size, feature_dim]
+        y (torch.Tensor): Second set of vectors [batch_size, feature_dim]
+
+    Returns:
+        torch.Tensor: Manhattan distances [batch_size]
+    """
+    # Input validation
+    if x.shape != y.shape:
+        raise ValueError(f"Input tensors must have the same shape, got {x.shape} and {y.shape}")
+    
+    if x.dim() != 2:
+        raise ValueError(f"Input tensors must be 2D, got {x.dim()}D")
+    
+    # Compute Manhattan distance: Σ|x_i - y_i|
+    manhattan_dist = torch.sum(torch.abs(x - y), dim=1)
+    
+    return manhattan_dist
+
+batch_size = 256
+feature_dim = 512
+
+def get_inputs():
+# Generate two sets of vectors
+x = torch.randn(batch_size, feature_dim)
+y = torch.randn(batch_size, feature_dim)
+return [x, y]
+
+def get_init_inputs():
+return [] # No special initialization inputs needed
+
+
+
+**EXPECTED OUTPUT STRUCTURE**:
+Generate two files:
+1. `manhattan_gelu_cudacode.py` - Contains ModelNew class with Manhattan+GELU fusion using pure CUDA
+2. `manhattan_gelu_torchcode.py` - Contains the reference PyTorch implementation with GELU fusion
+
+**KEY REQUIREMENTS**:
+- The CUDA implementation must use pure CUDA functions only
+- Must implement GELU activation before Manhattan distance computation
+- Must use float4 vectorization for maximum performance
+- Must use warp-level optimization for maximum performance
+- Must use efficient sum reduction algorithm
+- Must handle arbitrary tensor shapes (not just fixed dimensions)
+- Must maintain mathematical precision with PyTorch implementation
+- Must use optimal block configuration (256 threads, 8 warps)
+- Expected speedup: 1.4-2.0x over PyTorch baseline
+- Must use fast math optimizations for better performance
+- Must be robust and handle edge cases properly
+- Must use only pure CUDA functions (no PyTorch internal functions)
+- Must use fabsf for absolute value computation
+- Must use tanhf for GELU approximation
+- Must implement exact mathematical formulas for GELU and Manhattan distance
+- Must generate both distance and activated output tensors
+- Must use shared memory efficiently for warp-level sum reduction
+- Must ensure coalesced memory access patterns
+- Must eliminate intermediate tensor storage for maximum fusion benefits
+- Must implement the complete fusion in a single CUDA kernel
+- Must use float4 vectorization as the primary optimization strategy
+- Must use the standard tanh approximation for GELU: gelu(x) = 0.5x * (1 + tanh(√(2/π) * (x + 0.044715x³)))

S1/wut0n_#63/run_code.py

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