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

我只是想這可能是在這裏提供的解決方案，以及有用的。