From 133f7269ab39ff0cfb5827549cd01bdb1c498f00 Mon Sep 17 00:00:00 2001
From: Ljy123 <Ljy123@example.org>
Date: Tue, 2 Dec 2025 15:26:08 +0800
Subject: [PATCH] finish swish2bias #25

---
 S1/Ljy123_#25/cudacode.py  | 53 ++++++++++++++++++++++++++++++++++++
 S1/Ljy123_#25/prompt.txt   |  1 +
 S1/Ljy123_#25/run_code.py  | 56 ++++++++++++++++++++++++++++++++++++++
 S1/Ljy123_#25/torchcode.py | 24 ++++++++++++++++
 4 files changed, 134 insertions(+)
 create mode 100644 S1/Ljy123_#25/cudacode.py
 create mode 100644 S1/Ljy123_#25/prompt.txt
 create mode 100644 S1/Ljy123_#25/run_code.py
 create mode 100644 S1/Ljy123_#25/torchcode.py

diff --git a/S1/Ljy123_#25/cudacode.py b/S1/Ljy123_#25/cudacode.py
new file mode 100644
index 00000000..571e87c3
--- /dev/null
+++ b/S1/Ljy123_#25/cudacode.py
@@ -0,0 +1,53 @@
+import torch
+from torch.utils.cpp_extension import load_inline
+
+source = """
+#include <torch/extension.h>
+#include <cuda_runtime.h>
+
+__global__ void swish2_bias_kernel(const float* x, const float* bias, float* y, int dim, long long total) {
+    long long idx = blockIdx.x * blockDim.x + threadIdx.x;
+    long long stride = blockDim.x * gridDim.x;
+    for (long long i = idx; i < total; i += stride) {
+        int j = (int)(i % dim);
+        float z = x[i] + bias[j];
+        float s = 1.0f / (1.0f + expf(-z));
+        float v = z * s;
+        y[i] = v * v;
+    }
+}
+
+torch::Tensor swish2_bias_cuda(torch::Tensor x, torch::Tensor bias) {
+    auto x_contig = x.contiguous();
+    auto b_contig = bias.contiguous();
+    auto y = torch::empty_like(x_contig);
+    long long total = x_contig.numel();
+    int dim = (int)x_contig.size(-1);
+    int block = 512;
+    long long grid = (total + block - 1) / block;
+    grid = grid > 65535 ? 65535 : grid;
+    swish2_bias_kernel<<<(int)grid, block>>>(x_contig.data_ptr<float>(), b_contig.data_ptr<float>(), y.data_ptr<float>(), dim, total);
+    return y;
+}
+"""
+
+cpp_source = """
+torch::Tensor swish2_bias_cuda(torch::Tensor x, torch::Tensor bias);
+"""
+
+ops = load_inline(
+    name="swish2_bias",
+    cpp_sources=cpp_source,
+    cuda_sources=source,
+    functions=["swish2_bias_cuda"],
+    verbose=True
+)
+
+class ModelNew(torch.nn.Module):
+    def __init__(self, bias: torch.Tensor):
+        super(ModelNew, self).__init__()
+        self.ops = ops
+        self.register_buffer("bias", bias)
+
+    def forward(self, x):
+        return self.ops.swish2_bias_cuda(x, self.bias)
diff --git a/S1/Ljy123_#25/prompt.txt b/S1/Ljy123_#25/prompt.txt
new file mode 100644
index 00000000..6414bd5a
--- /dev/null
+++ b/S1/Ljy123_#25/prompt.txt
@@ -0,0 +1 @@
+融合算子：Swish^2+Bias，计算 y = (z*sigmoid(z))^2，z = x + bias，一次内核完成。
diff --git a/S1/Ljy123_#25/run_code.py b/S1/Ljy123_#25/run_code.py
new file mode 100644
index 00000000..8bb6bf24
--- /dev/null
+++ b/S1/Ljy123_#25/run_code.py
@@ -0,0 +1,56 @@
+import torch
+import time
+from torchcode import Model, get_inputs, get_init_inputs
+from cudacode import ModelNew
+
+def run_benchmark():
+    if not torch.cuda.is_available():
+        print("CUDA 不可用，请确保您有可用的 NVIDIA GPU 并已正确安装 PyTorch CUDA 版本。")
+        return
+    device = torch.device("cuda")
+
+    init_inputs = [x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in get_init_inputs()]
+    inputs = [x.cuda(device=device) if isinstance(x, torch.Tensor) else x for x in get_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("❌ 精度不一致！")
+        diff = (output_torch - output_cuda).abs().max().item()
+        print(f"最大绝对误差: {diff}")
+        print(f"输出张量形状: torch={tuple(output_torch.shape)}, cuda={tuple(output_cuda.shape)}")
+        print(f"数据类型: torch={output_torch.dtype}, cuda={output_cuda.dtype}")
+        print(f"设备: torch={output_torch.device}, cuda={output_cuda.device}")
+
+    print("\n-------------------- 性能加速比测试 --------------------")
+    num_iterations = 100
+    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
+
+    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 Swish^2+Bias 平均执行时间: {torch_time:.6f} 秒")
+    print(f"自定义 CUDA 融合内核 平均执行时间: {cuda_time:.6f} 秒")
+    speedup = torch_time / cuda_time if cuda_time > 0 else 0
+    if cuda_time > 0:
+        print(f"加速比 (Speedup): {speedup:.2f}x")
+    else:
+        print("CUDA 内核执行时间为0，无法计算加速比。")
+    return precision_flag, speedup
+
+if __name__ == "__main__":
+    run_benchmark()
diff --git a/S1/Ljy123_#25/torchcode.py b/S1/Ljy123_#25/torchcode.py
new file mode 100644
index 00000000..b805440f
--- /dev/null
+++ b/S1/Ljy123_#25/torchcode.py
@@ -0,0 +1,24 @@
+import torch
+import torch.nn as nn
+import torch.nn.functional as F
+
+class Model(nn.Module):
+    def __init__(self, bias: torch.Tensor):
+        super(Model, self).__init__()
+        self.register_buffer("bias", bias)
+
+    def forward(self, x: torch.Tensor) -> torch.Tensor:
+        z = x + self.bias
+        s = torch.sigmoid(z)
+        return (z * s) * (z * s)
+
+batch_size = 16
+dim = 16384
+
+def get_inputs():
+    x = torch.randn(batch_size, dim)
+    return [x]
+
+def get_init_inputs():
+    bias = torch.randn(dim)
+    return [bias]