Programa CUDA faz com que o driver nvidia para falhar

Meu programa CUDA de cálculo Monte Carlo pi está causando meu driver nvidia para falhar quando eu exceder cerca de 500 tentativas e 256 blocos completos. Parece estar acontecendo na function do kernel monteCarlo. Qualquer ajuda é apreciada.

#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 um número menor de tentativas funcionar corretamente e se você estiver executando no MS Windows sem o driver TCC (Compaition Compute Cluster) e / ou a GPU que você está usando estiver conectada a um monitor, provavelmente você estará excedendo o watchdog do sistema operacional. ” tempo esgotado. Se o kernel ocupar o dispositivo de exibição (ou qualquer GPU no Windows sem o TCC) por muito tempo, o SO matará o kernel para que o sistema não se torne não interativo.

A solução é rodar em uma GPU sem display e se você estiver no Windows, use o driver TCC. Caso contrário, você precisará reduzir o número de tentativas em seu kernel e executar o kernel várias vezes para calcular o número de tentativas necessárias.

EDIT: De acordo com os docs curand CUDA 4.0 (página 15, “Performance Notes”), você pode melhorar o desempenho copiando o estado de um gerador para armazenamento local dentro de seu kernel, em seguida, armazenando o estado de volta (se você precisar dele novamente) quando Você terminou:

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

Em seguida, menciona que a configuração é cara e sugere que você mova curand_init para um kernel separado. Isso pode ajudar a manter baixo o custo do seu kernel MC, para que você não enfrente o watchdog.

Eu recomendo ler essa seção dos documentos, existem várias orientações úteis.

Para aqueles que têm uma GPU geforce que não suporta o driver TCC, há outra solução baseada em:

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

  1. inicie o regedit,
  2. navegue até HKEY_LOCAL_MACHINE \ System \ CurrentControlSet \ Control \ GraphicsDrivers
  3. crie uma nova chave DWORD chamada TdrLevel, defina o valor como 0,
  4. reinicie o PC.

Agora seus kernels de longa duração não devem ser terminados. Esta resposta é baseada em:

Modificando o Registro para aumentar o tempo limite da GPU, o Windows 7

Eu apenas pensei que poderia ser útil fornecer a solução aqui também.

    Intereting Posts