Le programme CUDA provoque le plantage du pilote nvidia

Question

Mon calcul de monte carlo pi Le programme CUDA provoque le crash de mon pilote nvidia lorsque je dépasse environ 500 essais et 256 blocs complets. Il semble se passer dans la fonction noyau de MonteCarlo. Toute aide est appréciée.

#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

Si un plus petit nombre d'essais fonctionnent correctement, et si vous utilisez MS Windows sans le pilote NVIDIA Tesla Compute Cluster (TCC) et/ou le GPU que vous utilisez est attaché à un écran, alors vous êtes dépassant probablement le "time-out" du système d'exploitation. Si le noyau occupe le périphérique d'affichage (ou n'importe quel GPU sous Windows sans TCC) pendant trop longtemps, le système d'exploitation va tuer le noyau pour que le système ne devienne pas interactif.

La solution consiste à exécuter sur un GPU non-display-attached et si vous utilisez Windows, utilisez le pilote TCC. Sinon, vous devrez réduire le nombre d'essais dans votre noyau et exécuter le noyau plusieurs fois pour calculer le nombre d'essais dont vous avez besoin. EDIT: Selon le CUDA 4.0 curand docs (page 15, «Notes sur les performances»), vous pouvez améliorer les performances en copiant l'état d'un générateur sur le stockage local dans votre noyau, puis en le reprenant (si vous en avez besoin) lorsque vous avez terminé:

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++; 
} 

Ensuite, il mentionne que l'installation est coûteuse, et suggère que vous vous déplacez curand_init dans un noyau séparé. Cela peut vous aider à réduire le coût de votre noyau MC afin de ne pas vous heurter au chien de garde.

Je recommande de lire cette section des documents, il y a plusieurs directives utiles.

Answer 2

Pour ceux d'entre vous ayant un GPU GeForce®, ne supporte pas pilote TCC il y a une autre solution basée sur:

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

1. regedit début,
2. accédez à HKEY_LOCAL_MACHINE \ System \ CurrentControlSet \ Control \ GraphicsDrivers
3. Créez une nouvelle clé DWORD appelée TdrLevel, définissez la valeur sur 0,
4. restart PC.

Maintenant, vos noyaux de longue durée ne doivent pas être terminés. Cette réponse est basée sur:

Modifying registry to increase GPU timeout, windows 7

Je pensais qu'il pourrait être utile de fournir la solution ici.