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

Question

我的蒙特卡罗pi计算CUDA程序导致我的nvidia驱动程序崩溃时，我超过了大约500次试验和256个满块。它似乎发生在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; 
} 

Answer 1

如果试验的更小的数字正确地工作，并且如果要在MS Windows而不NVIDIA特斯拉计算群集（TCC）驱动器和/或正在使用的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内核的成本，因此您不会碰到看门狗。

我推荐阅读文档的那一部分，有几条有用的指导原则。

Answer 2

对于那些你有一块GeForce GPU不支持TCC驱动程序还有另一种解决方案基于：

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

1. 启动注册表编辑器，
2. 定位到HKEY_LOCAL_MACHINE \系统\ CurrentControlSet \ Control \ GraphicsDrivers
3. 创建新的DWORD密钥，称为TdrLevel，将值设置为0,
4. 重新启动PC。

现在你的长时间运行的内核不应该被终止。这个答案是基于：

Modifying registry to increase GPU timeout, windows 7

我只是想这可能是在这里提供的解决方案，以及有用的。