diff --git a/S1/Ljy123_#44/cudacode.py b/S1/Ljy123_#44/cudacode.py new file mode 100644 index 00000000..8c225adc --- /dev/null +++ b/S1/Ljy123_#44/cudacode.py @@ -0,0 +1,88 @@ +import torch +from torch.utils.cpp_extension import load_inline + +source = """ +#include +#include + +__global__ void sin_affine_gate_kernel(const float* x, const float* scale, const float* bias, float* y, int B, int D, float alpha, float beta){ + int b = blockIdx.x; + int tid = threadIdx.x; + int stride = blockDim.x; + int row_start = b * D; + const float* xr = x + row_start; + float* yr = y + row_start; + int aligned = ((((long long)xr & 15LL) == 0) && (((long long)yr & 15LL) == 0) && (((long long)scale & 15LL) == 0) && (((long long)bias & 15LL) == 0) && ((D & 3) == 0)); + if(aligned){ + int D4 = (D / 4) * 4; + #pragma unroll 4 + for(int i = tid * 4; i < D4; i += stride * 4){ + float4 xv = reinterpret_cast(xr)[i / 4]; + float4 sv = reinterpret_cast(scale)[i / 4]; + float4 bv = reinterpret_cast(bias)[i / 4]; + 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); + yv.x = xv.x * sinf(alpha * z0 + beta); + yv.y = xv.y * sinf(alpha * z1 + beta); + yv.z = xv.z * sinf(alpha * z2 + beta); + yv.w = xv.w * sinf(alpha * z3 + beta); + reinterpret_cast(yr)[i / 4] = yv; + } + #pragma unroll 4 + for(int i = D4 + tid; i < D; i += stride){ + float z = fmaf(xr[i], scale[i], bias[i]); + yr[i] = xr[i] * sinf(alpha * z + beta); + } + } else { + #pragma unroll 4 + for(int i = tid; i < D; i += stride){ + float z = fmaf(xr[i], scale[i], bias[i]); + yr[i] = xr[i] * sinf(alpha * z + beta); + } + } +} + +torch::Tensor sin_affine_gate_cuda(torch::Tensor x, torch::Tensor scale, torch::Tensor bias, torch::Tensor alpha, torch::Tensor 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 = alpha.item(); + float be = beta.item(); + int block = 1024; + int grid = B; + sin_affine_gate_kernel<<>>(xc.data_ptr(), sc.data_ptr(), bc.data_ptr(), y.data_ptr(), B, D, a, be); + return y; +} +""" + +cpp_source = """ +torch::Tensor sin_affine_gate_cuda(torch::Tensor x, torch::Tensor scale, torch::Tensor bias, torch::Tensor alpha, torch::Tensor beta); +""" + +ops = load_inline( + name="sin_affine_gate", + cpp_sources=cpp_source, + cuda_sources=source, + functions=["sin_affine_gate_cuda"], + extra_cuda_cflags=["-O3","--use_fast_math"], + 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.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): + return self.ops.sin_affine_gate_cuda(x, self.scale, self.bias, self.alpha, self.beta) + diff --git a/S1/Ljy123_#44/prompt.txt b/S1/Ljy123_#44/prompt.txt new file mode 100644 index 00000000..e3e2d146 --- /dev/null +++ b/S1/Ljy123_#44/prompt.txt @@ -0,0 +1,28 @@ +融合算子:Sin-Affine-Gate(一次核内完成仿射与正弦门控,直接返回 y = x * sin(α*z + β),其中 z = x*scale + bias)。通过融合计算减少显存往返与内核启动,适合周期性信号或相位敏感的逐维门控。 + +目标与定义 +- 输入张量:`x[B, D]` +- 逐维参数:`scale[D]`、`bias[D]` +- 标量超参:`alpha`、`beta` +- 计算流程:`z = x*scale + bias`,`g = sin(alpha*z + beta)`,`y = x * g` + +参考实现(文件要求) +- `torchcode.py`:提供 PyTorch 参考 `Model`,实现上述公式;`get_inputs()` 与 `get_init_inputs()` 保持统一接口 +- `cudacode.py`:使用 `load_inline` 实现单核融合(仿射+sin 门控+乘法),编译参数 `-O3 --use_fast_math` +- `run_code.py`:100 次迭代统计耗时;精度判定:`rtol=1e-03, atol=1e-06`;打印加速比 + +CUDA 实现要点 +- 并行策略:`grid = B`(每行一个 block),块内线程沿 D 连续访问,保证合并读写 +- 对齐向量化:`float4` 路径需 16 字节对齐且 `D%4==0`,否则标量路径;两路径逻辑一致 +- 指令优化:仿射使用 `fmaf`;正弦函数采用 `--use_fast_math`;循环 `#pragma unroll 4` +- 访存与写回:一次遍历计算并写回 `y`,避免中间张量与多核启动 +- 线程配置:以 `block=1024` 为基准,可按设备与 D 调整到 `256/512/1024` + +评估与目标 +- 精度对齐:`torch.allclose(..., rtol=1e-03, atol=1e-06)` 通过 +- 性能:基准规模下期望 ≥1.0x;更大规模或更好对齐通常进一步增益 + +加分项(可选) +- 统一处理对齐与非对齐路径的尾部元素,减少分支与尾循环开销 +- 在 D 很大时,每线程批量处理多个元素以提升指令与带宽利用 + diff --git a/S1/Ljy123_#44/run_code.py b/S1/Ljy123_#44/run_code.py new file mode 100644 index 00000000..143c4319 --- /dev/null +++ b/S1/Ljy123_#44/run_code.py @@ -0,0 +1,57 @@ +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, atol=1e-06) + 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 Sin-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_#44/torchcode.py b/S1/Ljy123_#44/torchcode.py new file mode 100644 index 00000000..187a71eb --- /dev/null +++ b/S1/Ljy123_#44/torchcode.py @@ -0,0 +1,28 @@ +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 + g = torch.sin(self.alpha * z + 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] +