From 4b2fd825cf7f5ef13ec09cfc04b835fec3f29c3a Mon Sep 17 00:00:00 2001
From: wut0n <3455534242@qq.com>
Date: Wed, 10 Dec 2025 23:53:23 +0800
Subject: [PATCH] FEAT:add bce_sigmoid #111

---
 S1/wut0n_#111/bce_sigmoid_cudacode.py  | 123 +++++++++++++++++++++++++
 S1/wut0n_#111/bce_sigmoid_torchcode.py |  46 +++++++++
 S1/wut0n_#111/prompt.txt               | 108 ++++++++++++++++++++++
 S1/wut0n_#111/run_code.py              |  84 +++++++++++++++++
 4 files changed, 361 insertions(+)
 create mode 100644 S1/wut0n_#111/bce_sigmoid_cudacode.py
 create mode 100644 S1/wut0n_#111/bce_sigmoid_torchcode.py
 create mode 100644 S1/wut0n_#111/prompt.txt
 create mode 100644 S1/wut0n_#111/run_code.py

diff --git a/S1/wut0n_#111/bce_sigmoid_cudacode.py b/S1/wut0n_#111/bce_sigmoid_cudacode.py
new file mode 100644
index 00000000..89f591a6
--- /dev/null
+++ b/S1/wut0n_#111/bce_sigmoid_cudacode.py
@@ -0,0 +1,123 @@
+import torch
+from torch.utils.cpp_extension import load_inline
+
+bce_sigmoid_source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+// 修复后的融合sigmoid + BCE kernel
+__global__ void bce_sigmoid_fused_kernel(
+    const float* __restrict__ logits,
+    const float* __restrict__ targets,
+    float* __restrict__ loss,
+    int size
+) {
+    int idx = blockIdx.x * blockDim.x + threadIdx.x;
+    
+    // 局部累加器
+    float local_sum = 0.0f;
+    
+    // 每个线程处理多个元素
+    int stride = blockDim.x * gridDim.x;
+    for (int i = idx; i < size; i += stride) {
+        float logit = logits[i];
+        float target = targets[i];
+        
+        // 修复：正确的BCE with logits公式
+        // BCE = log(1 + exp(-logit)) + (1-target)*logit 当 logit >= 0
+        // BCE = -logit + log(1 + exp(logit)) + target*logit 当 logit < 0
+        float bce_loss;
+        if (logit >= 0.0f) {
+            // 对于正logit：使用数值稳定的计算
+            float exp_neg_logit = expf(-logit);
+            bce_loss = log1pf(exp_neg_logit) + (1.0f - target) * logit;
+        } else {
+            // 对于负logit：使用数值稳定的计算
+            float exp_logit = expf(logit);
+            bce_loss = -logit + log1pf(exp_logit) + target * logit;
+        }
+        
+        local_sum += bce_loss;
+    }
+    
+    // 使用共享内存进行块内归约
+    extern __shared__ float shared_mem[];
+    shared_mem[threadIdx.x] = local_sum;
+    
+    __syncthreads();
+    
+    // 块内归约
+    for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
+        if (threadIdx.x < stride) {
+            shared_mem[threadIdx.x] += shared_mem[threadIdx.x + stride];
+        }
+        __syncthreads();
+    }
+    
+    // 第一个线程将结果写入全局内存
+    if (threadIdx.x == 0) {
+        atomicAdd(loss, shared_mem[0]);
+    }
+}
+
+torch::Tensor bce_sigmoid_cuda(torch::Tensor logits, torch::Tensor targets) {
+    TORCH_CHECK(logits.scalar_type() == torch::kFloat32, "Input must be float32");
+    TORCH_CHECK(targets.scalar_type() == torch::kFloat32, "Target must be float32");
+    TORCH_CHECK(logits.sizes() == targets.sizes(), "Input and target must have same shape");
+    
+    auto logits_contig = logits.contiguous();
+    auto targets_contig = targets.contiguous();
+    int size = logits_contig.numel();
+    
+    // 创建输出tensor，初始化为0
+    auto loss = torch::zeros({1}, logits.options());
+    
+    // 优化的kernel配置
+    const int block_size = 256;
+    int num_blocks = min(65535, (size + block_size - 1) / block_size);
+    
+    size_t shared_mem = block_size * sizeof(float);
+    
+    // 启动修复后的kernel
+    bce_sigmoid_fused_kernel<<<num_blocks, block_size, shared_mem>>>(
+        logits_contig.data_ptr<float>(),
+        targets_contig.data_ptr<float>(),
+        loss.data_ptr<float>(),
+        size
+    );
+    
+    // 检查CUDA错误
+    cudaError_t err = cudaGetLastError();
+    if (err != cudaSuccess) {
+        AT_ERROR("CUDA kernel failed: ", cudaGetErrorString(err));
+    }
+    
+    return loss;
+}
+"""
+
+bce_sigmoid_cpp_source = """
+torch::Tensor bce_sigmoid_cuda(torch::Tensor logits, torch::Tensor targets);
+"""
+
+# 编译修复后的CUDA代码
+bce_sigmoid = load_inline(
+    name="bce_sigmoid_fixed",
+    cpp_sources=bce_sigmoid_cpp_source,
+    cuda_sources=bce_sigmoid_source,
+    functions=["bce_sigmoid_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.bce_sigmoid = bce_sigmoid
+
+    def forward(self, logits, targets):
+        return self.bce_sigmoid.bce_sigmoid_cuda(logits, targets)
diff --git a/S1/wut0n_#111/bce_sigmoid_torchcode.py b/S1/wut0n_#111/bce_sigmoid_torchcode.py
new file mode 100644
index 00000000..aa712bc1
--- /dev/null
+++ b/S1/wut0n_#111/bce_sigmoid_torchcode.py
@@ -0,0 +1,46 @@
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+
+class Model(nn.Module):
+    """
+    BCE with Sigmoid implementation for binary classification.
+    使用binary_cross_entropy_with_logits作为基准
+    """
+    def __init__(self, reduction='sum'):
+        super(Model, self).__init__()
+        self.reduction = reduction
+    
+    def forward(self, inputs: torch.Tensor, targets: torch.Tensor) -> torch.Tensor:
+        """
+        Compute BCE Loss with logits.
+        
+        Args:
+            inputs (torch.Tensor): Predicted logits
+            targets (torch.Tensor): Ground truth labels (0 or 1)
+        
+        Returns:
+            torch.Tensor: Computed BCE loss
+        """
+        # 确保类型一致
+        inputs = inputs.to(torch.float32)
+        targets = targets.to(torch.float32)
+        
+        # 使用PyTorch的标准实现
+        return F.binary_cross_entropy_with_logits(
+            inputs, 
+            targets, 
+            reduction=self.reduction
+        )
+
+batch_size = 256
+num_features = 2000
+
+def get_inputs():
+    # 生成logits（不是概率）
+    input_logits = torch.randn(batch_size, num_features, dtype=torch.float32)
+    target_labels = torch.randint(0, 2, (batch_size, num_features), dtype=torch.float32)
+    return [input_logits, target_labels]
+
+def get_init_inputs():
+    return []
diff --git a/S1/wut0n_#111/prompt.txt b/S1/wut0n_#111/prompt.txt
new file mode 100644
index 00000000..006b2cf1
--- /dev/null
+++ b/S1/wut0n_#111/prompt.txt
@@ -0,0 +1,108 @@
+Write a custom CUDA kernel to replace PyTorch's Focal Loss with Label Smoothing implementation for binary classification.
+
+You are given the following PyTorch architecture:
+
+python
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+
+class Model(nn.Module):
+"""
+Focal Loss with Label Smoothing implementation for binary classification.
+Combines label smoothing with focal loss for better generalization.
+Focal Loss = -α * (1-pt)^γ * log(pt_smoothed)
+where pt = p if target=1, else (1-p), p = sigmoid(logit)
+"""
+def init(self, alpha=0.25, gamma=2.0, smoothing=0.1, reduction='mean'):
+super(Model, self).init()
+self.alpha = alpha
+self.gamma = gamma
+self.smoothing = smoothing
+self.reduction = reduction
+
+def forward(self, inputs: torch.Tensor, targets: torch.Tensor) -> torch.Tensor:
+    """
+    Compute Focal Loss with Label Smoothing.
+    
+    Args:
+        inputs (torch.Tensor): Predicted logits of shape (batch_size, num_classes)
+        targets (torch.Tensor): Ground truth labels of shape (batch_size,)
+    
+    Returns:
+        torch.Tensor: Computed focal loss with label smoothing
+    """
+    # Ensure input types are consistent
+    inputs = inputs.to(torch.float32)
+    targets = targets.to(torch.float32)
+    
+    # Handle shape matching
+    if inputs.dim() == 2 and inputs.size(1) == 1:
+        inputs = inputs.squeeze(1)
+    
+    # Apply label smoothing to targets
+    # For binary classification: smooth_target = (1-smoothing)*target + smoothing/2
+    smoothed_targets = (1.0 - self.smoothing) * targets + self.smoothing / 2.0
+    
+    # Compute probabilities with sigmoid
+    probs = torch.sigmoid(inputs)
+    
+    # Compute pt based on original targets (not smoothed)
+    pt = torch.where(targets == 1, probs, 1 - probs)
+    
+    # Compute focal weight
+    focal_weight = self.alpha * torch.pow(1 - pt, self.gamma)
+    
+    # Compute binary cross entropy with smoothed targets
+    bce = F.binary_cross_entropy_with_logits(inputs, smoothed_targets, reduction='none')
+    
+    # Apply focal weight
+    focal_loss = focal_weight * bce
+    
+    # Apply reduction
+    if self.reduction == 'mean':
+        return focal_loss.mean()
+    elif self.reduction == 'sum':
+        return focal_loss.sum()
+    else:
+        return focal_loss
+batch_size = 32
+num_classes = 1
+
+def get_inputs():
+# Generate random logits with explicit float32
+inputs = torch.randn(batch_size, num_classes, dtype=torch.float32)
+# Generate random binary targets (0 or 1) with explicit float32
+targets = torch.randint(0, 2, (batch_size,), dtype=torch.float32)
+return [inputs, targets]
+
+def get_init_inputs():
+return [0.25, 2.0, 0.1] # alpha, gamma, smoothing
+
+
+
+Your task is to optimize this Focal Loss with Label Smoothing implementation by:
+
+1. **Complete Operator Fusion**: Combine the label smoothing, sigmoid computation, and focal loss calculation into a single CUDA kernel to eliminate intermediate tensor storage and multiple computation passes.
+
+2. **Enhanced Numerical Stability**: Implement numerically stable sigmoid computation with conditional branches for positive/negative logits, use optimized BCE computation with log1p for better precision, and add proper epsilon handling (1e-8).
+
+3. **Label Smoothing Integration**: Directly compute smoothed targets within the kernel using the formula: smoothed_target = (1-smoothing)*target + smoothing/2, avoiding separate tensor operations.
+
+4. **Memory Access Optimization**: Minimize global memory access by keeping all intermediate computations (smoothed_target, sigmoid, pt, focal_weight, bce) in registers, and ensure coalesced memory access patterns.
+
+5. **Optimized BCE Computation**: Implement numerically stable binary cross entropy computation using logits directly with log1p function, avoiding intermediate probability calculations for better precision.
+
+The optimized CUDA kernel should:
+- Take logits and targets as input (both float32)
+- Compute smoothed targets internally: smoothed_target = (1-smoothing)*target + smoothing/2
+- Compute sigmoid, pt, focal weight, and BCE loss in a single fused kernel
+- Use optimized sigmoid computation with numerical stability for both positive and negative logits
+- Implement stable BCE computation using logits and log1p function for enhanced precision
+- Compute pt based on original targets (not smoothed) for focal weight calculation
+- Output the fused focal loss values with label smoothing
+- Support both 'mean' and 'sum' reduction modes
+- Use optimized compilation flags (-O3, --use_fast_math)
+- Achieve significant speedup (1.5-2.0x) over the PyTorch implementation through complete fusion and reduced memory overhead
+
+Follow the inline CUDA extension syntax example provided in the reference. The kernel should demonstrate performance improvements through complete operator fusion, enhanced numerical stability with log1p, and optimized memory access patterns.
diff --git a/S1/wut0n_#111/run_code.py b/S1/wut0n_#111/run_code.py
new file mode 100644
index 00000000..80a5974d
--- /dev/null
+++ b/S1/wut0n_#111/run_code.py
@@ -0,0 +1,84 @@
+###########################################################
+# 性能和精度验证程序
+###########################################################
+import torch
+import torch.nn as nn
+import time
+from bce_sigmoid_torchcode import Model, get_inputs, get_init_inputs
+from bce_sigmoid_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, atol=1e-05)
+    max_diff = torch.max(torch.abs(output_torch - output_cuda)).item()
+    mean_diff = torch.mean(torch.abs(output_torch - output_cuda)).item()
+    
+    if precision_flag:
+        print(f"✅ 精度对齐：两个模型的输出结果非常接近。")
+        print(f"最大误差: {max_diff:.8f}, 平均误差: {mean_diff:.8f}")
+    else:
+        print(f"❌ 精度不一致！最大误差: {max_diff:.8f}, 平均误差: {mean_diff:.8f}")
+        
+    print("\n-------------------- 性能加速比测试 --------------------")
+    num_iterations = 100
+    
+    # GPU 预热
+    for _ in range(10):
+        _ = torch_model(*inputs)
+        _ = cuda_model(*inputs)
+    
+    # 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 bce_sigmoid 平均执行时间: {torch_time:.6f} 秒")
+    print(f"自定义 CUDA bce_sigmoid 平均执行时间: {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()