forked from ccf-ai-infra/GPUCodeForces
optimize power sigmoid (affine-gate) #46
This commit is contained in:
parent
10eed82956
commit
6ea60a5ddf
|
|
@ -0,0 +1,105 @@
|
|||
import torch
|
||||
from torch.utils.cpp_extension import load_inline
|
||||
|
||||
source = """
|
||||
#include <torch/extension.h>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
__global__ void power_sigmoid_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, float power){
|
||||
int b = blockIdx.x;
|
||||
int lane = blockIdx.y * blockDim.x + threadIdx.x;
|
||||
int stride = blockDim.x * gridDim.y;
|
||||
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 = lane * 4; i < D4; i += stride * 4){
|
||||
float4 xv = reinterpret_cast<const float4*>(xr)[i / 4];
|
||||
float4 sv = reinterpret_cast<const float4*>(scale)[i / 4];
|
||||
float4 bv = reinterpret_cast<const float4*>(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);
|
||||
float a0 = fabsf(z0), a1 = fabsf(z1), a2 = fabsf(z2), a3 = fabsf(z3);
|
||||
float v0 = exp2f(power * log2f(a0));
|
||||
float v1 = exp2f(power * log2f(a1));
|
||||
float v2 = exp2f(power * log2f(a2));
|
||||
float v3 = exp2f(power * log2f(a3));
|
||||
float g0 = 1.0f / (1.0f + expf(-(alpha * v0 + beta)));
|
||||
float g1 = 1.0f / (1.0f + expf(-(alpha * v1 + beta)));
|
||||
float g2 = 1.0f / (1.0f + expf(-(alpha * v2 + beta)));
|
||||
float g3 = 1.0f / (1.0f + expf(-(alpha * v3 + beta)));
|
||||
yv.x = xv.x * g0;
|
||||
yv.y = xv.y * g1;
|
||||
yv.z = xv.z * g2;
|
||||
yv.w = xv.w * g3;
|
||||
reinterpret_cast<float4*>(yr)[i / 4] = yv;
|
||||
}
|
||||
#pragma unroll 4
|
||||
for(int i = D4 + lane; i < D; i += stride){
|
||||
float z = fmaf(xr[i], scale[i], bias[i]);
|
||||
float a = fabsf(z);
|
||||
float v = exp2f(power * log2f(a));
|
||||
float g = 1.0f / (1.0f + expf(-(alpha * v + beta)));
|
||||
yr[i] = xr[i] * g;
|
||||
}
|
||||
} else {
|
||||
#pragma unroll 4
|
||||
for(int i = lane; i < D; i += stride){
|
||||
float z = fmaf(xr[i], scale[i], bias[i]);
|
||||
float a = fabsf(z);
|
||||
float v = exp2f(power * log2f(a));
|
||||
float g = 1.0f / (1.0f + expf(-(alpha * v + beta)));
|
||||
yr[i] = xr[i] * g;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
torch::Tensor power_sigmoid_affine_gate_cuda(torch::Tensor x, torch::Tensor scale, torch::Tensor bias, torch::Tensor alpha, torch::Tensor beta, torch::Tensor power){
|
||||
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>();
|
||||
float be = beta.item<float>();
|
||||
float pw = power.item<float>();
|
||||
int block = 256;
|
||||
int gy = max(1, min((D + 4095) / 4096, 8));
|
||||
dim3 grid(B, gy);
|
||||
power_sigmoid_affine_gate_kernel<<<grid, block>>>(xc.data_ptr<float>(), sc.data_ptr<float>(), bc.data_ptr<float>(), y.data_ptr<float>(), B, D, a, be, pw);
|
||||
return y;
|
||||
}
|
||||
"""
|
||||
|
||||
cpp_source = """
|
||||
torch::Tensor power_sigmoid_affine_gate_cuda(torch::Tensor x, torch::Tensor scale, torch::Tensor bias, torch::Tensor alpha, torch::Tensor beta, torch::Tensor power);
|
||||
"""
|
||||
|
||||
ops = load_inline(
|
||||
name="power_sigmoid_affine_gate",
|
||||
cpp_sources=cpp_source,
|
||||
cuda_sources=source,
|
||||
functions=["power_sigmoid_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, power: 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))
|
||||
self.register_buffer("power", torch.tensor(float(power), dtype=torch.float32))
|
||||
|
||||
def forward(self, x):
|
||||
return self.ops.power_sigmoid_affine_gate_cuda(x, self.scale, self.bias, self.alpha, self.beta, self.power)
|
||||
|
|
@ -0,0 +1,28 @@
|
|||
融合算子：Power-Sigmoid-Affine-Gate（一次核内完成仿射、幂变换与 Sigmoid 门控，返回 y = x * σ(α * |z|^p + β)，其中 z = x*scale + bias，p>0）。通过控制 p 可调节非线性强度，适合幅值敏感的门控。
|
||||
|
||||
目标与定义
|
||||
- 输入张量：`x[B, D]`
|
||||
- 逐维参数：`scale[D]`、`bias[D]`
|
||||
- 标量超参：`alpha`、`beta`、`power=p`
|
||||
- 计算流程：`z = x*scale + bias`，`v = |z|^p`（建议 `exp2(power * log2(|z|))` 实现），`g = sigmoid(alpha*v + beta)`，`y = x * g`
|
||||
|
||||
参考实现（文件要求）
|
||||
- `torchcode.py`：参考 `Model` 含 `power`；`get_inputs()`/`get_init_inputs()` 统一接口
|
||||
- `cudacode.py`：单核融合（仿射+幂+sigmoid+乘法）；建议 `exp2f(power*log2f(|z|))` 以获得 `--use_fast_math` 下更高吞吐；`-O3 --use_fast_math`
|
||||
- `run_code.py`：100 次迭代；精度 `rtol=1e-03, atol=1e-06`；输出加速比
|
||||
|
||||
CUDA 实现要点
|
||||
- 行并行：`grid = B`，块内沿 D 连续访存；只遍历一次并写 `y`
|
||||
- 对齐向量化：满足 16 字节对齐且 `D%4==0` 时走 `float4`；否则标量路径；两路径计算保持一致
|
||||
- 指令优化：仿射用 `fmaf`；幂变换走 `log2f/exp2f`；sigmoid 用 `expf`；循环 `#pragma unroll 4`
|
||||
- 零值处理：当 `|z|=0` 时，`log2f(0)=-inf`，`exp2f(-inf)=0`，最终 `v=0`，该分支数值稳定
|
||||
- 线程配置：推荐 `block=1024` 起步，按设备试探最佳值
|
||||
|
||||
评估与目标
|
||||
- 精度：与 PyTorch 参考对齐（`rtol=1e-03, atol=1e-06`）
|
||||
- 性能：融合与向量化后期望 ≥1.0x；大尺度更优
|
||||
|
||||
加分项（可选）
|
||||
- 针对未对齐路径减少尾循环与分支开销
|
||||
- 在 p 为常量时进行常量传播或近似优化（谨慎控制误差）
|
||||
|
||||
|
|
@ -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 Power-Sigmoid-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()
|
||||
|
||||
|
|
@ -0,0 +1,31 @@
|
|||
import torch
|
||||
import torch.nn as nn
|
||||
|
||||
class Model(nn.Module):
|
||||
def __init__(self, scale: torch.Tensor, bias: torch.Tensor, alpha: float, beta: float, power: 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))
|
||||
self.register_buffer("power", torch.tensor(float(power), dtype=torch.float32))
|
||||
|
||||
def forward(self, x: torch.Tensor) -> torch.Tensor:
|
||||
z = x * self.scale + self.bias
|
||||
a = torch.abs(z)
|
||||
v = torch.pow(a, self.power)
|
||||
g = torch.sigmoid(self.alpha * v + 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, 1.5]
|
||||
|
||||
Loading…
Reference in New Issue