From e8d0056a8ac7b49f1ddbcd969a67986db0eb533b Mon Sep 17 00:00:00 2001 From: Ljy123 Date: Wed, 10 Dec 2025 23:35:01 +0800 Subject: [PATCH] finish SigmoidSlopeGate #110 --- S1/Ljy123_#110/cudacode.py | 132 ++++++++++++++++++++++++++++++++++++ S1/Ljy123_#110/prompt.txt | 10 +++ S1/Ljy123_#110/run_code.py | 50 ++++++++++++++ S1/Ljy123_#110/torchcode.py | 29 ++++++++ 4 files changed, 221 insertions(+) create mode 100644 S1/Ljy123_#110/cudacode.py create mode 100644 S1/Ljy123_#110/prompt.txt create mode 100644 S1/Ljy123_#110/run_code.py create mode 100644 S1/Ljy123_#110/torchcode.py diff --git a/S1/Ljy123_#110/cudacode.py b/S1/Ljy123_#110/cudacode.py new file mode 100644 index 00000000..6fd66df3 --- /dev/null +++ b/S1/Ljy123_#110/cudacode.py @@ -0,0 +1,132 @@ +import torch +from torch.utils.cpp_extension import load_inline + +source = """ +#include +#include +#include + +__device__ __forceinline__ float sigmoidf(float x){ + return __fdividef(1.0f, 1.0f + __expf(-x)); +} + +__global__ __launch_bounds__(256) void sigmoid_slope_affine_gate_kernel(const float* __restrict__ x, const float* __restrict__ scale, const float* __restrict__ bias, float* __restrict__ y, int B, int D, float alpha, float beta){ + int row = blockIdx.x; + int tid = threadIdx.x; + int col_block = blockIdx.y; + int stride4 = blockDim.x * 4; + int col_idx = (col_block * stride4) + tid * 4; + const float* xr = x + row * D; + float* yr = y + row * D; + for(int base = col_idx; base < D; base += stride4 * 2){ + if (base < D){ + float4 xv = reinterpret_cast(xr + base)[0]; + float4 sv = reinterpret_cast(scale + base)[0]; + float4 bv = reinterpret_cast(bias + base)[0]; + float4 yv; + float z0 = fmaf(xv.x, sv.x, bv.x); + float z1 = fmaf(xv.y, sv.y, bv.y); + float z2 = fmaf(xv.z, sv.z, bv.z); + float z3 = fmaf(xv.w, sv.w, bv.w); + float s0 = sigmoidf(z0); + float s1 = sigmoidf(z1); + float s2 = sigmoidf(z2); + float s3 = sigmoidf(z3); + float m0 = s0 * (1.0f - s0); + float m1 = s1 * (1.0f - s1); + float m2 = s2 * (1.0f - s2); + float m3 = s3 * (1.0f - s3); + float t0 = fmaf(alpha, m0, beta); + float t1 = fmaf(alpha, m1, beta); + float t2 = fmaf(alpha, m2, beta); + float t3 = fmaf(alpha, m3, beta); + float g0 = sigmoidf(t0); + float g1 = sigmoidf(t1); + float g2 = sigmoidf(t2); + float g3 = sigmoidf(t3); + yv.x = xv.x * g0; + yv.y = xv.y * g1; + yv.z = xv.z * g2; + yv.w = xv.w * g3; + reinterpret_cast(yr + base)[0] = yv; + } + int base2 = base + stride4; + if (base2 < D){ + float4 xv2 = reinterpret_cast(xr + base2)[0]; + float4 sv2 = reinterpret_cast(scale + base2)[0]; + float4 bv2 = reinterpret_cast(bias + base2)[0]; + float4 yv2; + float z0b = fmaf(xv2.x, sv2.x, bv2.x); + float z1b = fmaf(xv2.y, sv2.y, bv2.y); + float z2b = fmaf(xv2.z, sv2.z, bv2.z); + float z3b = fmaf(xv2.w, sv2.w, bv2.w); + float sb0 = sigmoidf(z0b); + float sb1 = sigmoidf(z1b); + float sb2 = sigmoidf(z2b); + float sb3 = sigmoidf(z3b); + float mb0 = sb0 * (1.0f - sb0); + float mb1 = sb1 * (1.0f - sb1); + float mb2 = sb2 * (1.0f - sb2); + float mb3 = sb3 * (1.0f - sb3); + float tb0 = fmaf(alpha, mb0, beta); + float tb1 = fmaf(alpha, mb1, beta); + float tb2 = fmaf(alpha, mb2, beta); + float tb3 = fmaf(alpha, mb3, beta); + float gb0 = sigmoidf(tb0); + float gb1 = sigmoidf(tb1); + float gb2 = sigmoidf(tb2); + float gb3 = sigmoidf(tb3); + yv2.x = xv2.x * gb0; + yv2.y = xv2.y * gb1; + yv2.z = xv2.z * gb2; + yv2.w = xv2.w * gb3; + reinterpret_cast(yr + base2)[0] = yv2; + } + } +} + +torch::Tensor sigmoid_slope_affine_gate_cuda(torch::Tensor x, torch::Tensor scale, torch::Tensor bias, double alpha, double beta){ + auto xc = x.contiguous(); + auto sc = scale.contiguous(); + auto bc = bias.contiguous(); + auto y = torch::empty_like(xc); + int B = (int)xc.size(0); + int D = (int)xc.size(1); + float a = (float)alpha; + float be = (float)beta; + int block = 256; + int elements_per_thread = 4; + int elements_per_block = block * elements_per_thread; + int gy = (D + elements_per_block - 1) / elements_per_block; + dim3 grid(B, gy); + sigmoid_slope_affine_gate_kernel<<>>(xc.data_ptr(), sc.data_ptr(), bc.data_ptr(), y.data_ptr(), B, D, a, be); + return y; +} +""" + +cpp_source = """ +#include +torch::Tensor sigmoid_slope_affine_gate_cuda(torch::Tensor x, torch::Tensor scale, torch::Tensor bias, double alpha, double beta); +""" + +ops = load_inline( + name="sigmoid_slope_affine_gate", + cpp_sources=cpp_source, + cuda_sources=source, + functions=["sigmoid_slope_affine_gate_cuda"], + extra_cflags=["-O3","-std=c++17"], + extra_cuda_cflags=["-O3","--use_fast_math","-std=c++17","-Xptxas","-O3,-dlcm=ca","-maxrregcount=64"], + verbose=True +) + +class ModelNew(torch.nn.Module): + def __init__(self, scale: torch.Tensor, bias: torch.Tensor, alpha: float, beta: float): + super(ModelNew, self).__init__() + self.ops = ops + self.register_buffer("scale", scale) + self.register_buffer("bias", bias) + self.alpha = float(alpha) + self.beta = float(beta) + + def forward(self, x): + return self.ops.sigmoid_slope_affine_gate_cuda(x, self.scale, self.bias, self.alpha, self.beta) diff --git a/S1/Ljy123_#110/prompt.txt b/S1/Ljy123_#110/prompt.txt new file mode 100644 index 00000000..7be45012 --- /dev/null +++ b/S1/Ljy123_#110/prompt.txt @@ -0,0 +1,10 @@ +Operator: SigmoidSlope-Affine-Gate (Fused CUDA Kernel) + +Definition +- z = x * scale + bias +- m = sigmoid(z) * (1 - sigmoid(z)) +- g = sigmoid(alpha * m + beta) +- y = x * g + +Goal +- Fuse derivative-like feature; target ≥1.30x speedup. diff --git a/S1/Ljy123_#110/run_code.py b/S1/Ljy123_#110/run_code.py new file mode 100644 index 00000000..abf1461b --- /dev/null +++ b/S1/Ljy123_#110/run_code.py @@ -0,0 +1,50 @@ +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().eval() + cuda_model = ModelNew(*init_inputs).cuda().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 + 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 SigmoidSlope-Affine-Gate 平均执行时间: {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_#110/torchcode.py b/S1/Ljy123_#110/torchcode.py new file mode 100644 index 00000000..1a6e6b76 --- /dev/null +++ b/S1/Ljy123_#110/torchcode.py @@ -0,0 +1,29 @@ +import torch +import torch.nn as nn + +class Model(nn.Module): + def __init__(self, scale: torch.Tensor, bias: torch.Tensor, alpha: float, beta: float): + super(Model, self).__init__() + self.register_buffer("scale", scale) + self.register_buffer("bias", bias) + self.register_buffer("alpha", torch.tensor(float(alpha), dtype=torch.float32)) + self.register_buffer("beta", torch.tensor(float(beta), dtype=torch.float32)) + + def forward(self, x: torch.Tensor) -> torch.Tensor: + z = x * self.scale + self.bias + s = torch.sigmoid(z) + m = s * (1.0 - s) + g = torch.sigmoid(self.alpha * m + self.beta) + return x * g + +batch_size = 16 +dim = 16384 + +def get_inputs(): + x = torch.randn(batch_size, dim) + return [x] + +def get_init_inputs(): + scale = torch.randn(dim) + bias = torch.randn(dim) + return [scale, bias, 1.0, 0.0]