c++ - 用于确定素数的 CUDA 内核比 OpenMP 代码慢 - 我该如何优化它？

Question

练习使用 C++ 进行 CUDA 编程。我做了一个练习，其中包括显示小于 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

这是我与 openMP 并行化的代码，运行时间为 1,306 秒:

#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);
    


}

直觉上，我认为减小网格的大小并增加每个网格的 block 数将使程序更加高效，但事实恰恰相反。 这里我在 cuda 中的程序比我在 OpenMP 中的程序效率低，如何解释它以及如何修复它？

Answer 1

首先，简单的观察是:您甚至不应该在运行时计算它。只需在编译时设置结果，运行时除了打印之外什么也没有。

另外，第二个半琐碎的观察:确保你没有计时打印!终端 I/O 非常慢。也许这就是您在两种选择中花费大部分时间的原因？

最后，更重要的是:线程的分歧会降低你的性能。一半的线程甚至不进入循环，另外 1/6 只执行一次迭代 - 而平均而言，线程迭代 Theta(log((N)) 次。

但是解决这个问题非常不简单。您最好的“选择”是完全切换您的算法(甚至可能在 CPU 上)，并使用类似 Erathostenes' sieve 的东西。 :从因子开始，而不是潜在的素数，并标记非素数。实际上，这也不是一个微不足道的并行化挑战。

无论如何，按照算法的原样，考虑更仔细的分离。说...而不是提前返回甚至 nbrs - 不要从一开始就迭代它们:不要检查 gtid 的素数，而是检查 2 *gtid + 1;并将网格大小减半。当然，您可以对 CPU 执行相同的操作 - 但 GPU 受益更多。

更一般地说，考虑将 block 大小设置为一个值，您可以先排除一些非素数，然后仅检查可疑的值。例如，如果 block 大小为 5(不是一个好主意!只是一个说明)，那么您可以检查 长度为 10 的数字的素性。怎么会？因为在序列 10k,10k+1,10k+2,10k+3,10k+4,10k+5,10k+6,10k+7,10k+8,10k+9 中 - 第 1 个、第 3 个、第5个、第7个、第9个数字是2的倍数；第 6 个(和第 1 个)是 5 的倍数。因此，线程数减半，并且发散有所减少。更进一步(运行 235 个值，或 235*7)，使用一个小的常量查找表来选择要检查的值，您将看到一些改进。

警告:这可能仍然不是您最大的敌人。始终使用分析(例如使用 nsight-compute)来查看您大部分时间都花在哪里以及如何花费。