问题描述
|
我的蒙特卡洛pi计算CUDA程序在我超过500次试用和256个完整块时导致nvidia驱动程序崩溃。它似乎发生在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,&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); // 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);
dim3 dimBlock1(256,1);
reduce0 <<<dimGrid1,dimBlock1,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;
}
解决方法
如果较小数量的试验可以正常工作，并且如果您在不带NVIDIA Tesla Compute Cluster（TCC）驱动程序和/或所使用GPU的MS Windows上运行显示器，则可能是操作系统超出了限制。的“看门狗”超时。如果内核占用显示设备（或Windows中没有TCC的Windows上的任何GPU）的时间过长，则OS将终止内核，因此系统不会变得非交互式。
该解决方案将在非显示器连接的GPU上运行，如果您使用的是Windows，请使用TCC驱动程序。否则，您将需要减少内核中的试用次数，并多次运行内核以计算所需的试用次数。
编辑：根据CUDA 4.0 curand文档（第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++;
}