网站首页 > 博客 > 正文

c++ 性能优化 cuda的block同步方法对比

Lyoune 博客 2024-05-20 1 0

先说结论：用最笨的加锁方法实现块同步就行。

cuda9之前都没提供块间同步功能，手工实现块间同步有两种思路：加锁、不加锁，如图所示：

加锁方法只有一个计数器，每个块执行完，用原子操作给计数器加一，然后等待计数器的值到达期望后，再进行下一步。

不加锁方法每个块分配一个计数器，每个块执行完，给自己的计数器加一，然后等待所有计数器都加一后，再进行下一步。

按照我的理解，不加锁方法避免了原子操作，开销会更小，实验验证一下。首先是随便写一个核函数，功能是：每个block计算0到1023的和，再把每个block的结果求和，代码如下：

// 无同步

__global__ void kernel_sum0(int *lock, int *sum) {

// 本地缓存赋值

__shared__ int localSum[1024];

localSum[threadIdx.x] = threadIdx.x; // 赋初值

// 规约

for (int offset=blockDim.x>>1; offset>0; offset>>=1) {

if (threadIdx.x < offset) localSum[threadIdx.x] += localSum[threadIdx.x+offset];

__syncthreads(); // 保证每轮循环前上轮循环都结束了

}

if (threadIdx.x == 0) sum[blockIdx.x] = localSum[0]; // block的结果

// 计算结果

if (threadIdx.x == 0 && blockIdx.x == 0) { // 最后得到结果

int sum0 = 0;

for (int i=0; i

sum[0] = sum0;

}

这个核函数得到的结果是错误的，因为进行block求和之前，有些block可能还没得到结果。所以要在block间同步，保证所有block都得到结果后再求和。用加锁的方法实现一下，代码如下：

// lock方法

__global__ void kernel_sum1(int *lock, int *sum) {

// 本地缓存赋值

__shared__ int localSum[1024];

localSum[threadIdx.x] = threadIdx.x; // 赋初值

// 规约

for (int offset=blockDim.x>>1; offset>0; offset>>=1) {

if (threadIdx.x < offset) localSum[threadIdx.x] += localSum[threadIdx.x+offset];

__syncthreads(); // 保证每轮循环前上轮循环都结束了

}

if (threadIdx.x == 0) sum[blockIdx.x] = localSum[0]; // block的结果

// 等待同步

if (threadIdx.x == 0) {

int lockTemp = lock[0];

lockTemp = (lockTemp/gridDim.x+1)*gridDim.x;

atomicAdd(&lock[0], 1); // 写锁

int volatile temp = lock[0];

while(temp != lockTemp) temp = lock[0]; // 直到所有结果相同

}

__syncthreads();

// 计算结果

if (threadIdx.x == 0 && blockIdx.x == 0) { // 最后得到结果

int sum0 = 0;

for (int i=0; i

sum[0] = sum0;

}

以上代码和原始版本的区别，只是加了等待同步的7行代码。加锁有原子操作，可能对性能有影响，改进一下，代码如下：

// lock-free的方法

__global__ void kernel_sum2(int *lock, int *sum) {

// 本地缓存赋值

__shared__ int localSum[1024];

localSum[threadIdx.x] = threadIdx.x; // 赋初值

// 规约

for (int offset=blockDim.x>>1; offset>0; offset>>=1) {

if (threadIdx.x < offset) localSum[threadIdx.x] += localSum[threadIdx.x+offset];

__syncthreads(); // 保证每轮循环前上轮循环都结束了

}

if (threadIdx.x == 0) sum[blockIdx.x] = localSum[0]; // block的结果

// 等待同步

if (threadIdx.x == 0) {

lock[blockIdx.x] += 1; // 加锁

int lockTemp = lock[blockIdx.x];

for (int i=0; i

int volatile temp = lock[i];

while(temp != lockTemp) temp = lock[i]; // 直到所有结果相同

}

__syncthreads();

// 计算结果

if (threadIdx.x == 0 && blockIdx.x == 0) { // 最后得到结果

int sum0 = 0;

for (int i=0; i

sum[0] = sum0;

}

以上代码和加锁版本相比，只是等待同步部分的8行代码不同。

有了许多版本，可以开始测试了，代码如下：

#include // cout

#include // system_clock

#include // time_t

//--------计时器--------//

class Timer { // 计时器

private:

std::chrono::system_clock::time_point now; // 当前时间

std::time_t time_now; // 时间戳的格式化版

std::chrono::steady_clock::time_point t1; // 计时开始时刻

std::chrono::steady_clock::time_point t2; // 计时结束时刻

std::chrono::duration duration; // 耗时

int p; // 是否暂停了

public:

void getTimeNow() { // 输出当前时间戳

now = std::chrono::system_clock::now();

time_now = std::chrono::system_clock::to_time_t(now);

std::cout << ctime(&time_now);

}

void start() { // 开始计时

t1 = std::chrono::steady_clock::now();

duration = std::chrono::duration(0);

p = 0; // 不暂停

}

void pause() { // 暂停计时

t2 = std::chrono::steady_clock::now();

duration += std::chrono::duration_cast

>(t2-t1);

p = 1;

}

void resume() { // 恢复计时

t1 = std::chrono::steady_clock::now();

p = 0;

}

void getDuration() { // 输出耗时

if (p == 0) pause(); // 没暂停需要先暂停

std::cout << duration.count() << " seconds.\n";

}

};

// 无同步

__global__ void kernel_sum0(int *lock, int *sum) {

// 本地缓存赋值

__shared__ int localSum[1024];

localSum[threadIdx.x] = threadIdx.x; // 赋初值

// 规约

for (int offset=blockDim.x>>1; offset>0; offset>>=1) {

if (threadIdx.x < offset) localSum[threadIdx.x] += localSum[threadIdx.x+offset];

__syncthreads(); // 保证每轮循环前上轮循环都结束了

}

if (threadIdx.x == 0) sum[blockIdx.x] = localSum[0]; // block的结果

// 计算结果

if (threadIdx.x == 0 && blockIdx.x == 0) { // 最后得到结果

int sum0 = 0;

for (int i=0; i

sum[0] = sum0;

}

// lock方法

__global__ void kernel_sum1(int *lock, int *sum) {

// 本地缓存赋值

__shared__ int localSum[1024];

localSum[threadIdx.x] = threadIdx.x; // 赋初值

// 规约

for (int offset=blockDim.x>>1; offset>0; offset>>=1) {

if (threadIdx.x < offset) localSum[threadIdx.x] += localSum[threadIdx.x+offset];

__syncthreads(); // 保证每轮循环前上轮循环都结束了

}

if (threadIdx.x == 0) sum[blockIdx.x] = localSum[0]; // block的结果

// 等待同步

if (threadIdx.x == 0) {

int lockTemp = lock[0];

lockTemp = (lockTemp/gridDim.x+1)*gridDim.x;

atomicAdd(&lock[0], 1); // 写锁

int volatile temp = lock[0];

while(temp != lockTemp) temp = lock[0]; // 直到所有结果相同

}

__syncthreads();

// 计算结果

if (threadIdx.x == 0 && blockIdx.x == 0) { // 最后得到结果

int sum0 = 0;

for (int i=0; i

sum[0] = sum0;

}

// lock-free的方法

__global__ void kernel_sum2(int *lock, int *sum) {

// 本地缓存赋值

__shared__ int localSum[1024];

localSum[threadIdx.x] = threadIdx.x; // 赋初值

// 规约

for (int offset=blockDim.x>>1; offset>0; offset>>=1) {

if (threadIdx.x < offset) localSum[threadIdx.x] += localSum[threadIdx.x+offset];

__syncthreads(); // 保证每轮循环前上轮循环都结束了

}

if (threadIdx.x == 0) sum[blockIdx.x] = localSum[0]; // block的结果

// 等待同步

if (threadIdx.x == 0) {

lock[blockIdx.x] += 1; // 加锁

int lockTemp = lock[blockIdx.x];

for (int i=0; i

int volatile temp = lock[i];

while(temp != lockTemp) temp = lock[i]; // 直到所有结果相同

}

__syncthreads();

// 计算结果

if (threadIdx.x == 0 && blockIdx.x == 0) { // 最后得到结果

int sum0 = 0;

for (int i=0; i

sum[0] = sum0;

}

int main() {

// 声明变量

Timer timer; // 随便用什么自己实现的计时器

int *lock_h, *lock_d, *sum_h, *sum_d;

cudaMallocHost((void**)&lock_h, sizeof(int)*65536);

cudaMalloc((void**)&lock_d, sizeof(int)*65536);

cudaMallocHost((void**)&sum_h, sizeof(int)*65536);

cudaMalloc((void**)&sum_d, sizeof(int)*65536);

// 测试原始函数

for (int i=0; i<1000; i++) kernel_sum0<<<1024, 1024>>>(lock_d, sum_d); // 热身

cudaDeviceSynchronize();

std::cout << "no sync\n";

for (int block=64; block<65537; block*=2) {

// 计时

timer.start();

for (int i=0; i<1000*10; i++) kernel_sum0<<>>(lock_d, sum_d);

cudaDeviceSynchronize();

std::cout << block << ":\t"; timer.getDuration();

}

// 测试加锁函数

for (int i=0; i<1000; i++) kernel_sum1<<<1024, 1024>>>(lock_d, sum_d); // 热身

cudaDeviceSynchronize();

std::cout << "lock sync\n";

for (int block=64; block<65537; block*=2) {

// 计时

timer.start();

for (int i=0; i<1000*10; i++) kernel_sum1<<>>(lock_d, sum_d);

cudaDeviceSynchronize();

std::cout << block << ":\t"; timer.getDuration();

}

// 测试无锁函数

for (int i=0; i<1000; i++) kernel_sum2<<<1024, 1024>>>(lock_d, sum_d); // 热身

cudaDeviceSynchronize();

std::cout << "lock-free sync\n";

for (int block=64; block<65537; block*=2) {

// 计时

timer.start();

for (int i=0; i<1000*10; i++) kernel_sum2<<>>(lock_d, sum_d);

cudaDeviceSynchronize();

std::cout << block << ":\t"; timer.getDuration();

}

// 收尾

cudaMemcpy(sum_h, sum_d, sizeof(int)*1, cudaMemcpyDeviceToHost);

std::cout << "sum:\t" << sum_h[0] << "\n";

return 0;

}

我的测试平台是3090，有82个sm，threads数都是1024，block数冲64到65536，每次翻倍，测试结果如下：

block641282565121024204840968192163843276865536no sync0.040.050.080.130.220.430.841.663.296.5613.10lock sync0.040.050.080.130.230.440.851.683.346.6613.29lock-free sync0.040.050.080.130.250.470.911.803.597.1514.28

上面表格中时间单位是秒，可以看到，最简单的加锁块同步，性能开销很小，实现简单，直接用就行了。

另外如果是cuda9以上的话，还可以用Cooperative Groups协作组，还没仔细学习，看手册上说是可以块同步，性能应该更好，不过代码的可移植性就没了。

查看原文

本文由用户于 2024-05-20 发布在夸智网，如有疑问，请联系我们。
本文链接：https://www.kuazhi.com/post/387943.html

夸智网

c++ 性能优化 cuda的block同步方法对比

性能优化压力测试 MySQL：数据库级监控及常用计数器解析（下）

性能优化硬件架构 DDR理论带宽计算

发表评论取消回复

夸智网

c++ 性能优化 cuda的block同步方法对比

性能优化 压力测试 MySQL：数据库级监控及常用计数器解析（下）

性能优化 硬件架构 DDR理论带宽计算

相关文章

发表评论取消回复

性能优化压力测试 MySQL：数据库级监控及常用计数器解析（下）

性能优化硬件架构 DDR理论带宽计算