diff --git a/S1/Ljy123_#21/cudacode.py b/S1/Ljy123_#21/cudacode.py new file mode 100644 index 00000000..3b7d42ac --- /dev/null +++ b/S1/Ljy123_#21/cudacode.py @@ -0,0 +1,66 @@ +import torch +from torch.utils.cpp_extension import load_inline + +source = """ +#include +#include + +__global__ void power_sigmoid_gate_kernel(const float* x, float* y, long long size, float alpha, float beta, float p) { + long long idx = blockIdx.x * blockDim.x + threadIdx.x; + long long stride = blockDim.x * gridDim.x; + for (long long i = idx; i < size; i += stride) { + float u = x[i]; + float r = fmaxf(u, 0.0f); + float v = __fmaf_rn(alpha, u, beta); + float s = 1.0f / (1.0f + expf(-v)); + float pw; + if (fabsf(p - 2.0f) < 1e-7f) { + pw = r * r; + } else if (fabsf(p - 3.0f) < 1e-7f) { + pw = r * r * r; + } else if (fabsf(p - 1.0f) < 1e-7f) { + pw = r; + } else { + pw = powf(r, p); + } + y[i] = pw * s; + } +} + +torch::Tensor power_sigmoid_gate_cuda(torch::Tensor x, torch::Tensor alpha, torch::Tensor beta, torch::Tensor p) { + auto x_contig = x.contiguous(); + auto y = torch::empty_like(x_contig); + float a = alpha.item().toFloat(); + float b = beta.item().toFloat(); + float pe = p.item().toFloat(); + long long total = x_contig.numel(); + int block = 512; + long long grid = (total + block - 1) / block; + grid = grid > 65535 ? 65535 : grid; + power_sigmoid_gate_kernel<<<(int)grid, block>>>(x_contig.data_ptr(), y.data_ptr(), total, a, b, pe); + return y; +} +""" + +cpp_source = """ +torch::Tensor power_sigmoid_gate_cuda(torch::Tensor x, torch::Tensor alpha, torch::Tensor beta, torch::Tensor p); +""" + +ops = load_inline( + name="power_sigmoid_gate", + cpp_sources=cpp_source, + cuda_sources=source, + functions=["power_sigmoid_gate_cuda"], + verbose=True +) + +class ModelNew(torch.nn.Module): + def __init__(self, alpha: torch.Tensor, beta: torch.Tensor, p: torch.Tensor): + super(ModelNew, self).__init__() + self.ops = ops + self.register_buffer("alpha", alpha) + self.register_buffer("beta", beta) + self.register_buffer("p", p) + + def forward(self, x): + return self.ops.power_sigmoid_gate_cuda(x, self.alpha, self.beta, self.p) diff --git a/S1/Ljy123_#21/prompt.txt b/S1/Ljy123_#21/prompt.txt new file mode 100644 index 00000000..f1e07297 --- /dev/null +++ b/S1/Ljy123_#21/prompt.txt @@ -0,0 +1,5 @@ +独特融合算子：Power-Sigmoid-Gate。一次内核完成 ReLU 的幂、以及对输入的仿射 Sigmoid 门控，适合需要控制非线性强度与门控强度的场景。 + +torchcode.py：参考实现 `y = relu(x)^p * sigmoid(alpha*x + beta)`。 +cudacode.py：`__global__ void power_sigmoid_gate_kernel(...)` 完成融合计算。 +run_code.py：比较精度与性能（100 次迭代，`rtol=1e-03`）。 diff --git a/S1/Ljy123_#21/run_code.py b/S1/Ljy123_#21/run_code.py new file mode 100644 index 00000000..ee4a0e1e --- /dev/null +++ b/S1/Ljy123_#21/run_code.py @@ -0,0 +1,76 @@ +########################################################### +# 性能和精度验证程序 +########################################################### +import torch +import torch.nn as nn +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 + 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("❌ 精度不一致！") + 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 Power-Sigmoid-Gate 平均执行时间: {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() diff --git a/S1/Ljy123_#21/torchcode.py b/S1/Ljy123_#21/torchcode.py new file mode 100644 index 00000000..046b0ba5 --- /dev/null +++ b/S1/Ljy123_#21/torchcode.py @@ -0,0 +1,28 @@ +import torch +import torch.nn as nn +import torch.nn.functional as F + +class Model(nn.Module): + def __init__(self, alpha: torch.Tensor, beta: torch.Tensor, p: torch.Tensor): + super(Model, self).__init__() + self.register_buffer("alpha", alpha) + self.register_buffer("beta", beta) + self.register_buffer("p", p) + + def forward(self, x: torch.Tensor) -> torch.Tensor: + g = torch.sigmoid(self.alpha * x + self.beta) + r = torch.relu(x) + return torch.pow(r, float(self.p.item())) * g + +batch_size = 16 +dim = 16384 + +def get_inputs(): + x = torch.randn(batch_size, dim) + return [x] + +def get_init_inputs(): + alpha = torch.tensor(1.0, dtype=torch.float32) + beta = torch.tensor(0.0, dtype=torch.float32) + p = torch.tensor(2.0, dtype=torch.float32) + return [alpha, beta, p]