CUTBoolean initGL(int argc, char **argv)
{
glutInit(&argc, argv);
glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE);
glutInitWindowSize(10, 10);
int bla = glutCreateWindow("Cuda GL Interop (VBO)");
glutDisplayFunc(dumm_display);
// initialize necessary OpenGL extensions
glewInit();
if (! glewIsSupported("GL_VERSION_2_0 ")) {
fprintf(stderr, "ERROR: Support for necessary OpenGL extensions missing.");
fflush(stderr);
return CUTFalse;
}
// default initialization
glClearColor(0.0, 0.0, 0.0, 1.0);
glDisable(GL_DEPTH_TEST);
// viewport
glViewport(0, 0, 10, 10);
// projection
glMatrixMode(GL_PROJECTION);
glLoadIdentity();
gluPerspective(60.0, (GLfloat)10 / (GLfloat) 10, 0.1, 10.0);
CUT_CHECK_ERROR_GL();
// start gui for the main application
//cudaError_t error = cudaGLSetGLDevice(0);
//cutilGLDeviceInit(argc, argv);
int deviceCount;
cutilSafeCallNoSync(cudaGetDeviceCount(&deviceCount));
if (deviceCount == 0) {
fprintf(stderr, "CUTIL CUDA error: no devices supporting CUDA.\n");
exit(-1);
}
int dev = 0;
cudaDeviceProp deviceProp;
cutilSafeCallNoSync(cudaGetDeviceProperties(&deviceProp, dev));
if (deviceProp.major < 1) {
fprintf(stderr, "cutil error: device does not support CUDA.\n");
exit(-1);
}
printf("gpu=%s\n", deviceProp.name);
cutilSafeCall(cudaGLSetGLDevice(dev));
glutDestroyWindow(bla);
return CUTTrue;
}
void SetSpots(unsigned int* cuda_int_dest)
{
int kk = 0;
for(int i=0; i < mmGridSize; i++)
{
if(mmGrid[i].ind[kk] >= 0 && mmGrid[i].v[kk]*(1-2*kk) >= 0)
{
unsigned int val = 255<<8;
cutilSafeCallNoSync( cudaMemcpy( cuda_int_dest + mmGrid[i].ind[kk], &val, sizeof(int), cudaMemcpyHostToDevice) );
}
}
int delta = 128;
for(unsigned int i=0; i < sim_height/delta; i++)
{
unsigned int val = 255;
cutilSafeCallNoSync( cudaMemset(cuda_int_dest + i*delta*sim_width, val, sim_width*sizeof(int)) );
}
}
bool
runTestMax( int argc, char** argv, ReduceType datatype)
{
int size = 1<<24; // number of elements to reduce
int maxThreads = 256; // number of threads per block
int whichKernel = 6;
int maxBlocks = 64;
bool cpuFinalReduction = false;
int cpuFinalThreshold = 1;
cutGetCmdLineArgumenti( argc, (const char**) argv, "n", &size);
cutGetCmdLineArgumenti( argc, (const char**) argv, "threads", &maxThreads);
cutGetCmdLineArgumenti( argc, (const char**) argv, "kernel", &whichKernel);
cutGetCmdLineArgumenti( argc, (const char**) argv, "maxblocks", &maxBlocks);
shrLog("METHOD: MAX\n");
shrLog("%d elements\n", size);
shrLog("%d threads (max)\n", maxThreads);
cpuFinalReduction = (cutCheckCmdLineFlag( argc, (const char**) argv, "cpufinal") == CUTTrue);
cutGetCmdLineArgumenti( argc, (const char**) argv, "cputhresh", &cpuFinalThreshold);
bool runShmoo = (cutCheckCmdLineFlag(argc, (const char**) argv, "shmoo") == CUTTrue);
if (runShmoo)
{
shmoo<T>(1, 33554432, maxThreads, maxBlocks, datatype);
}
else
{
// create random input data on CPU
unsigned int bytes = size * sizeof(T);
T *h_idata = (T *) malloc(bytes);
for(int i=0; i<size; i++)
{
// Keep the numbers small so we don't get truncation error in the sum
if (datatype == REDUCE_INT)
h_idata[i] = (T)(rand() & 0xFF);
else
h_idata[i] = (rand() & 0xFF) / (T)RAND_MAX;
}
int numBlocks = 0;
int numThreads = 0;
getNumBlocksAndThreads(whichKernel, size, maxBlocks, maxThreads, numBlocks, numThreads);
if (numBlocks == 1) cpuFinalThreshold = 1;
// allocate mem for the result on host side
T* h_odata = (T*) malloc(numBlocks*sizeof(T));
shrLog("%d blocks\n\n", numBlocks);
// allocate device memory and data
T* d_idata = NULL;
T* d_odata = NULL;
cutilSafeCallNoSync( cudaMalloc((void**) &d_idata, bytes) );
cutilSafeCallNoSync( cudaMalloc((void**) &d_odata, numBlocks*sizeof(T)) );
// copy data directly to device memory
cutilSafeCallNoSync( cudaMemcpy(d_idata, h_idata, bytes, cudaMemcpyHostToDevice) );
cutilSafeCallNoSync( cudaMemcpy(d_odata, h_idata, numBlocks*sizeof(T), cudaMemcpyHostToDevice) );
// warm-up
maxreduce<T>(size, numThreads, numBlocks, whichKernel, d_idata, d_odata);
int testIterations = 100;
unsigned int timer = 0;
cutilCheckError( cutCreateTimer( &timer));
T gpu_result = 0;
gpu_result = benchmarkReduceMax<T>(size, numThreads, numBlocks, maxThreads, maxBlocks,
whichKernel, testIterations, cpuFinalReduction,
cpuFinalThreshold, timer,
h_odata, d_idata, d_odata);
double reduceTime = cutGetAverageTimerValue(timer) * 1e-3;
shrLogEx(LOGBOTH | MASTER, 0, "Reduction, Throughput = %.4f GB/s, Time = %.5f s, Size = %u Elements, NumDevsUsed = %d, Workgroup = %u\n",
1.0e-9 * ((double)bytes)/reduceTime, reduceTime, size, 1, numThreads);
// compute reference solution
T cpu_result = maxreduceCPU<T>(h_idata, size);
double threshold = 1e-12;
double diff = 0;
if (datatype == REDUCE_INT)
{
shrLog("\nGPU result = %d\n", gpu_result);
shrLog("CPU result = %d\n\n", cpu_result);
}
else
{
shrLog("\nGPU result = %f\n", gpu_result);
shrLog("CPU result = %f\n\n", cpu_result);
if (datatype == REDUCE_FLOAT)
threshold = 1e-8 * size;
diff = fabs((double)gpu_result - (double)cpu_result);
}
// cleanup
cutilCheckError( cutDeleteTimer(timer) );
free(h_idata);
free(h_odata);
cutilSafeCallNoSync(cudaFree(d_idata));
cutilSafeCallNoSync(cudaFree(d_odata));
if (datatype == REDUCE_INT) {
return (gpu_result == cpu_result);
} else {
return (diff < threshold);
}
}
return true;
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int
main( int argc, char** argv)
{
shrQAStart( argc, argv );
shrSetLogFileName ("reduction.txt");
char *reduceMethod;
cutGetCmdLineArgumentstr( argc, (const char**) argv, "method", &reduceMethod);
char *typeChoice;
cutGetCmdLineArgumentstr( argc, (const char**) argv, "type", &typeChoice);
if (0 == typeChoice)
{
typeChoice = (char*)malloc(4 * sizeof(char));
strcpy(typeChoice, "int");
}
ReduceType datatype = REDUCE_INT;
if (!strcasecmp(typeChoice, "float"))
datatype = REDUCE_FLOAT;
else if (!strcasecmp(typeChoice, "double"))
datatype = REDUCE_DOUBLE;
else
datatype = REDUCE_INT;
cudaDeviceProp deviceProp;
deviceProp.major = 1;
deviceProp.minor = 0;
int minimumComputeVersion = 10;
if (datatype == REDUCE_DOUBLE)
{
deviceProp.minor = 3;
minimumComputeVersion = 13;
}
int dev;
if(!cutCheckCmdLineFlag(argc, (const char**)argv, "method") )
{
fprintf(stderr, "MISSING --method FLAG.\nYou must provide --method={ SUM | MIN | MAX }.\n");
exit(1);
}
if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") )
{
cutilDeviceInit(argc, argv);
cutilSafeCallNoSync(cudaGetDevice(&dev));
}
else
{
cutilSafeCallNoSync(cudaChooseDevice(&dev, &deviceProp));
}
cutilSafeCallNoSync(cudaGetDeviceProperties(&deviceProp, dev));
if((deviceProp.major * 10 + deviceProp.minor) >= minimumComputeVersion)
{
shrLog("Using Device %d: %s\n\n", dev, deviceProp.name);
cutilSafeCallNoSync(cudaSetDevice(dev));
}
else
{
shrLog("Error: the selected device does not support the minimum compute capability of %d.%d.\n\n",
minimumComputeVersion / 10, minimumComputeVersion % 10);
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_WAIVED);
}
shrLog("Reducing array of type %s\n\n", typeChoice);
bool bResult = false;
switch (datatype)
{
default:
case REDUCE_INT:
if (strcmp("SUM", reduceMethod) == 0) {
bResult = runTestSum<int>( argc, argv, datatype);
} else if ( strcmp("MAX", reduceMethod) == 0 ) {
bResult = runTestMax<int>( argc, argv, datatype);
} else if ( strcmp("MIN", reduceMethod) == 0 ) {
bResult = runTestMin<int>( argc, argv, datatype);
} else {
fprintf(stderr, "No --method specified!\n");
exit(1);
}
break;
case REDUCE_FLOAT:
if (strcmp("SUM", reduceMethod) == 0) {
bResult = runTestSum<float>( argc, argv, datatype);
} else if ( strcmp("MAX", reduceMethod) == 0 ) {
bResult = runTestMax<float>( argc, argv, datatype);
} else if ( strcmp("MIN", reduceMethod) == 0 ) {
bResult = runTestMin<float>( argc, argv, datatype);
} else {
fprintf(stderr, "No --method specified!\n");
exit(1);
}
break;
case REDUCE_DOUBLE:
if (strcmp("SUM", reduceMethod) == 0) {
bResult = runTestSum<double>( argc, argv, datatype);
} else if ( strcmp("MAX", reduceMethod) == 0 ) {
bResult = runTestMax<double>( argc, argv, datatype);
} else if ( strcmp("MIN", reduceMethod) == 0 ) {
bResult = runTestMin<double>( argc, argv, datatype);
} else {
fprintf(stderr, "No --method specified!\n");
exit(1);
}
break;
}
cutilDeviceReset();
shrQAFinishExit(argc, (const char**)argv, (bResult ? QA_PASSED : QA_FAILED));
}
void shmoo(int minN, int maxN, int maxThreads, int maxBlocks, ReduceType datatype)
{
fprintf(stderr, "Shmoo wasn't implemented in this modified kernel!\n");
exit(1);
// create random input data on CPU
unsigned int bytes = maxN * sizeof(T);
T *h_idata = (T*) malloc(bytes);
for(int i = 0; i < maxN; i++) {
// Keep the numbers small so we don't get truncation error in the sum
if (datatype == REDUCE_INT)
h_idata[i] = (T)(rand() & 0xFF);
else
h_idata[i] = (rand() & 0xFF) / (T)RAND_MAX;
}
int maxNumBlocks = MIN( maxN / maxThreads, MAX_BLOCK_DIM_SIZE);
// allocate mem for the result on host side
T* h_odata = (T*) malloc(maxNumBlocks*sizeof(T));
// allocate device memory and data
T* d_idata = NULL;
T* d_odata = NULL;
cutilSafeCallNoSync( cudaMalloc((void**) &d_idata, bytes) );
cutilSafeCallNoSync( cudaMalloc((void**) &d_odata, maxNumBlocks*sizeof(T)) );
// copy data directly to device memory
cutilSafeCallNoSync( cudaMemcpy(d_idata, h_idata, bytes, cudaMemcpyHostToDevice) );
cutilSafeCallNoSync( cudaMemcpy(d_odata, h_idata, maxNumBlocks*sizeof(T), cudaMemcpyHostToDevice) );
// warm-up
for (int kernel = 0; kernel < 7; kernel++)
{
sumreduce<T>(maxN, maxThreads, maxNumBlocks, kernel, d_idata, d_odata);
}
int testIterations = 100;
unsigned int timer = 0;
cutilCheckError( cutCreateTimer( &timer));
// print headers
shrLog("Time in milliseconds for various numbers of elements for each kernel\n\n\n");
shrLog("Kernel");
for (int i = minN; i <= maxN; i *= 2)
{
shrLog(", %d", i);
}
for (int kernel = 0; kernel < 7; kernel++)
{
shrLog("\n%d", kernel);
for (int i = minN; i <= maxN; i *= 2)
{
cutResetTimer(timer);
int numBlocks = 0;
int numThreads = 0;
getNumBlocksAndThreads(kernel, i, maxBlocks, maxThreads, numBlocks, numThreads);
float reduceTime;
if( numBlocks <= MAX_BLOCK_DIM_SIZE ) {
benchmarkReduceSum(i, numThreads, numBlocks, maxThreads, maxBlocks, kernel,
testIterations, false, 1, timer, h_odata, d_idata, d_odata);
reduceTime = cutGetAverageTimerValue(timer);
} else {
reduceTime = -1.0;
}
shrLog(", %.5f", reduceTime);
}
}
// cleanup
cutilCheckError(cutDeleteTimer(timer));
free(h_idata);
free(h_odata);
cutilSafeCallNoSync(cudaFree(d_idata));
cutilSafeCallNoSync(cudaFree(d_odata));
}
T benchmarkReduceMax(int n,
int numThreads,
int numBlocks,
int maxThreads,
int maxBlocks,
int whichKernel,
int testIterations,
bool cpuFinalReduction,
int cpuFinalThreshold,
unsigned int timer,
T* h_odata,
T* d_idata,
T* d_odata)
{
T gpu_result = 0;
bool needReadBack = true;
for (int i = 0; i < testIterations; ++i)
{
gpu_result = 0;
cutilDeviceSynchronize();
cutilCheckError( cutStartTimer( timer));
// execute the kernel
maxreduce<T>(n, numThreads, numBlocks, whichKernel, d_idata, d_odata);
// check if kernel execution generated an error
cutilCheckMsg("Kernel execution failed");
if (cpuFinalReduction)
{
// sum partial sums from each block on CPU
// copy result from device to host
cutilSafeCallNoSync( cudaMemcpy( h_odata, d_odata, numBlocks*sizeof(T), cudaMemcpyDeviceToHost) );
for(int i=0; i<numBlocks; i++)
{
gpu_result += h_odata[i];
}
needReadBack = false;
}
else
{
// sum partial block sums on GPU
int s=numBlocks;
int kernel = whichKernel;
while(s > cpuFinalThreshold)
{
int threads = 0, blocks = 0;
getNumBlocksAndThreads(kernel, s, maxBlocks, maxThreads, blocks, threads);
maxreduce<T>(s, threads, blocks, kernel, d_odata, d_odata);
if (kernel < 3)
s = (s + threads - 1) / threads;
else
s = (s + (threads*2-1)) / (threads*2);
}
if (s > 1)
{
// copy result from device to host
cutilSafeCallNoSync( cudaMemcpy( h_odata, d_odata, s * sizeof(T), cudaMemcpyDeviceToHost) );
for(int i=0; i < s; i++)
{
gpu_result += h_odata[i];
}
needReadBack = false;
}
}
cutilDeviceSynchronize();
cutilCheckError( cutStopTimer(timer) );
}
if (needReadBack)
{
// copy final sum from device to host
cutilSafeCallNoSync( cudaMemcpy( &gpu_result, d_odata, sizeof(T), cudaMemcpyDeviceToHost) );
}
return gpu_result;
}