c语言 - 在 CUDA 内核中不起作用时执行

Question

好吧，我是CUDA的新手，我有点迷失了，真的迷失了。

我试图用蒙特卡罗方法计算圆周率，最后我只得到一个加法，而不是50。

我不想"做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

您的代码存在各种问题。

1. 我建议使用适当的cuda错误检查，并使用cuda-memcheck运行代码来发现任何运行时错误。为了简洁起见，我在下面的代码中省略了正确的错误检查，但我使用cuda-memcheck运行它以指示没有运行时错误。

2. 您对curand()的使用可能不正确（它返回的整数范围很大）。为了使此代码正常工作，您需要一个介于0和1之间的浮点量。正确的调用是curand_uniform()。

3. 由于您希望所有线程都使用相同的值，因此必须防止这些线程相互踩踏。一种方法是使用有问题的变量的原子更新。

4. 不需要在每次迭代中重新运行curand_init。每个线程一次就足够了。

5. 我们不对__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
$

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

请注意，有许多批评可能会针对此代码提出。我在这里的目的是，因为这显然是一个学习练习，使用你概述的算法，指出获得函数代码的最小更改次数。举个例子，您可能希望将内核启动配置（<<<1,32>>>）更改为其他更大的数字，以便更充分地利用GPU。