为了练习用CUDA在C++中编程,我做了一个练习,其中包括显示小于N的质数。对于每个代码,我注解掉了最后一个显示循环,以便只比较计算时间。
生成文件:
all: sum sum_cpu
nothing:
g++ -O3 -std=c++17 -o premier.exe premier.cpp -Wall
cpu:
g++ -O3 -std=c++17 -o cpu_premier.exe premier.cpp -Wall -fopenmp
gpu:
nvcc --compile --compiler-options -O3 -o gpu_premier.o gpu_premier.cu -gencode arch=compute_50,code=sm_50
nvcc --link --compiler-options -O3 -o gpu_premier.exe gpu_premier.o
clear:
rm *.exe *.o
下面是我的代码,用于与1,306秒内运行的OpenMP并行化:
#include <math.h>
#include <iostream>
const int N = 2<<22;
bool * premiers;
bool est_premier(int nbr){
if ( nbr==1 || nbr == 0) return false;
else if (nbr == 2) return true;
else if (nbr % 2 == 0) return false;
else {
for (int i=3;i<=sqrt(nbr);++i){
if (nbr % i == 0){
return false;
}
}
}
return true;
}
int main(){
premiers = new bool[N+1];
# pragma omp parallel for
for (int i = 0;i<N;++i){
premiers[i] = est_premier(i);
}
/*
for (int i = 0;i<N;++i){
if (premiers[i])
std::cout<<i<<",";
} std::cout<<std::endl;
*/
delete[] premiers;
}
下面是对应的cuda代码,运行时间为1,613秒:
#include <cuda.h>
#include <iostream>
const int N = 2<<22;
bool * premiers_cpu;
bool * premiers_gpu;
__device__
bool est_premier(int nbr){
if ( nbr==1 || nbr == 0) return false;
else if (nbr == 2) return true;
else if (nbr % 2 == 0) return false;
else {
for (int i=3;i * i <= nbr ;++i){
if (nbr % i == 0){
return false;
}
}
}
return true;
}
__global__ void kernel_premier(bool * premiers, int size){
int gtid = blockIdx.x * blockDim.x + threadIdx.x ;
while(gtid < size){
premiers[gtid] = est_premier(gtid);
gtid += blockDim.x * gridDim.x;
}
}
int main(){
bool * premiers_cpu = new bool[N];
dim3 block (256,1);
dim3 grid (2048,1,1);
cudaMalloc(( void **) &premiers_gpu, N * sizeof(bool));
cudaMemcpy(premiers_gpu,premiers_cpu,N * sizeof(bool),cudaMemcpyHostToDevice);
kernel_premier<<<grid,block>>>(premiers_gpu,N);
cudaMemcpy(premiers_cpu,premiers_gpu,N * sizeof(bool),cudaMemcpyDeviceToHost);
/*
for (int i = 0;i<N;++i){
if (premiers_cpu[i])
std::cout<<i<<",";
} std::cout<<std::endl;
*/
delete[] premiers_cpu;
cudaFree(premiers_gpu);
}
直觉上,我认为降低网格的大小和增加每个网格的块的数量会使程序更高效,但事实恰恰相反,在这里,我用cuda编写的程序比用OpenMP编写的程序效率低,如何解释,如何修复?
1条答案
按热度按时间wljmcqd81#
首先,我们来看看这个琐碎的观察:你甚至不应该在运行时计算它,只要在编译时设置结果就行了,因为运行时除了打印什么都没有。
此外,还有第二个半平凡的观察结果:确保你没有计时打印!终端I/O是相当慢的。也许这是什么占用了你的大部分时间,两种选择?
最后,更重要的是:线程的分散会降低性能,一半的线程甚至不进入循环,另外1/6的线程只执行一次迭代--而平均而言,一个线程迭代Theta(log((N))次。
但是解决这个问题并不容易,你最好的“选择”是完全改变你的算法(甚至可能在CPU上),并使用类似Erathostenes' sieve的东西:从因式开始,而不是从潜在的素数开始,然后标记非素数。实际上,这也不是一个微不足道的并行化挑战。
无论如何,对于算法,考虑一个更仔细的分离,比如......即使
nbr
s,也不要提前返回--不要迭代它们开始:不要检查gtid
的素性,而应该检查2*gtid + 1
;并将你的网格大小减半。当然,你也可以对CPU做同样的事情--但GPU会受益更多。更一般地说,考虑将块大小设置为这样一个值,即您可以从该值开始排除一些非素数,并且只检查可疑的部分。(不是一个好主意!只是一个说明)那么你可以检查长度为10的数字的素数。为什么呢?因为在序列10 k,10 k +1,10 k +2,10 k +3,10 k +4,10 k +5中,10 k +6、10 k +7、10 k +8、10 k +9 -第1、3、5、7和9个数字是2的倍数;第6个(和第1个)是5的倍数。因此,线程减少了一半,散度也有所降低。进一步(运行235个值,或235*7),使用一个小的常量查找表来选择要检查的值,您将看到一些改进。
警告:这可能还不是你最大的敌人。总是使用分析(例如使用nsight-compute)来查看你在哪里花费了大部分时间以及如何花费。