CUDA 程序导致 nvidia 驱动程序崩溃

Posted 2023-04-15

【中文标题】CUDA 程序导致 nvidia 驱动程序崩溃

【英文标题】：CUDA program causes nvidia driver to crash

【发布时间】：2011-09-05 03:29:38

【问题描述】：

当我超过大约 500 次试验和 256 个完整块时，我的 monte carlo pi 计算 CUDA 程序导致我的 nvidia 驱动程序崩溃。它似乎发生在 monteCarlo 核函数中。任何帮助表示赞赏。

#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <curand.h>
#include <curand_kernel.h>


#define NUM_THREAD 256
#define NUM_BLOCK 256



///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////

// Function to sum an array
__global__ void reduce0(float *g_odata) 
extern __shared__ int sdata[];

// each thread loads one element from global to shared mem
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
sdata[tid] = g_odata[i];
__syncthreads();

// do reduction in shared mem
for (unsigned int s=1; s < blockDim.x; s *= 2)  // step = s x 2
    if (tid % (2*s) == 0)  // only threadIDs divisible by the step participate
        sdata[tid] += sdata[tid + s];
    
    __syncthreads();


// write result for this block to global mem
if (tid == 0) g_odata[blockIdx.x] = sdata[0];


///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////
__global__ void monteCarlo(float *g_odata, int  trials, curandState *states)
//  unsigned int tid = threadIdx.x;
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int incircle, k;
    float x, y, z;
    incircle = 0;

    curand_init(1234, i, 0, &states[i]);

    for(k = 0; k < trials; k++)
        x = curand_uniform(&states[i]);
        y = curand_uniform(&states[i]);
        z =(x*x + y*y);
        if (z <= 1.0f) incircle++;
    
    __syncthreads();
    g_odata[i] = incircle;

///////////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////////////////////////////////////////////////////////////////
int main() 

    float* solution = (float*)calloc(100, sizeof(float));
    float *sumDev, *sumHost, total;
    const char *error;
    int trials; 
    curandState *devStates;

    trials = 500;
    total = trials*NUM_THREAD*NUM_BLOCK;

    dim3 dimGrid(NUM_BLOCK,1,1); // Grid dimensions
    dim3 dimBlock(NUM_THREAD,1,1); // Block dimensions
    size_t size = NUM_BLOCK*NUM_THREAD*sizeof(float); //Array memory size
    sumHost = (float*)calloc(NUM_BLOCK*NUM_THREAD, sizeof(float));

    cudaMalloc((void **) &sumDev, size); // Allocate array on device
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);


    cudaMalloc((void **) &devStates, (NUM_THREAD*NUM_BLOCK)*sizeof(curandState));
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);


    // Do calculation on device by calling CUDA kernel
    monteCarlo <<<dimGrid, dimBlock>>> (sumDev, trials, devStates);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);

        // call reduction function to sum
    reduce0 <<<dimGrid, dimBlock, (NUM_THREAD*sizeof(float))>>> (sumDev);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);

    dim3 dimGrid1(1,1,1);
    dim3 dimBlock1(256,1,1);
    reduce0 <<<dimGrid1, dimBlock1, (NUM_THREAD*sizeof(float))>>> (sumDev);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);

    // Retrieve result from device and store it in host array
    cudaMemcpy(sumHost, sumDev, sizeof(float), cudaMemcpyDeviceToHost);
    error = cudaGetErrorString(cudaGetLastError());
    printf("%s\n", error);


    *solution = 4*(sumHost[0]/total);
    printf("%.*f\n", 1000, *solution);
    free (solution);
    free(sumHost);
    cudaFree(sumDev);
    cudaFree(devStates);
    //*solution = NULL;
    return 0;

【参考方案1】：

如果少量试验正常工作，并且如果您在没有 NVIDIA Tesla Compute Cluster (TCC) 驱动程序的 MS Windows 上运行，并且/或者您使用的 GPU 连接到显示器，那么您可能超出了操作系统的“看门狗”超时。如果内核占用显示设备（或没有 TCC 的 Windows 上的任何 GPU）时间过长，操作系统将杀死内核，使系统不会变为非交互状态。

解决方案是在未连接显示器的 GPU 上运行，如果您使用的是 Windows，请使用 TCC 驱动程序。否则，您将需要减少内核中的试验次数并多次运行内核来计算所需的试验次数。

编辑：根据CUDA 4.0 curand docs（第 15 页，“性能说明”），您可以通过将生成器的状态复制到内核中的本地存储，然后将状态存储回来（如果您再次需要它）来提高性能) 完成后：

curandState state = states[i];

for(k = 0; k < trials; k++)
    x = curand_uniform(&state);
    y = curand_uniform(&state);
    z =(x*x + y*y);
    if (z <= 1.0f) incircle++;

接下来，它提到设置很昂贵，并建议您将 curand_init 移到单独的内核中。这可能有助于降低 MC 内核的成本，因此您不会遇到看门狗。

我建议阅读文档的该部分，有几个有用的指南。

【讨论】：

我正在运行 Windows，我的 GPU 连接到显示器。我仍然很惊讶内核需要这么长时间才能完成。 curand_init 和 curand_uniform 调用可能是原因吗？

应该很容易找到——将curand_uniform的调用替换为1.0f，并注释掉curand_init。顺便说一句，你不需要那个__syncthreads()。

感谢您通知我有关同步的信息。此外，你的 curand_uniform 似乎使内核需要更长的时间才能完成。这也很遗憾，因为我什至无法与当前的试验数量很好地融合。运行更多的内核可以让我获得更好的精度，但程序需要更长的时间才能获得不满意的正确数字。

我在我的回答中添加了文档中的一些性能提示——我敢打赌你可以节省时间，这不应该是一个昂贵的内核——curand_uniform 只有几个失败，如果你将状态保存在局部变量中，它将保存在寄存器中。我猜真正的开销是 curand_init()，当您注释掉 curand_uniform() 时，编译器可能会消除死代码，这使得 curand_uniform 看起来很昂贵。将 curand_init 移动到一个单独的内核中，并将状态移动到一个局部变量中，你应该会好很多。不过，您可能需要 x 和 y 的单独状态...

谢谢！这些提示帮助很大。将 curand_init 放入单独的内核使我可以将试验次数增加几个数量级。我还为 y 创建了一个单独的状态数组，具有不同的种子值和 curand_init 调用。这稍微增加了运行时间，但比以前多给了我至少 1 个数字。虽然看起来这个蒙特卡罗收敛速度仍然非常慢，因为我只有 4 个正确的数字，总试验次数超过 13 亿次。

【参考方案2】：

对于那些拥有不支持 TCC 驱动程序的 geforce GPU 的人，还有另一种基于以下解决方案：

http://msdn.microsoft.com/en-us/library/windows/hardware/ff569918(v=vs.85).aspx

开始注册， 导航到 HKEY_LOCAL_MACHINE\System\CurrentControlSet\Control\GraphicsDrivers 创建名为 TdrLevel 的新 DWORD 键，将值设置为 0， 重启电脑。

现在您的长时间运行的内核不应该被终止。此答案基于：

Modifying registry to increase GPU timeout, windows 7

我只是认为在这里也提供解决方案可能很有用。

【讨论】：

如果显示器连接到此 GPU，是否会使系统/图形挂起？

@SergeRogatch 是的，我想。