Feat:add softmax+mask #85

S1/wut0n_#85/prompt.txt

@@ -0,0 +1,307 @@
+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 softmax+mask), or algorithmic changes . 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
+
+class Model(nn.Module):
+
+    def __init__(self):
+        super(Model, self).__init__()
+    
+    def forward(self, scores: torch.Tensor, mask: torch.Tensor) -> torch.Tensor:
+        if mask.dtype != torch.float32:
+            mask = mask.float()
+
+        mask_scores = scores.masked_fill(mask == 0, float('-inf'))
+        output = torch.softmax(mask_scores, dim=-1)
+        return output
+
+N = 1024      # batch size
+SEQ_LEN = 512 # sequence length
+
+def get_inputs():
+    scores = torch.randn(N, SEQ_LEN)
+    mask = torch.randint(0, 2, (N, SEQ_LEN)).float() 
+    return [scores, mask]
+
+def get_init_inputs():
+    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
+from torch.utils.cpp_extension import load_inline
+
+softmax_mask_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+#include <math.h>
+
+__device__ __forceinline__ float warp_reduce_max(float val) {
+    #pragma unroll
+    for (int mask = 16; mask > 0; mask >>= 1) {
+        val = fmaxf(val, __shfl_xor_sync(0xffffffff, val, mask));
+    }
+    return val;
+}
+
+__device__ __forceinline__ float warp_reduce_sum(float val) {
+    #pragma unroll
+    for (int mask = 16; mask > 0; mask >>= 1) {
+        val += __shfl_xor_sync(0xffffffff, val, mask);
+    }
+    return val;
+}
+
+
+__global__ void softmax_mask_kernel(
+    const float* __restrict__ scores,
+    const float* __restrict__ mask,
+    float* __restrict__ output,
+    int N, int seq_len
+) {
+
+    int row = blockIdx.x * blockDim.y + threadIdx.y;
+    int lane_id = threadIdx.x;
+    
+    if (row >= N) return;
+    
+    const float* scores_row = scores + row * seq_len;
+    const float* mask_row = mask + row * seq_len;
+    float* output_row = output + row * seq_len;
+    float max_val = -FLT_MAX;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        max_val = fmaxf(max_val, masked_score);
+    }
+    
+    // Warp reduction for max
+    max_val = warp_reduce_max(max_val);
+    
+    float sum = 0.0f;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        float exp_val = expf(masked_score - max_val);
+        sum += exp_val;
+
+        output_row[col] = exp_val;
+    }
+    
+    sum = warp_reduce_sum(sum);
+    float inv_sum = 1.0f / sum;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        output_row[col] *= inv_sum;
+    }
+}
+
+// 向量化优化的Softmax + Mask融合kernel
+__global__ void softmax_mask_kernel_vectorized(
+    const float* __restrict__ scores,
+    const float* __restrict__ mask,
+    float* __restrict__ output,
+    int N, int seq_len
+) {
+    int row = blockIdx.x * blockDim.y + threadIdx.y;
+    int lane_id = threadIdx.x;
+    
+    if (row >= N) return;
+    
+    const float* scores_row = scores + row * seq_len;
+    const float* mask_row = mask + row * seq_len;
+    float* output_row = output + row * seq_len;
+
+    float max_val = -FLT_MAX;
+    int vec_elems = seq_len / 4;
+    
+    // 向量化处理
+    for (int i = lane_id; i < vec_elems; i += 32) {
+        int col_start = i * 4;
+        
+        // 加载4个元素
+        float score0 = scores_row[col_start + 0];
+        float score1 = scores_row[col_start + 1];
+        float score2 = scores_row[col_start + 2];
+        float score3 = scores_row[col_start + 3];
+        
+        float mask0 = mask_row[col_start + 0];
+        float mask1 = mask_row[col_start + 1];
+        float mask2 = mask_row[col_start + 2];
+        float mask3 = mask_row[col_start + 3];
+        
+        float masked_score0 = (mask0 > 0.5f) ? score0 : -FLT_MAX;
+        float masked_score1 = (mask1 > 0.5f) ? score1 : -FLT_MAX;
+        float masked_score2 = (mask2 > 0.5f) ? score2 : -FLT_MAX;
+        float masked_score3 = (mask3 > 0.5f) ? score3 : -FLT_MAX;
+        
+        max_val = fmaxf(max_val, fmaxf(fmaxf(masked_score0, masked_score1), fmaxf(masked_score2, masked_score3)));
+    }
+    
+    for (int col = vec_elems * 4 + lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        max_val = fmaxf(max_val, masked_score);
+    }
+    
+    max_val = warp_reduce_max(max_val);
+
+    float sum = 0.0f;
+
+    for (int i = lane_id; i < vec_elems; i += 32) {
+        int col_start = i * 4;
+        
+        float score0 = scores_row[col_start + 0];
+        float score1 = scores_row[col_start + 1];
+        float score2 = scores_row[col_start + 2];
+        float score3 = scores_row[col_start + 3];
+        
+        float mask0 = mask_row[col_start + 0];
+        float mask1 = mask_row[col_start + 1];
+        float mask2 = mask_row[col_start + 2];
+        float mask3 = mask_row[col_start + 3];
+        
+        float masked_score0 = (mask0 > 0.5f) ? score0 : -FLT_MAX;
+        float masked_score1 = (mask1 > 0.5f) ? score1 : -FLT_MAX;
+        float masked_score2 = (mask2 > 0.5f) ? score2 : -FLT_MAX;
+        float masked_score3 = (mask3 > 0.5f) ? score3 : -FLT_MAX;
+        
+        float exp0 = expf(masked_score0 - max_val);
+        float exp1 = expf(masked_score1 - max_val);
+        float exp2 = expf(masked_score2 - max_val);
+        float exp3 = expf(masked_score3 - max_val);
+        
+        sum += exp0 + exp1 + exp2 + exp3;
+        
+        output_row[col_start + 0] = exp0;
+        output_row[col_start + 1] = exp1;
+        output_row[col_start + 2] = exp2;
+        output_row[col_start + 3] = exp3;
+    }
+    
+    for (int col = vec_elems * 4 + lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        float exp_val = expf(masked_score - max_val);
+        sum += exp_val;
+        output_row[col] = exp_val;
+    }
+    
+    sum = warp_reduce_sum(sum);
+    float inv_sum = 1.0f / sum;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        output_row[col] *= inv_sum;
+    }
+}
+
+torch::Tensor softmax_mask_cuda(
+    torch::Tensor scores,
+    torch::Tensor mask
+) {
+    TORCH_CHECK(scores.is_cuda() && scores.is_contiguous());
+    TORCH_CHECK(mask.is_cuda() && mask.is_contiguous());
+    TORCH_CHECK(scores.dim() == 2 && mask.dim() == 2);
+    TORCH_CHECK(scores.sizes() == mask.sizes());
+
+    int N = scores.size(0);
+    int seq_len = scores.size(1);
+    
+    auto output = torch::empty_like(scores);
+
+    if (seq_len % 4 == 0) {
+        // 向量化版本
+        dim3 block_size(32, 8);  
+        int grid_size = (N + 7) / 8;  
+        
+        softmax_mask_kernel_vectorized<<<grid_size, block_size>>>(
+            scores.data_ptr<float>(),
+            mask.data_ptr<float>(),
+            output.data_ptr<float>(),
+            N, seq_len
+        );
+    } else {
+        dim3 block_size(32, 8);
+        int grid_size = (N + 7) / 8;
+        
+        softmax_mask_kernel<<<grid_size, block_size>>>(
+            scores.data_ptr<float>(),
+            mask.data_ptr<float>(),
+            output.data_ptr<float>(),
+            N, seq_len
+        );
+    }
+    
+    return output;
+}
+"""
+
+softmax_mask_cpp_source = """
+torch::Tensor softmax_mask_cuda(
+    torch::Tensor scores,
+    torch::Tensor mask
+);
+"""
+
+softmax_mask_module = load_inline(
+    name="softmax_mask",
+    cpp_sources=softmax_mask_cpp_source,
+    cuda_sources=softmax_mask_source,
+    functions=["softmax_mask_cuda"],
+    verbose=True,
+    extra_cuda_cflags=["-O3", "--use_fast_math", "-maxrregcount=64"]
+)
+
+class ModelNew(torch.nn.Module):
+    def __init__(self):
+        super(ModelNew, self).__init__()
+        self.softmax_mask = softmax_mask_module
+
+    def forward(self, scores, mask):
+        if not scores.is_contiguous(): scores = scores.contiguous()
+        if not mask.is_contiguous(): mask = mask.contiguous()
+        
+        return self.softmax_mask.softmax_mask_cuda(scores, mask)
+
+```

S1/wut0n_#85/run_code.py

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

@@ -0,0 +1,240 @@
+import torch
+from torch.utils.cpp_extension import load_inline
+
+softmax_mask_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+#include <math.h>
+
+__device__ __forceinline__ float warp_reduce_max(float val) {
+    #pragma unroll
+    for (int mask = 16; mask > 0; mask >>= 1) {
+        val = fmaxf(val, __shfl_xor_sync(0xffffffff, val, mask));
+    }
+    return val;
+}
+
+__device__ __forceinline__ float warp_reduce_sum(float val) {
+    #pragma unroll
+    for (int mask = 16; mask > 0; mask >>= 1) {
+        val += __shfl_xor_sync(0xffffffff, val, mask);
+    }
+    return val;
+}
+
+
+__global__ void softmax_mask_kernel(
+    const float* __restrict__ scores,
+    const float* __restrict__ mask,
+    float* __restrict__ output,
+    int N, int seq_len
+) {
+
+    int row = blockIdx.x * blockDim.y + threadIdx.y;
+    int lane_id = threadIdx.x;
+    
+    if (row >= N) return;
+    
+    const float* scores_row = scores + row * seq_len;
+    const float* mask_row = mask + row * seq_len;
+    float* output_row = output + row * seq_len;
+    float max_val = -FLT_MAX;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        max_val = fmaxf(max_val, masked_score);
+    }
+    
+    // Warp reduction for max
+    max_val = warp_reduce_max(max_val);
+    
+    float sum = 0.0f;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        float exp_val = expf(masked_score - max_val);
+        sum += exp_val;
+
+        output_row[col] = exp_val;
+    }
+    
+    sum = warp_reduce_sum(sum);
+    float inv_sum = 1.0f / sum;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        output_row[col] *= inv_sum;
+    }
+}
+
+// 向量化优化的Softmax + Mask融合kernel
+__global__ void softmax_mask_kernel_vectorized(
+    const float* __restrict__ scores,
+    const float* __restrict__ mask,
+    float* __restrict__ output,
+    int N, int seq_len
+) {
+    int row = blockIdx.x * blockDim.y + threadIdx.y;
+    int lane_id = threadIdx.x;
+    
+    if (row >= N) return;
+    
+    const float* scores_row = scores + row * seq_len;
+    const float* mask_row = mask + row * seq_len;
+    float* output_row = output + row * seq_len;
+
+    float max_val = -FLT_MAX;
+    int vec_elems = seq_len / 4;
+    
+    // 向量化处理
+    for (int i = lane_id; i < vec_elems; i += 32) {
+        int col_start = i * 4;
+        
+        // 加载4个元素
+        float score0 = scores_row[col_start + 0];
+        float score1 = scores_row[col_start + 1];
+        float score2 = scores_row[col_start + 2];
+        float score3 = scores_row[col_start + 3];
+        
+        float mask0 = mask_row[col_start + 0];
+        float mask1 = mask_row[col_start + 1];
+        float mask2 = mask_row[col_start + 2];
+        float mask3 = mask_row[col_start + 3];
+        
+        float masked_score0 = (mask0 > 0.5f) ? score0 : -FLT_MAX;
+        float masked_score1 = (mask1 > 0.5f) ? score1 : -FLT_MAX;
+        float masked_score2 = (mask2 > 0.5f) ? score2 : -FLT_MAX;
+        float masked_score3 = (mask3 > 0.5f) ? score3 : -FLT_MAX;
+        
+        max_val = fmaxf(max_val, fmaxf(fmaxf(masked_score0, masked_score1), fmaxf(masked_score2, masked_score3)));
+    }
+    
+    for (int col = vec_elems * 4 + lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        max_val = fmaxf(max_val, masked_score);
+    }
+    
+    max_val = warp_reduce_max(max_val);
+
+    float sum = 0.0f;
+
+    for (int i = lane_id; i < vec_elems; i += 32) {
+        int col_start = i * 4;
+        
+        float score0 = scores_row[col_start + 0];
+        float score1 = scores_row[col_start + 1];
+        float score2 = scores_row[col_start + 2];
+        float score3 = scores_row[col_start + 3];
+        
+        float mask0 = mask_row[col_start + 0];
+        float mask1 = mask_row[col_start + 1];
+        float mask2 = mask_row[col_start + 2];
+        float mask3 = mask_row[col_start + 3];
+        
+        float masked_score0 = (mask0 > 0.5f) ? score0 : -FLT_MAX;
+        float masked_score1 = (mask1 > 0.5f) ? score1 : -FLT_MAX;
+        float masked_score2 = (mask2 > 0.5f) ? score2 : -FLT_MAX;
+        float masked_score3 = (mask3 > 0.5f) ? score3 : -FLT_MAX;
+        
+        float exp0 = expf(masked_score0 - max_val);
+        float exp1 = expf(masked_score1 - max_val);
+        float exp2 = expf(masked_score2 - max_val);
+        float exp3 = expf(masked_score3 - max_val);
+        
+        sum += exp0 + exp1 + exp2 + exp3;
+        
+        output_row[col_start + 0] = exp0;
+        output_row[col_start + 1] = exp1;
+        output_row[col_start + 2] = exp2;
+        output_row[col_start + 3] = exp3;
+    }
+    
+    for (int col = vec_elems * 4 + lane_id; col < seq_len; col += 32) {
+        float score = scores_row[col];
+        float mask_val = mask_row[col];
+        float masked_score = (mask_val > 0.5f) ? score : -FLT_MAX;
+        float exp_val = expf(masked_score - max_val);
+        sum += exp_val;
+        output_row[col] = exp_val;
+    }
+    
+    sum = warp_reduce_sum(sum);
+    float inv_sum = 1.0f / sum;
+    
+    for (int col = lane_id; col < seq_len; col += 32) {
+        output_row[col] *= inv_sum;
+    }
+}
+
+torch::Tensor softmax_mask_cuda(
+    torch::Tensor scores,
+    torch::Tensor mask
+) {
+    TORCH_CHECK(scores.is_cuda() && scores.is_contiguous());
+    TORCH_CHECK(mask.is_cuda() && mask.is_contiguous());
+    TORCH_CHECK(scores.dim() == 2 && mask.dim() == 2);
+    TORCH_CHECK(scores.sizes() == mask.sizes());
+
+    int N = scores.size(0);
+    int seq_len = scores.size(1);
+    
+    auto output = torch::empty_like(scores);
+
+    if (seq_len % 4 == 0) {
+        // 向量化版本
+        dim3 block_size(32, 8);  
+        int grid_size = (N + 7) / 8;  
+        
+        softmax_mask_kernel_vectorized<<<grid_size, block_size>>>(
+            scores.data_ptr<float>(),
+            mask.data_ptr<float>(),
+            output.data_ptr<float>(),
+            N, seq_len
+        );
+    } else {
+        dim3 block_size(32, 8);
+        int grid_size = (N + 7) / 8;
+        
+        softmax_mask_kernel<<<grid_size, block_size>>>(
+            scores.data_ptr<float>(),
+            mask.data_ptr<float>(),
+            output.data_ptr<float>(),
+            N, seq_len
+        );
+    }
+    
+    return output;
+}
+"""
+
+softmax_mask_cpp_source = """
+torch::Tensor softmax_mask_cuda(
+    torch::Tensor scores,
+    torch::Tensor mask
+);
+"""
+
+softmax_mask_module = load_inline(
+    name="softmax_mask",
+    cpp_sources=softmax_mask_cpp_source,
+    cuda_sources=softmax_mask_source,
+    functions=["softmax_mask_cuda"],
+    verbose=True,
+    extra_cuda_cflags=["-O3", "--use_fast_math", "-maxrregcount=64"]
+)
+
+class ModelNew(torch.nn.Module):
+    def __init__(self):
+        super(ModelNew, self).__init__()
+        self.softmax_mask = softmax_mask_module
+
+    def forward(self, scores, mask):
+        if not scores.is_contiguous(): scores = scores.contiguous()
+        if not mask.is_contiguous(): mask = mask.contiguous()
+        
+        return self.softmax_mask.softmax_mask_cuda(scores, mask)

S1/wut0n_#85/softmax_mask_torchcode.py

@@ -0,0 +1,26 @@
+import torch
+import torch.nn as nn
+
+class Model(nn.Module):
+
+    def __init__(self):
+        super(Model, self).__init__()
+    
+    def forward(self, scores: torch.Tensor, mask: torch.Tensor) -> torch.Tensor:
+        if mask.dtype != torch.float32:
+            mask = mask.float()
+
+        mask_scores = scores.masked_fill(mask == 0, float('-inf'))
+        output = torch.softmax(mask_scores, dim=-1)
+        return output
+
+N = 1024      # batch size
+SEQ_LEN = 512 # sequence length
+
+def get_inputs():
+    scores = torch.randn(N, SEQ_LEN)
+    mask = torch.randint(0, 2, (N, SEQ_LEN)).float() 
+    return [scores, mask]
+
+def get_init_inputs():
+    return []