static CUT_THREADPROC solverThread(TOptionPlan *plan) {
//Init GPU
cutilSafeCall( cudaSetDevice(plan->device) );
cudaDeviceProp deviceProp;
cutilSafeCall(cudaGetDeviceProperties(&deviceProp, plan->device));
int version = deviceProp.major * 10 + deviceProp.minor;
if(useDoublePrecision && version < 13) {
printf("Double precision is not supported on device %i.\n", plan->device);
exit(0);
}
//Allocate memory for normally distributed samples
cutilSafeCall( cudaMalloc(
(void **)&plan->d_Samples,
plan->pathN * sizeof(float)
) );
//Generate normally distributed samples
if(useDoublePrecision)
inverseCND_SM13(plan->d_Samples, NULL, plan->pathN);
else
inverseCND_SM10(plan->d_Samples, NULL, plan->pathN);
//Allocate intermediate memory for MC integrator
if(useDoublePrecision)
initMonteCarlo_SM13(plan);
else
initMonteCarlo_SM10(plan);
//Main computations
if(useDoublePrecision)
MonteCarlo_SM13(plan);
else
MonteCarlo_SM10(plan);
cutilSafeCall( cudaThreadSynchronize() );
//Shut down this GPU
if(useDoublePrecision)
closeMonteCarlo_SM13(plan);
else
closeMonteCarlo_SM10(plan);
cutilSafeCall( cudaFree(plan->d_Samples) );
cudaThreadExit();
CUT_THREADEND;
}
int main(int argc, char **argv)
{
// Start logs
printf("%s Starting...\n\n", argv[0]);
unsigned int useDoublePrecision;
char *precisionChoice;
getCmdLineArgumentString(argc, (const char **)argv, "type", &precisionChoice);
if (precisionChoice == NULL)
{
useDoublePrecision = 0;
}
else
{
if (!STRCASECMP(precisionChoice, "double"))
{
useDoublePrecision = 1;
}
else
{
useDoublePrecision = 0;
}
}
unsigned int tableCPU[QRNG_DIMENSIONS][QRNG_RESOLUTION];
float *h_OutputGPU, *d_Output;
int dim, pos;
double delta, ref, sumDelta, sumRef, L1norm, gpuTime;
StopWatchInterface *hTimer = NULL;
if (sizeof(INT64) != 8)
{
printf("sizeof(INT64) != 8\n");
return 0;
}
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
int dev = findCudaDevice(argc, (const char **)argv);
sdkCreateTimer(&hTimer);
int deviceIndex;
checkCudaErrors(cudaGetDevice(&deviceIndex));
cudaDeviceProp deviceProp;
checkCudaErrors(cudaGetDeviceProperties(&deviceProp, deviceIndex));
int version = deviceProp.major * 10 + deviceProp.minor;
if (useDoublePrecision && version < 13)
{
printf("Double precision not supported.\n");
// cudaDeviceReset causes the driver to clean up all state. While
// not mandatory in normal operation, it is good practice. It is also
// needed to ensure correct operation when the application is being
// profiled. Calling cudaDeviceReset causes all profile data to be
// flushed before the application exits
cudaDeviceReset();
return 0;
}
printf("Allocating GPU memory...\n");
checkCudaErrors(cudaMalloc((void **)&d_Output, QRNG_DIMENSIONS * N * sizeof(float)));
printf("Allocating CPU memory...\n");
h_OutputGPU = (float *)malloc(QRNG_DIMENSIONS * N * sizeof(float));
printf("Initializing QRNG tables...\n\n");
initQuasirandomGenerator(tableCPU);
if (useDoublePrecision)
{
initTable_SM13(tableCPU);
}
else
{
initTable_SM10(tableCPU);
}
printf("Testing QRNG...\n\n");
checkCudaErrors(cudaMemset(d_Output, 0, QRNG_DIMENSIONS * N * sizeof(float)));
int numIterations = 20;
for (int i = -1; i < numIterations; i++)
{
if (i == 0)
{
checkCudaErrors(cudaDeviceSynchronize());
sdkResetTimer(&hTimer);
sdkStartTimer(&hTimer);
}
if (useDoublePrecision)
{
quasirandomGenerator_SM13(d_Output, 0, N);
}
else
{
quasirandomGenerator_SM10(d_Output, 0, N);
}
}
checkCudaErrors(cudaDeviceSynchronize());
sdkStopTimer(&hTimer);
gpuTime = sdkGetTimerValue(&hTimer)/(double)numIterations*1e-3;
printf("quasirandomGenerator, Throughput = %.4f GNumbers/s, Time = %.5f s, Size = %u Numbers, NumDevsUsed = %u, Workgroup = %u\n",
(double)QRNG_DIMENSIONS * (double)N * 1.0E-9 / gpuTime, gpuTime, QRNG_DIMENSIONS*N, 1, 128*QRNG_DIMENSIONS);
printf("\nReading GPU results...\n");
checkCudaErrors(cudaMemcpy(h_OutputGPU, d_Output, QRNG_DIMENSIONS * N * sizeof(float), cudaMemcpyDeviceToHost));
printf("Comparing to the CPU results...\n\n");
sumDelta = 0;
sumRef = 0;
for (dim = 0; dim < QRNG_DIMENSIONS; dim++)
for (pos = 0; pos < N; pos++)
{
ref = getQuasirandomValue63(pos, dim);
delta = (double)h_OutputGPU[dim * N + pos] - ref;
sumDelta += fabs(delta);
sumRef += fabs(ref);
}
printf("L1 norm: %E\n", sumDelta / sumRef);
printf("\nTesting inverseCNDgpu()...\n\n");
checkCudaErrors(cudaMemset(d_Output, 0, QRNG_DIMENSIONS * N * sizeof(float)));
for (int i = -1; i < numIterations; i++)
{
if (i == 0)
{
checkCudaErrors(cudaDeviceSynchronize());
sdkResetTimer(&hTimer);
sdkStartTimer(&hTimer);
}
if (useDoublePrecision)
{
inverseCND_SM13(d_Output, NULL, QRNG_DIMENSIONS * N);
}
else
{
inverseCND_SM10(d_Output, NULL, QRNG_DIMENSIONS * N);
}
}
checkCudaErrors(cudaDeviceSynchronize());
sdkStopTimer(&hTimer);
gpuTime = sdkGetTimerValue(&hTimer)/(double)numIterations*1e-3;
printf("quasirandomGenerator-inverse, Throughput = %.4f GNumbers/s, Time = %.5f s, Size = %u Numbers, NumDevsUsed = %u, Workgroup = %u\n",
(double)QRNG_DIMENSIONS * (double)N * 1E-9 / gpuTime, gpuTime, QRNG_DIMENSIONS*N, 1, 128);
printf("Reading GPU results...\n");
checkCudaErrors(cudaMemcpy(h_OutputGPU, d_Output, QRNG_DIMENSIONS * N * sizeof(float), cudaMemcpyDeviceToHost));
printf("\nComparing to the CPU results...\n");
sumDelta = 0;
sumRef = 0;
unsigned int distance = ((unsigned int)-1) / (QRNG_DIMENSIONS * N + 1);
for (pos = 0; pos < QRNG_DIMENSIONS * N; pos++)
{
unsigned int d = (pos + 1) * distance;
ref = MoroInvCNDcpu(d);
delta = (double)h_OutputGPU[pos] - ref;
sumDelta += fabs(delta);
sumRef += fabs(ref);
}
printf("L1 norm: %E\n\n", L1norm = sumDelta / sumRef);
printf("Shutting down...\n");
sdkDeleteTimer(&hTimer);
free(h_OutputGPU);
checkCudaErrors(cudaFree(d_Output));
// cudaDeviceReset causes the driver to clean up all state. While
// not mandatory in normal operation, it is good practice. It is also
// needed to ensure correct operation when the application is being
// profiled. Calling cudaDeviceReset causes all profile data to be
// flushed before the application exits
cudaDeviceReset();
exit(L1norm < 1e-6 ? EXIT_SUCCESS : EXIT_FAILURE);
}
int main(int argc, char **argv)
{
// Start logs
shrSetLogFileName ("quasirandomGenerator.txt");
shrLog("%s Starting...\n\n", argv[0]);
unsigned int useDoublePrecision;
char *precisionChoice;
cutGetCmdLineArgumentstr(argc, (const char **)argv, "type", &precisionChoice);
if(precisionChoice == NULL)
useDoublePrecision = 0;
else{
if(!strcasecmp(precisionChoice, "double"))
useDoublePrecision = 1;
else
useDoublePrecision = 0;
}
unsigned int tableCPU[QRNG_DIMENSIONS][QRNG_RESOLUTION];
float
*h_OutputGPU;
float
*d_Output;
int
dim, pos;
double
delta, ref, sumDelta, sumRef, L1norm, gpuTime;
unsigned int hTimer;
if(sizeof(INT64) != 8){
shrLog("sizeof(INT64) != 8\n");
return 0;
}
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") )
cutilDeviceInit(argc, argv);
else
cudaSetDevice( cutGetMaxGflopsDeviceId() );
cutilCheckError(cutCreateTimer(&hTimer));
int deviceIndex;
cutilSafeCall(cudaGetDevice(&deviceIndex));
cudaDeviceProp deviceProp;
cutilSafeCall(cudaGetDeviceProperties(&deviceProp, deviceIndex));
int version = deviceProp.major * 10 + deviceProp.minor;
if(useDoublePrecision && version < 13){
shrLog("Double precision not supported.\n");
cudaThreadExit();
return 0;
}
shrLog("Allocating GPU memory...\n");
cutilSafeCall( cudaMalloc((void **)&d_Output, QRNG_DIMENSIONS * N * sizeof(float)) );
shrLog("Allocating CPU memory...\n");
h_OutputGPU = (float *)malloc(QRNG_DIMENSIONS * N * sizeof(float));
shrLog("Initializing QRNG tables...\n\n");
initQuasirandomGenerator(tableCPU);
if(useDoublePrecision)
initTable_SM13(tableCPU);
else
initTable_SM10(tableCPU);
shrLog("Testing QRNG...\n\n");
cutilSafeCall( cudaMemset(d_Output, 0, QRNG_DIMENSIONS * N * sizeof(float)) );
int numIterations = 20;
for (int i = -1; i < numIterations; i++)
{
if (i == 0)
{
cutilSafeCall( cudaThreadSynchronize() );
cutilCheckError( cutResetTimer(hTimer) );
cutilCheckError( cutStartTimer(hTimer) );
}
if(useDoublePrecision)
quasirandomGenerator_SM13(d_Output, 0, N);
else
quasirandomGenerator_SM10(d_Output, 0, N);
}
cutilSafeCall( cudaThreadSynchronize() );
cutilCheckError(cutStopTimer(hTimer));
gpuTime = cutGetTimerValue(hTimer)/(double)numIterations*1e-3;
shrLogEx(LOGBOTH | MASTER, 0, "quasirandomGenerator, Throughput = %.4f GNumbers/s, Time = %.5f s, Size = %u Numbers, NumDevsUsed = %u, Workgroup = %u\n",
(double)QRNG_DIMENSIONS * (double)N * 1.0E-9 / gpuTime, gpuTime, QRNG_DIMENSIONS*N, 1, 128*QRNG_DIMENSIONS);
shrLog("\nReading GPU results...\n");
cutilSafeCall( cudaMemcpy(h_OutputGPU, d_Output, QRNG_DIMENSIONS * N * sizeof(float), cudaMemcpyDeviceToHost) );
shrLog("Comparing to the CPU results...\n\n");
sumDelta = 0;
sumRef = 0;
for(dim = 0; dim < QRNG_DIMENSIONS; dim++)
for(pos = 0; pos < N; pos++){
ref = getQuasirandomValue63(pos, dim);
delta = (double)h_OutputGPU[dim * N + pos] - ref;
sumDelta += fabs(delta);
sumRef += fabs(ref);
}
shrLog("L1 norm: %E\n", sumDelta / sumRef);
shrLog("\nTesting inverseCNDgpu()...\n\n");
cutilSafeCall( cudaMemset(d_Output, 0, QRNG_DIMENSIONS * N * sizeof(float)) );
for (int i = -1; i < numIterations; i++)
{
if (i == 0)
{
cutilSafeCall( cudaThreadSynchronize() );
cutilCheckError( cutResetTimer(hTimer) );
cutilCheckError( cutStartTimer(hTimer) );
}
if(useDoublePrecision)
inverseCND_SM13(d_Output, NULL, QRNG_DIMENSIONS * N);
else
inverseCND_SM10(d_Output, NULL, QRNG_DIMENSIONS * N);
}
cutilSafeCall( cudaThreadSynchronize() );
cutilCheckError(cutStopTimer(hTimer));
gpuTime = cutGetTimerValue(hTimer)/(double)numIterations*1e-3;
shrLogEx(LOGBOTH | MASTER, 0, "quasirandomGenerator-inverse, Throughput = %.4f GNumbers/s, Time = %.5f s, Size = %u Numbers, NumDevsUsed = %u, Workgroup = %u\n",
(double)QRNG_DIMENSIONS * (double)N * 1E-9 / gpuTime, gpuTime, QRNG_DIMENSIONS*N, 1, 128);
shrLog("Reading GPU results...\n");
cutilSafeCall( cudaMemcpy(h_OutputGPU, d_Output, QRNG_DIMENSIONS * N * sizeof(float), cudaMemcpyDeviceToHost) );
shrLog("\nComparing to the CPU results...\n");
sumDelta = 0;
sumRef = 0;
for(pos = 0; pos < QRNG_DIMENSIONS * N; pos++){
double p = (double)(pos + 1) / (double)(QRNG_DIMENSIONS * N + 1);
ref = MoroInvCNDcpu(p);
delta = (double)h_OutputGPU[pos] - ref;
sumDelta += fabs(delta);
sumRef += fabs(ref);
}
shrLog("L1 norm: %E\n\n", L1norm = sumDelta / sumRef);
shrLog((L1norm < 1E-6) ? "PASSED\n\n" : "FAILED\n\n");
shrLog("Shutting down...\n");
cutilCheckError(cutDeleteTimer(hTimer));
free(h_OutputGPU);
cutilSafeCall( cudaFree(d_Output) );
cudaThreadExit();
shrEXIT(argc, (const char**)argv);
}