Do While 不在 CUDA 内核中工作

Question

好吧，我是 CUDA 的新手，我有点迷茫，真的很迷茫。

我正在尝试使用 Monte Carlo 方法计算 pi，最后我只得到一个加法而不是 50。

我不想 "do while" 调用内核，因为它太慢了。我的问题是，我的代码没有循环，它只在内核中执行一次。

而且，我希望所有线程都访问相同的 niter 和 pi，这样当某个线程命中计数器时，所有其他线程都会停止。

#define SEED 35791246

__shared__ int niter;
__shared__ double pi;

__global__ void calcularPi(){

    double x;
    double y;
    int count;
    double z;

    count = 0;
    niter = 0;

    //keep looping
    do{

        niter = niter + 1;

        //Generate random number
        curandState state;
        curand_init(SEED,(int)niter, 0, &state);
        x = curand(&state);
        y = curand(&state);

        z = x*x+y*y;
        if (z<=1) count++;
     pi =(double)count/niter*4;

    }while(niter < 50);

}

int main(void){

    float tempoTotal;
    //Start timer
    clock_t t;
    t = clock();

    //call kernel
    calcularPi<<<1,32>>>();

    //wait while kernel finish
    cudaDeviceSynchronize();

    typeof(pi) piFinal;
    cudaMemcpyFromSymbol(&piFinal, "pi", sizeof(piFinal),0, cudaMemcpyDeviceToHost);

    typeof(niter) niterFinal;
    cudaMemcpyFromSymbol(&niterFinal, "niter", sizeof(niterFinal),0, cudaMemcpyDeviceToHost);

    //Ends timer
    t = clock() - t;
    tempoTotal = ((double)t)/CLOCKS_PER_SEC;
    printf("Pi: %g \n", piFinal);
    printf("Adds: %d \n", niterFinal);
    printf("Total time: %f \n", tempoTotal);

}

Answer 1

您的代码存在各种问题。

我建议使用 proper cuda error checking 和运行您的代码以及 cuda-memcheck 来发现任何运行时间错误。为了简洁起见，我在下面的代码中省略了正确的错误检查，但我用 cuda-memcheck 运行表示没有运行时间错误。
您对 curand() 的用法可能不正确（它是 returns 大范围内的整数）。为了使此代码正常工作，您需要一个介于 0 和 1 之间的浮点数。correct call for that 是 curand_uniform()。
由于您希望所有线程都处理相同的值，因此必须防止这些线程相互踩踏。一种方法是使用相关变量的原子更新。
应该没有必要在每次迭代时重新运行 curand_init。每个线程一次就足够了。
我们不对 __shared__ 变量使用 cudaMemcpy..Symbol 操作。为方便起见，并保留与原始代码相似的内容，我选择将它们转换为 __device__ 变量。

这是您的代码的修改版本，已修复上述大部分问题：

$ cat t978.cu
#include <curand.h>
#include <curand_kernel.h>
#include <stdio.h>

#define ITER_MAX 5000
#define SEED 35791246

__device__ int niter;
__device__ int count;

__global__ void calcularPi(){

    double x;
    double y;
    double z;
    int lcount;
    curandState state;
    curand_init(SEED,threadIdx.x, 0, &state);
    //keep looping
    do{

        lcount = atomicAdd(&niter, 1);

        //Generate random number
        x = curand_uniform(&state);
        y = curand_uniform(&state);

        z = x*x+y*y;
        if (z<=1) atomicAdd(&count, 1);

    }while(lcount < ITER_MAX);

}

int main(void){

    float tempoTotal;
    //Start timer
    clock_t t;
    t = clock();
    int count_final = 0;
    int niter_final = 0;
    cudaMemcpyToSymbol(niter, &niter_final, sizeof(int));
    cudaMemcpyToSymbol(count, &count_final, sizeof(int));
    //call kernel
    calcularPi<<<1,32>>>();

    //wait while kernel finish
    cudaDeviceSynchronize();
    cudaMemcpyFromSymbol(&count_final, count, sizeof(int));
    cudaMemcpyFromSymbol(&niter_final, niter, sizeof(int));

    //Ends timer
    double pi = count_final/(double)niter_final*4;
    t = clock() - t;
    tempoTotal = ((double)t)/CLOCKS_PER_SEC;
    printf("Pi: %g \n", pi);
    printf("Adds: %d \n", niter_final);
    printf("Total time: %f \n", tempoTotal);

}
$ nvcc -o t978 t978.cu -lcurand
$ cuda-memcheck ./t978
========= CUDA-MEMCHECK
Pi: 3.12083
Adds: 5032
Total time: 0.558463
========= ERROR SUMMARY: 0 errors
$

我已将迭代次数修改为更大的数字，但如果需要，可以使用 50 ITER_MAX。

请注意，可以针对此代码提出许多批评。我的目标是，因为这显然是一个学习练习，所以我要指出使用您概述的算法获得功能代码所需的最少更改数量。仅作为一个示例，您可能希望将内核启动配置 (<<<1,32>>>) 更改为其他更大的数字，以便更充分地利用 GPU。

Do While 不在 CUDA 内核中工作

Do While don't work inside CUDA Kernel

c

cuda

gpgpu