forked from ccf-ai-infra/GPUCodeForces
feat:add high performance Logbeta #29
This commit is contained in:
parent
f989885dde
commit
a136ac20d5
|
|
@ -0,0 +1,109 @@
|
|||
import torch
|
||||
from torch.utils.cpp_extension import load_inline
|
||||
|
||||
logbeta_source = """
|
||||
#include <torch/extension.h>
|
||||
#include <cuda_runtime.h>
|
||||
#include <cmath>
|
||||
|
||||
// VECTORIZED版本：4元素向量化处理（最终优化版本）
|
||||
__global__ void logbeta_vectorized_kernel(
|
||||
const float* __restrict__ x,
|
||||
const float* __restrict__ y,
|
||||
float* __restrict__ z,
|
||||
int size
|
||||
) {
|
||||
int tid = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int stride = blockDim.x * gridDim.x;
|
||||
|
||||
// 每次处理4个元素
|
||||
int vec_size = size / 4;
|
||||
|
||||
for (int i = tid; i < vec_size; i += stride) {
|
||||
int base_idx = i * 4;
|
||||
|
||||
// 加载4个元素对
|
||||
float x0 = x[base_idx + 0];
|
||||
float x1 = x[base_idx + 1];
|
||||
float x2 = x[base_idx + 2];
|
||||
float x3 = x[base_idx + 3];
|
||||
|
||||
float y0 = y[base_idx + 0];
|
||||
float y1 = y[base_idx + 1];
|
||||
float y2 = y[base_idx + 2];
|
||||
float y3 = y[base_idx + 3];
|
||||
|
||||
// 计算4个log beta值（使用标准库函数确保精度）
|
||||
z[base_idx + 0] = lgammaf(x0) + lgammaf(y0) - lgammaf(x0 + y0);
|
||||
z[base_idx + 1] = lgammaf(x1) + lgammaf(y1) - lgammaf(x1 + y1);
|
||||
z[base_idx + 2] = lgammaf(x2) + lgammaf(y2) - lgammaf(x2 + y2);
|
||||
z[base_idx + 3] = lgammaf(x3) + lgammaf(y3) - lgammaf(x3 + y3);
|
||||
}
|
||||
|
||||
// 处理剩余元素
|
||||
int remainder = size % 4;
|
||||
if (tid == 0 && remainder > 0) {
|
||||
int start_idx = vec_size * 4;
|
||||
for (int i = start_idx; i < size; i++) {
|
||||
float xi = x[i];
|
||||
float yi = y[i];
|
||||
z[i] = lgammaf(xi) + lgammaf(yi) - lgammaf(xi + yi);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
torch::Tensor logbeta_cuda(
|
||||
torch::Tensor x,
|
||||
torch::Tensor y,
|
||||
std::string mode = "vectorized"
|
||||
) {
|
||||
auto x_contig = x.contiguous();
|
||||
auto y_contig = y.contiguous();
|
||||
auto z = torch::empty_like(x_contig);
|
||||
int size = x_contig.numel();
|
||||
|
||||
if (mode == "vectorized") {
|
||||
// 向量化版本
|
||||
const int block_size = 256;
|
||||
int num_blocks = (size / 4 + block_size - 1) / block_size;
|
||||
int grid_size = std::min(num_blocks, 65535);
|
||||
|
||||
logbeta_vectorized_kernel<<<grid_size, block_size>>>(
|
||||
x_contig.data_ptr<float>(),
|
||||
y_contig.data_ptr<float>(),
|
||||
z.data_ptr<float>(),
|
||||
size
|
||||
);
|
||||
}
|
||||
|
||||
return z;
|
||||
}
|
||||
"""
|
||||
|
||||
logbeta_cpp_source = """
|
||||
torch::Tensor logbeta_cuda(torch::Tensor x, torch::Tensor y, std::string mode);
|
||||
"""
|
||||
|
||||
# 编译CUDA代码
|
||||
logbeta = load_inline(
|
||||
name="logbeta_vectorized",
|
||||
cpp_sources=logbeta_cpp_source,
|
||||
cuda_sources=logbeta_source,
|
||||
functions=["logbeta_cuda"],
|
||||
extra_cuda_cflags=[
|
||||
"-O3",
|
||||
"--use_fast_math",
|
||||
"-std=c++17",
|
||||
"-maxrregcount=64"
|
||||
],
|
||||
verbose=True
|
||||
)
|
||||
|
||||
class ModelNew(torch.nn.Module):
|
||||
def __init__(self, mode="vectorized"):
|
||||
super(ModelNew, self).__init__()
|
||||
self.mode = mode
|
||||
self.logbeta = logbeta # The module containing the kernel
|
||||
|
||||
def forward(self, x, y):
|
||||
return self.logbeta.logbeta_cuda(x, y, self.mode)
|
||||
|
|
@ -0,0 +1,37 @@
|
|||
import torch
|
||||
import torch.nn as nn
|
||||
import math
|
||||
|
||||
class Model(nn.Module):
|
||||
"""
|
||||
Log Beta operator implementation.
|
||||
"""
|
||||
def __init__(self):
|
||||
super(Model, self).__init__()
|
||||
|
||||
def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
|
||||
"""
|
||||
Computes the logarithm of the beta function: log(beta(x, y))
|
||||
|
||||
The beta function is defined as: beta(x, y) = gamma(x) * gamma(y) / gamma(x + y)
|
||||
So log(beta(x, y)) = lgamma(x) + lgamma(y) - lgamma(x + y)
|
||||
|
||||
Args:
|
||||
x (torch.Tensor): First input tensor of any shape.
|
||||
y (torch.Tensor): Second input tensor of same shape as x.
|
||||
|
||||
Returns:
|
||||
torch.Tensor: Output tensor with log beta applied, same shape as input.
|
||||
"""
|
||||
return torch.lgamma(x) + torch.lgamma(y) - torch.lgamma(x + y)
|
||||
|
||||
batch_size = 512
|
||||
dim = 16384
|
||||
|
||||
def get_inputs():
|
||||
x = torch.rand(batch_size, dim).cuda() * 10.0 + 0.1 # 避免零值
|
||||
y = torch.rand(batch_size, dim).cuda() * 10.0 + 0.1
|
||||
return [x, y]
|
||||
|
||||
def get_init_inputs():
|
||||
return []
|
||||
|
|
@ -0,0 +1,98 @@
|
|||
You write custom CUDA kernels to replace the pytorch operators in the given architecture to get speedups.
|
||||
|
||||
You have complete freedom to choose the set of operators you want to replace. You may make the decision to replace some operators with custom CUDA kernels and leave others unchanged. You may replace multiple operators with custom implementations, consider operator fusion opportunities (combining multiple operators into a single kernel, for example, combining matmul+relu), or algorithmic changes (such as online softmax). You are only limited by your imagination.
|
||||
|
||||
Here's an example to show you the syntax of inline embedding custom CUDA operators in torch: The example given architecture is:
|
||||
|
||||
python
|
||||
import torch
|
||||
import torch.nn as nn
|
||||
import torch.nn.functional as F
|
||||
|
||||
class Model(nn.Module):
|
||||
def init(self) -> None:
|
||||
super().init()
|
||||
|
||||
def forward(self, a, b):
|
||||
return a + b
|
||||
def get_inputs():
|
||||
# randomly generate input tensors based on the model architecture
|
||||
a = torch.randn(1, 128).cuda()
|
||||
b = torch.randn(1, 128).cuda()
|
||||
return [a, b]
|
||||
|
||||
def get_init_inputs():
|
||||
# randomly generate tensors required for initialization based on the model architecture
|
||||
return []
|
||||
|
||||
|
||||
|
||||
The example new arch with custom CUDA kernels looks like this:
|
||||
python
|
||||
import torch
|
||||
import torch.nn as nn
|
||||
import torch.nn.functional as F
|
||||
|
||||
class Model(nn.Module):
|
||||
def init(self) -> None:
|
||||
super().init()
|
||||
|
||||
def forward(self, a, b):
|
||||
return a + b
|
||||
def get_inputs():
|
||||
# randomly generate input tensors based on the model architecture
|
||||
a = torch.randn(1, 128).cuda()
|
||||
b = torch.randn(1, 128).cuda()
|
||||
return [a, b]
|
||||
|
||||
def get_init_inputs():
|
||||
# randomly generate tensors required for initialization based on the model architecture
|
||||
return []
|
||||
|
||||
|
||||
|
||||
You are given the following architecture:
|
||||
|
||||
python
|
||||
import torch
|
||||
import torch.nn as nn
|
||||
|
||||
class Model(nn.Module):
|
||||
“”"
|
||||
Log Beta operator implementation.
|
||||
“”"
|
||||
def init(self):
|
||||
super(Model, self).init()
|
||||
|
||||
def forward(self, x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
|
||||
"""
|
||||
Computes the logarithm of the beta function: log(beta(x, y))
|
||||
|
||||
The beta function is defined as: beta(x, y) = gamma(x) * gamma(y) / gamma(x + y)
|
||||
So log(beta(x, y)) = lgamma(x) + lgamma(y) - lgamma(x + y)
|
||||
|
||||
Args:
|
||||
x (torch.Tensor): First input tensor of any shape.
|
||||
y (torch.Tensor): Second input tensor of same shape as x.
|
||||
|
||||
Returns:
|
||||
torch.Tensor: Output tensor with log beta applied, same shape as input.
|
||||
"""
|
||||
return torch.lgamma(x) + torch.lgamma(y) - torch.lgamma(x + y)
|
||||
batch_size = 16
|
||||
dim = 16384
|
||||
|
||||
def get_inputs():
|
||||
x = torch.rand(batch_size, dim).cuda() * 10.0 + 0.1 # avoid zero values
|
||||
y = torch.rand(batch_size, dim).cuda() * 10.0 + 0.1
|
||||
return [x, y]
|
||||
|
||||
def get_init_inputs():
|
||||
return [] # No special initialization inputs needed
|
||||
Optimization Requirements:
|
||||
1. Implement vectorized CUDA kernel for log beta computation
|
||||
2. Use 4-element vectorization for better memory bandwidth utilization
|
||||
3. Ensure numerical precision by using CUDA's lgammaf function
|
||||
4. Optimize for both performance and accuracy
|
||||
5. Handle edge cases and boundary conditions properly
|
||||
6. Provide multiple optimization modes (precise, vectorized, high-performance)
|
||||
|
|
@ -0,0 +1,74 @@
|
|||
###########################################################
|
||||
# 性能和精度验证程序
|
||||
###########################################################
|
||||
import torch
|
||||
import torch.nn as nn
|
||||
import time
|
||||
from logbeta_torchcode import Model, get_inputs, get_init_inputs
|
||||
from logbeta_cudacode import ModelNew
|
||||
|
||||
def run_benchmark():
|
||||
# 检查 CUDA 是否可用
|
||||
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("❌ 精度不一致！")
|
||||
|
||||
print("\n-------------------- 性能加速比测试 --------------------")
|
||||
num_iterations = 100
|
||||
|
||||
# PyTorch 模型计时
|
||||
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
|
||||
|
||||
# 自定义 CUDA 内核计时
|
||||
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 logbeta 平均执行时间: {torch_time:.6f} 秒")
|
||||
print(f"自定义 CUDA logbeta 平均执行时间: {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()
|
||||
Loading…
Reference in New Issue