c++ CUDA内核确定质数的速度比OpenMP代码慢-如何优化它?

rmbxnbpk  于 2023-02-01  发布在  其他
关注(0)|答案(1)|浏览(120)

为了练习用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编写的程序效率低,如何解释,如何修复?

wljmcqd8

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)来查看你在哪里花费了大部分时间以及如何花费。

相关问题