Il programma CUDA causa l’arresto anomalo del driver nvidia

Il mio programma CUDA per il calcolo del monte carlo pi sta causando il crash del mio driver nvidia quando ho superato circa 500 prove e 256 blocchi completi. Sembra che stia accadendo nella funzione del kernel di monteCarlo. Qualsiasi aiuto è apprezzato.

#include  #include  #include  #include  #include  #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 = sx 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 <<>> (sumDev, trials, devStates); error = cudaGetErrorString(cudaGetLastError()); printf("%s\n", error); // call reduction function to sum reduce0 <<>> (sumDev); error = cudaGetErrorString(cudaGetLastError()); printf("%s\n", error); dim3 dimGrid1(1,1,1); dim3 dimBlock1(256,1,1); reduce0 <<>> (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; } 

Se numeri più piccoli di prove funzionano correttamente, e se si sta lavorando su MS Windows senza il driver NVIDIA Tesla Compute Cluster (TCC) e / o la GPU che si sta utilizzando è collegato a un display, probabilmente si sta superando il “watchdog” del sistema operativo. ” tempo scaduto. Se il kernel occupa il dispositivo di visualizzazione (o qualsiasi GPU su Windows senza TCC) per troppo tempo, il sistema operativo ucciderà il kernel in modo che il sistema non diventi non interattivo.

La soluzione è quella di girare su una GPU senza display e se sei su Windows, usa il driver TCC. Altrimenti, sarà necessario ridurre il numero di prove nel kernel ed eseguire il kernel più volte per calcolare il numero di prove necessarie.

EDIT: Secondo i documenti di curand CUDA 4.0 (pagina 15, “Note sulle prestazioni”), è ansible migliorare le prestazioni copiando lo stato di un generatore sulla memoria locale all’interno del kernel, quindi memorizzando lo stato indietro (se ne hai bisogno di nuovo) quando sei finito:

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

Successivamente, viene menzionato che l'installazione è costosa e suggerisce di spostare curand_init in un kernel separato. Ciò può aiutare a mantenere basso il costo del tuo MC MC in modo da non essere in grado di funzionare contro il watchdog.

Consiglio di leggere quella sezione dei documenti, ci sono diverse linee guida utili.

Per quelli di voi che hanno una GPU Geforce che non supporta il driver TCC, c’è un’altra soluzione basata su:

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

  1. inizia regedit,
  2. passare a HKEY_LOCAL_MACHINE \ System \ CurrentControlSet \ Control \ GraphicsDrivers
  3. crea una nuova chiave DWORD chiamata TdrLevel, imposta il valore su 0,
  4. riavviare PC.

Ora i tuoi kernel a lunga durata non dovrebbero essere terminati. Questa risposta è basata su:

Modifica del registro per aumentare il timeout della GPU, Windows 7

Ho solo pensato che potrebbe essere utile fornire la soluzione anche qui.