forked from ccf-ai-infra/GPUCodeForces
Merge pull request 'feat:add high performance Bceloss #8"' (#197) from wut0n/GPUCodeForces:Bceloss into main
This commit is contained in:
commit
dbabaacc7a
|
|
@ -0,0 +1,187 @@
|
|||
# 优化后的 bce_cudacode.py
|
||||
import torch
|
||||
from torch.utils.cpp_extension import load_inline
|
||||
|
||||
bce_source = """
|
||||
#include <torch/extension.h>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
// 高效归约版本 - 使用两级归约
|
||||
__global__ void bce_kernel_efficient(
|
||||
const float* __restrict__ input,
|
||||
const float* __restrict__ target,
|
||||
float* __restrict__ partial_sums,
|
||||
int size,
|
||||
float epsilon
|
||||
) {
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// 局部累加器
|
||||
float local_bce = 0.0f;
|
||||
|
||||
// 每个线程处理多个元素
|
||||
int stride = blockDim.x * gridDim.x;
|
||||
for (int i = idx; i < size; i += stride) {
|
||||
float xi = input[i];
|
||||
float yi = target[i];
|
||||
|
||||
// 数值稳定性处理
|
||||
xi = max(min(xi, 1.0f - epsilon), epsilon);
|
||||
|
||||
// BCE = -[y * log(x) + (1-y) * log(1-x)]
|
||||
float term1 = yi * logf(xi);
|
||||
float term2 = (1.0f - yi) * logf(1.0f - xi);
|
||||
local_bce -= (term1 + term2);
|
||||
}
|
||||
|
||||
// 使用共享内存进行块内归约
|
||||
extern __shared__ float shared_mem[];
|
||||
shared_mem[threadIdx.x] = local_bce;
|
||||
|
||||
__syncthreads();
|
||||
|
||||
// 块内归约
|
||||
for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
|
||||
if (threadIdx.x < stride) {
|
||||
shared_mem[threadIdx.x] += shared_mem[threadIdx.x + stride];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
// 每个block写入部分和到全局内存
|
||||
if (threadIdx.x == 0) {
|
||||
partial_sums[blockIdx.x] = shared_mem[0];
|
||||
}
|
||||
}
|
||||
|
||||
// 快速版本 - 使用快速log近似
|
||||
__global__ void bce_kernel_fast(
|
||||
const float* __restrict__ input,
|
||||
const float* __restrict__ target,
|
||||
float* __restrict__ bce_loss,
|
||||
int size,
|
||||
float epsilon
|
||||
) {
|
||||
int idx = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
|
||||
// 局部累加器
|
||||
float local_bce = 0.0f;
|
||||
|
||||
// 每个线程处理多个元素
|
||||
int stride = blockDim.x * gridDim.x;
|
||||
for (int i = idx; i < size; i += stride) {
|
||||
float xi = input[i];
|
||||
float yi = target[i];
|
||||
|
||||
// 数值稳定性处理
|
||||
xi = max(min(xi, 1.0f - epsilon), epsilon);
|
||||
|
||||
// 使用快速log近似
|
||||
float log_x = __logf(xi); // 使用CUDA内置快速log
|
||||
float log_1_x = __logf(1.0f - xi);
|
||||
|
||||
local_bce -= (yi * log_x + (1.0f - yi) * log_1_x);
|
||||
}
|
||||
|
||||
// 使用共享内存进行块内归约
|
||||
extern __shared__ float shared_mem[];
|
||||
shared_mem[threadIdx.x] = local_bce;
|
||||
|
||||
__syncthreads();
|
||||
|
||||
// 块内归约
|
||||
for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
|
||||
if (threadIdx.x < stride) {
|
||||
shared_mem[threadIdx.x] += shared_mem[threadIdx.x + stride];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
|
||||
// 只有第一个线程做原子操作
|
||||
if (threadIdx.x == 0) {
|
||||
atomicAdd(bce_loss, shared_mem[0]);
|
||||
}
|
||||
}
|
||||
|
||||
torch::Tensor bce_cuda(torch::Tensor input, torch::Tensor target, std::string mode = "fast") {
|
||||
TORCH_CHECK(input.scalar_type() == torch::kFloat32, "Input must be float32");
|
||||
TORCH_CHECK(target.scalar_type() == torch::kFloat32, "Target must be float32");
|
||||
TORCH_CHECK(input.sizes() == target.sizes(), "Input and target must have same shape");
|
||||
|
||||
auto input_contig = input.contiguous();
|
||||
auto target_contig = target.contiguous();
|
||||
int size = input_contig.numel();
|
||||
|
||||
if (mode == "efficient") {
|
||||
// 高效归约版本 - 两级归约
|
||||
const int block_size = 256;
|
||||
int num_blocks = min(1024, (size + block_size - 1) / block_size);
|
||||
|
||||
// 创建部分和数组
|
||||
auto partial_sums = torch::zeros({num_blocks}, input.options());
|
||||
|
||||
size_t shared_mem = block_size * sizeof(float);
|
||||
bce_kernel_efficient<<<num_blocks, block_size, shared_mem>>>(
|
||||
input_contig.data_ptr<float>(),
|
||||
target_contig.data_ptr<float>(),
|
||||
partial_sums.data_ptr<float>(),
|
||||
size,
|
||||
1e-8f
|
||||
);
|
||||
|
||||
// 在CPU上完成最终归约
|
||||
auto partial_sums_cpu = partial_sums.to(torch::kCPU);
|
||||
float total_loss = 0.0f;
|
||||
for (int i = 0; i < num_blocks; i++) {
|
||||
total_loss += partial_sums_cpu.data_ptr<float>()[i];
|
||||
}
|
||||
|
||||
return torch::tensor({total_loss}, input.options());
|
||||
|
||||
} else {
|
||||
// 快速版本 - 默认
|
||||
auto bce_loss = torch::zeros(1, input.options());
|
||||
|
||||
const int block_size = 256;
|
||||
int num_blocks = min(1024, (size + block_size - 1) / block_size);
|
||||
size_t shared_mem = block_size * sizeof(float);
|
||||
|
||||
bce_kernel_fast<<<num_blocks, block_size, shared_mem>>>(
|
||||
input_contig.data_ptr<float>(),
|
||||
target_contig.data_ptr<float>(),
|
||||
bce_loss.data_ptr<float>(),
|
||||
size,
|
||||
1e-8f
|
||||
);
|
||||
|
||||
return bce_loss;
|
||||
}
|
||||
}
|
||||
"""
|
||||
|
||||
bce_cpp_source = """
|
||||
torch::Tensor bce_cuda(torch::Tensor input, torch::Tensor target, std::string mode);
|
||||
"""
|
||||
|
||||
# 编译CUDA代码
|
||||
bce = load_inline(
|
||||
name="bce",
|
||||
cpp_sources=bce_cpp_source,
|
||||
cuda_sources=bce_source,
|
||||
functions=["bce_cuda"],
|
||||
extra_cuda_cflags=[
|
||||
"-O3",
|
||||
"--use_fast_math",
|
||||
"-gencode=arch=compute_80,code=sm_80"
|
||||
],
|
||||
verbose=True
|
||||
)
|
||||
|
||||
class ModelNew(torch.nn.Module):
|
||||
def __init__(self, mode="fast"):
|
||||
super(ModelNew, self).__init__()
|
||||
self.mode = mode
|
||||
self.bce = bce
|
||||
|
||||
def forward(self, input, target):
|
||||
return self.bce.bce_cuda(input, target, self.mode)
|
||||
|
|
@ -0,0 +1,43 @@
|
|||
import torch
|
||||
import torch.nn as nn
|
||||
|
||||
class Model(nn.Module):
|
||||
"""
|
||||
合理优化的PyTorch BCE Loss实现
|
||||
使用PyTorch内置函数，避免重复的数值处理
|
||||
"""
|
||||
def __init__(self):
|
||||
super(Model, self).__init__()
|
||||
|
||||
def forward(self, input: torch.Tensor, target: torch.Tensor) -> torch.Tensor:
|
||||
"""
|
||||
使用PyTorch内置的binary_cross_entropy函数
|
||||
让PyTorch自己处理数值稳定性
|
||||
|
||||
Args:
|
||||
input (torch.Tensor): 预测概率 (0-1之间)
|
||||
target (torch.Tensor): 真实标签 (0或1)
|
||||
|
||||
Returns:
|
||||
torch.Tensor: BCE Loss标量值
|
||||
"""
|
||||
# 直接使用内置函数，让PyTorch处理数值稳定性
|
||||
return torch.nn.functional.binary_cross_entropy(
|
||||
input,
|
||||
target,
|
||||
reduction='sum'
|
||||
)
|
||||
|
||||
batch_size = 256
|
||||
num_features = 2000
|
||||
|
||||
def get_inputs():
|
||||
"""
|
||||
生成合理的测试数据
|
||||
"""
|
||||
input_probs = torch.sigmoid(torch.randn(batch_size, num_features))
|
||||
target_labels = torch.randint(0, 2, (batch_size, num_features)).float()
|
||||
return [input_probs, target_labels]
|
||||
|
||||
def get_init_inputs():
|
||||
return [] # 没有特殊的初始化输入需求
|
||||
|
|
@ -0,0 +1,196 @@
|
|||
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):
|
||||
“”"
|
||||
合理优化的PyTorch BCE Loss实现
|
||||
使用PyTorch内置函数，避免重复的数值处理
|
||||
“”"
|
||||
def init(self):
|
||||
super(Model, self).init()
|
||||
|
||||
def forward(self, input: torch.Tensor, target: torch.Tensor) -> torch.Tensor:
|
||||
"""
|
||||
使用PyTorch内置的binary_cross_entropy函数
|
||||
让PyTorch自己处理数值稳定性
|
||||
|
||||
Args:
|
||||
input (torch.Tensor): 预测概率 (0-1之间)
|
||||
target (torch.Tensor): 真实标签 (0或1)
|
||||
|
||||
Returns:
|
||||
torch.Tensor: BCE Loss标量值
|
||||
"""
|
||||
# 直接使用内置函数，让PyTorch处理数值稳定性
|
||||
return torch.nn.functional.binary_cross_entropy(
|
||||
input,
|
||||
target,
|
||||
reduction='sum'
|
||||
)
|
||||
batch_size = 128
|
||||
num_features = 2000
|
||||
|
||||
def get_inputs():
|
||||
“”"
|
||||
生成合理的测试数据
|
||||
“”"
|
||||
input_probs = torch.sigmoid(torch.randn(batch_size, num_features))
|
||||
target_labels = torch.randint(0, 2, (batch_size, num_features)).float()
|
||||
return [input_probs, target_labels]
|
||||
|
||||
def get_init_inputs():
|
||||
return [] # 没有特殊的初始化输入需求
|
||||
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):
|
||||
“”"
|
||||
合理优化的PyTorch BCE Loss实现
|
||||
使用PyTorch内置函数，避免重复的数值处理
|
||||
“”"
|
||||
def init(self):
|
||||
super(Model, self).init()
|
||||
|
||||
def forward(self, input: torch.Tensor, target: torch.Tensor) -> torch.Tensor:
|
||||
"""
|
||||
使用PyTorch内置的binary_cross_entropy函数
|
||||
让PyTorch自己处理数值稳定性
|
||||
|
||||
Args:
|
||||
input (torch.Tensor): 预测概率 (0-1之间)
|
||||
target (torch.Tensor): 真实标签 (0或1)
|
||||
|
||||
Returns:
|
||||
torch.Tensor: BCE Loss标量值
|
||||
"""
|
||||
# 直接使用内置函数，让PyTorch处理数值稳定性
|
||||
return torch.nn.functional.binary_cross_entropy(
|
||||
input,
|
||||
target,
|
||||
reduction='sum'
|
||||
)
|
||||
batch_size = 128
|
||||
num_features = 2000
|
||||
|
||||
def get_inputs():
|
||||
“”"
|
||||
生成合理的测试数据
|
||||
“”"
|
||||
input_probs = torch.sigmoid(torch.randn(batch_size, num_features))
|
||||
target_labels = torch.randint(0, 2, (batch_size, num_features)).float()
|
||||
return [input_probs, target_labels]
|
||||
|
||||
def get_init_inputs():
|
||||
return [] # 没有特殊的初始化输入需求
|
||||
|
|
@ -0,0 +1,74 @@
|
|||
###########################################################
|
||||
# 性能和精度验证程序
|
||||
###########################################################
|
||||
import torch
|
||||
import torch.nn as nn
|
||||
import time
|
||||
from bce_torchcode import Model,get_inputs,get_init_inputs
|
||||
from bce_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 BCELoss 平均执行时间: {torch_time:.6f} 秒")
|
||||
print(f"自定义 CUDA BCELoss 平均执行时间: {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