CUDA程序导致NVIDIA驱动崩溃。

Question

CUDA程序导致NVIDIA驱动崩溃。

4

我的蒙特卡罗π计算CUDA程序在超过500次试验和256个完整块时会导致我的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;
}

- zetatr

2个回答

6

对于那些拥有不支持TCC驱动程序的Geforce GPU的人，可以基于以下解决方案：http://msdn.microsoft.com/en-us/library/windows/hardware/ff569918(v=vs.85).aspx。请按照以下步骤进行操作：1.启动regedit；2.导航到HKEY_LOCAL_MACHINE\System\CurrentControlSet\Control\GraphicsDrivers；3.创建名为TdrLevel的新DWORD键，并将值设置为0；4.重启电脑。现在，您的长时间运行的内核不应该被终止了。此答案基于：修改注册表以增加GPU超时，Windows 7。我认为在这里提供解决方案可能会有所帮助。

- Michal Hosala

如果将显示器连接到这个 GPU 上，会导致系统/图形界面挂起吗？ - Serge Rogatch

@SergeRogatch 是的，我认为是这样。 - Michal Hosala

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- harrism · Accepted Answer

如果较少数量的试验能够正确运行，并且您正在使用MS Windows而没有安装NVIDIA Tesla Compute Cluster（TCC）驱动程序，或者您使用的GPU连接到显示器，则您可能会超出操作系统的“看门狗”超时时间。如果内核占用显示设备（或任何Windows上没有TCC的GPU）时间太长，操作系统将终止内核以使系统保持可交互状态。

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

编辑：根据CUDA 4.0 curand文档（第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内核的成本，从而避免触发看门狗。

我建议阅读文档的这一部分，里面有几个有用的指导方针。