int main(int argc, char** argv)
{
shrQAStart(argc, argv);
// set logfile name and start logs
shrSetLogFileName ("oclBandwidthTest.txt");
shrLog("%s Starting...\n\n", argv[0]);
// run the main test
int iRetVal = runTest(argc, (const char **)argv);
// finish
shrQAFinishExit(argc, (const char **)argv, (iRetVal == 0) ? QA_PASSED : QA_FAILED);
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char** argv)
{
shrQAStart(argc, argv);
// start the logs
shrSetLogFileName ("oclMatrixMul.txt");
shrLog("%s Starting...\n\n", argv[0]);
// run the code
bool bOK = (runTest(argc, (const char **)argv) == CL_SUCCESS);
shrLog("%s\n\n", (bOK ? "PASSED" : "FAILED"));
// finish
shrQAFinishExit(argc, (const char **)argv, (bOK ? QA_PASSED : QA_FAILED));
}
int main(int argc, char **argv)
{
bool bTestResult = true;
shrQAStart(argc, argv);
// Start the log
shrSetLogFileName(shrLogFile);
shrLog("%s Starting...\n\n", argv[0]);
// Check help flag
if (shrCheckCmdLineFlag(argc, (const char **)argv, "help")) {
shrLog("Displaying help on console\n");
showHelp(argc, (const char **)argv);
} else {
// Execute
bTestResult = runTest(argc, (const char **)argv);
oclCheckErrorEX(bTestResult, true, NULL);
}
// Finish
shrQAFinishExit( argc, (const char **)argv, (bTestResult ? QA_PASSED : QA_FAILED) );
}
int main(int argc, char **argv)
{
int M = 0, N = 0, nz = 0, *I = NULL, *J = NULL;
float *val = NULL;
const float tol = 1e-5f;
const int max_iter = 10000;
float *x;
float *rhs;
float a, b, na, r0, r1;
int *d_col, *d_row;
float *d_val, *d_x, dot;
float *d_r, *d_p, *d_Ax;
int k;
float alpha, beta, alpham1;
shrQAStart(argc, argv);
// This will pick the best possible CUDA capable device
cudaDeviceProp deviceProp;
int devID = findCudaDevice(argc, (const char **)argv);
if (devID < 0) {
printf("exiting...\n");
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
exit(0);
}
checkCudaErrors( cudaGetDeviceProperties(&deviceProp, devID) );
// Statistics about the GPU device
printf("> GPU device has %d Multi-Processors, SM %d.%d compute capabilities\n\n",
deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
int version = (deviceProp.major * 0x10 + deviceProp.minor);
if(version < 0x11)
{
printf("%s: requires a minimum CUDA compute 1.1 capability\n", sSDKname);
cudaDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
/* Generate a random tridiagonal symmetric matrix in CSR format */
M = N = 1048576;
nz = (N-2)*3 + 4;
I = (int*)malloc(sizeof(int)*(N+1));
J = (int*)malloc(sizeof(int)*nz);
val = (float*)malloc(sizeof(float)*nz);
genTridiag(I, J, val, N, nz);
x = (float*)malloc(sizeof(float)*N);
rhs = (float*)malloc(sizeof(float)*N);
for (int i = 0; i < N; i++) {
rhs[i] = 1.0;
x[i] = 0.0;
}
/* Get handle to the CUBLAS context */
cublasHandle_t cublasHandle = 0;
cublasStatus_t cublasStatus;
cublasStatus = cublasCreate(&cublasHandle);
if ( checkCublasStatus (cublasStatus, "!!!! CUBLAS initialization error\n") ) return EXIT_FAILURE;
/* Get handle to the CUSPARSE context */
cusparseHandle_t cusparseHandle = 0;
cusparseStatus_t cusparseStatus;
cusparseStatus = cusparseCreate(&cusparseHandle);
if ( checkCusparseStatus (cusparseStatus, "!!!! CUSPARSE initialization error\n") ) return EXIT_FAILURE;
cusparseMatDescr_t descr = 0;
cusparseStatus = cusparseCreateMatDescr(&descr);
if ( checkCusparseStatus (cusparseStatus, "!!!! CUSPARSE cusparseCreateMatDescr error\n") ) return EXIT_FAILURE;
cusparseSetMatType(descr,CUSPARSE_MATRIX_TYPE_GENERAL);
cusparseSetMatIndexBase(descr,CUSPARSE_INDEX_BASE_ZERO);
checkCudaErrors( cudaMalloc((void**)&d_col, nz*sizeof(int)) );
checkCudaErrors( cudaMalloc((void**)&d_row, (N+1)*sizeof(int)) );
checkCudaErrors( cudaMalloc((void**)&d_val, nz*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_x, N*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_r, N*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_p, N*sizeof(float)) );
checkCudaErrors( cudaMalloc((void**)&d_Ax, N*sizeof(float)) );
cudaMemcpy(d_col, J, nz*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_row, I, (N+1)*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_val, val, nz*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_r, rhs, N*sizeof(float), cudaMemcpyHostToDevice);
alpha = 1.0;
alpham1 = -1.0;
beta = 0.0;
r0 = 0.;
cusparseScsrmv(cusparseHandle,CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, d_val, d_row, d_col, d_x, &beta, d_Ax);
cublasSaxpy(cublasHandle, N, &alpham1, d_Ax, 1, d_r, 1);
cublasStatus = cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
k = 1;
while (r1 > tol*tol && k <= max_iter) {
if (k > 1) {
b = r1 / r0;
cublasStatus = cublasSscal(cublasHandle, N, &b, d_p, 1);
cublasStatus = cublasSaxpy(cublasHandle, N, &alpha, d_r, 1, d_p, 1);
} else {
cublasStatus = cublasScopy(cublasHandle, N, d_r, 1, d_p, 1);
}
cusparseScsrmv(cusparseHandle, CUSPARSE_OPERATION_NON_TRANSPOSE, N, N, nz, &alpha, descr, d_val, d_row, d_col, d_p, &beta, d_Ax);
cublasStatus = cublasSdot(cublasHandle, N, d_p, 1, d_Ax, 1, &dot);
a = r1 / dot;
cublasStatus = cublasSaxpy(cublasHandle, N, &a, d_p, 1, d_x, 1);
na = -a;
cublasStatus = cublasSaxpy(cublasHandle, N, &na, d_Ax, 1, d_r, 1);
r0 = r1;
cublasStatus = cublasSdot(cublasHandle, N, d_r, 1, d_r, 1, &r1);
cudaThreadSynchronize();
printf("iteration = %3d, residual = %e\n", k, sqrt(r1));
k++;
}
cudaMemcpy(x, d_x, N*sizeof(float), cudaMemcpyDeviceToHost);
float rsum, diff, err = 0.0;
for (int i = 0; i < N; i++) {
rsum = 0.0;
for (int j = I[i]; j < I[i+1]; j++) {
rsum += val[j]*x[J[j]];
}
diff = fabs(rsum - rhs[i]);
if (diff > err) err = diff;
}
cusparseDestroy(cusparseHandle);
cublasDestroy(cublasHandle);
free(I);
free(J);
free(val);
free(x);
free(rhs);
cudaFree(d_col);
cudaFree(d_row);
cudaFree(d_val);
cudaFree(d_x);
cudaFree(d_r);
cudaFree(d_p);
cudaFree(d_Ax);
cudaDeviceReset();
printf("Test Summary: Error amount = %f\n", err);
shrQAFinishExit(argc, (const char **)argv, (k <= max_iter) ? QA_PASSED : QA_FAILED );
}
// Main program
//*****************************************************************************
int main(int argc, char** argv)
{
// Locals used with command line args
int p = 256; // workgroup X dimension
int q = 1; // workgroup Y dimension
pArgc = &argc;
pArgv = argv;
shrQAStart(argc, argv);
// latch the executable path for other funcs to use
cExecutablePath = argv[0];
// start logs and show command line help
shrSetLogFileName ("oclNbody.txt");
shrLog("%s Starting...\n\n", cExecutablePath);
shrLog("Command line switches:\n");
shrLog(" --qatest\t\tCheck correctness of GPU execution and measure performance)\n");
shrLog(" --noprompt\t\tQuit simulation automatically after a brief period\n");
shrLog(" --n=<numbodies>\tSpecify # of bodies to simulate (default = %d)\n", numBodies);
shrLog(" --double\t\tUse double precision floating point values for simulation\n");
shrLog(" --p=<workgroup X dim>\tSpecify X dimension of workgroup (default = %d)\n", p);
shrLog(" --q=<workgroup Y dim>\tSpecify Y dimension of workgroup (default = %d)\n\n", q);
// Get command line arguments if there are any and set vars accordingly
if (argc > 0)
{
shrGetCmdLineArgumenti(argc, (const char**)argv, "p", &p);
shrGetCmdLineArgumenti(argc, (const char**)argv, "q", &q);
shrGetCmdLineArgumenti(argc, (const char**)argv, "n", &numBodies);
bDouble = (shrTRUE == shrCheckCmdLineFlag(argc, (const char**)argv, "double"));
bNoPrompt = shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
bQATest = shrCheckCmdLineFlag(argc, (const char**)argv, "qatest");
}
//Get the NVIDIA platform
cl_int ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("clGetPlatformID...\n\n");
if (bDouble)
{
shrLog("Double precision execution...\n\n");
}
else
{
shrLog("Single precision execution...\n\n");
}
flopsPerInteraction = bDouble ? 30 : 20;
//Get all the devices
shrLog("Get the Device info and select Device...\n");
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 0, NULL, &uiNumDevices);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cdDevices = (cl_device_id *)malloc(uiNumDevices * sizeof(cl_device_id) );
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, uiNumDevices, cdDevices, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Set target device and Query number of compute units on uiTargetDevice
shrLog(" # of Devices Available = %u\n", uiNumDevices);
if(shrGetCmdLineArgumentu(argc, (const char**)argv, "device", &uiTargetDevice)== shrTRUE)
{
uiTargetDevice = CLAMP(uiTargetDevice, 0, (uiNumDevices - 1));
}
shrLog(" Using Device %u, ", uiTargetDevice);
oclPrintDevName(LOGBOTH, cdDevices[uiTargetDevice]);
cl_uint uiNumComputeUnits;
clGetDeviceInfo(cdDevices[uiTargetDevice], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(uiNumComputeUnits), &uiNumComputeUnits, NULL);
shrLog(" # of Compute Units = %u\n", uiNumComputeUnits);
//Create the context
shrLog("clCreateContext...\n");
cxContext = clCreateContext(0, uiNumDevsUsed, &cdDevices[uiTargetDevice], NULL, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Create a command-queue
shrLog("clCreateCommandQueue...\n\n");
cqCommandQueue = clCreateCommandQueue(cxContext, cdDevices[uiTargetDevice], CL_QUEUE_PROFILING_ENABLE, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Log and config for number of bodies
shrLog("Number of Bodies = %d\n", numBodies);
switch (numBodies)
{
case 1024:
activeParams.m_clusterScale = 1.52f;
activeParams.m_velocityScale = 2.f;
break;
case 2048:
activeParams.m_clusterScale = 1.56f;
activeParams.m_velocityScale = 2.64f;
break;
case 4096:
activeParams.m_clusterScale = 1.68f;
activeParams.m_velocityScale = 2.98f;
break;
case 7680:
case 8192:
activeParams.m_clusterScale = 1.98f;
activeParams.m_velocityScale = 2.9f;
break;
default:
case 15360:
case 16384:
activeParams.m_clusterScale = 1.54f;
activeParams.m_velocityScale = 8.f;
break;
case 30720:
case 32768:
activeParams.m_clusterScale = 1.44f;
activeParams.m_velocityScale = 11.f;
break;
}
if ((q * p) > 256)
{
p = 256 / q;
shrLog("Setting p=%d to maintain %d threads per block\n", p, 256);
}
if ((q == 1) && (numBodies < p))
{
p = numBodies;
shrLog("Setting p=%d because # of bodies < p\n", p);
}
shrLog("Workgroup Dims = (%d x %d)\n\n", p, q);
// Initialize OpenGL items if using GL
if (bQATest == shrFALSE)
{
assert(0);
/*
shrLog("Calling InitGL...\n");
InitGL(&argc, argv);
*/
}
else
{
shrLog("Skipping InitGL...\n");
}
// CL/GL interop disabled
bUsePBO = (false && (bQATest == shrFALSE));
InitNbody(cdDevices[uiTargetDevice], cxContext, cqCommandQueue, numBodies, p, q, bUsePBO, bDouble);
ResetSim(nbody, numBodies, NBODY_CONFIG_SHELL, bUsePBO);
// init timers
shrDeltaT(DEMOTIME); // timer 0 is for timing demo periods
shrDeltaT(FUNCTIME); // timer 1 is for logging function delta t's
shrDeltaT(FPSTIME); // timer 2 is for fps measurement
// Standard simulation
if (bQATest == shrFALSE)
{
assert(0);
/*
shrLog("Running standard oclNbody simulation...\n\n");
glutDisplayFunc(DisplayGL);
glutReshapeFunc(ReshapeGL);
glutMouseFunc(MouseGL);
glutMotionFunc(MotionGL);
glutKeyboardFunc(KeyboardGL);
glutSpecialFunc(SpecialGL);
glutIdleFunc(IdleGL);
glutMainLoop();
*/
}
// Compare to host, profile and write out file for regression analysis
if (bQATest == shrTRUE) {
bool bTestResults = false;
shrLog("Running oclNbody Results Comparison...\n\n");
bTestResults = CompareResults(numBodies);
//shrLog("Profiling oclNbody...\n\n");
//RunProfiling(100, (unsigned int)(p * q)); // 100 iterations
shrQAFinish(argc, (const char **)argv, bTestResults ? QA_PASSED : QA_FAILED);
} else {
// Cleanup/exit
bNoPrompt = shrTRUE;
shrQAFinish2(false, *pArgc, (const char **)pArgv, QA_PASSED);
}
Cleanup(EXIT_SUCCESS);
}
// Main function
// *********************************************************************
int main(int argc, char **argv)
{
gp_argc = &argc;
gp_argv = &argv;
shrQAStart(argc, argv);
// Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("clGetPlatformID...\n");
// Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("clGetPlatformID...\n");
//Get all the devices
cl_uint uiNumDevices = 0; // Number of devices available
cl_uint uiTargetDevice = 0; // Default Device to compute on
cl_uint uiNumComputeUnits; // Number of compute units (SM's on NV GPU)
shrLog("Get the Device info and select Device...\n");
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 0, NULL, &uiNumDevices);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
cdDevices = (cl_device_id *)malloc(uiNumDevices * sizeof(cl_device_id) );
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, uiNumDevices, cdDevices, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
// Get command line device options and config accordingly
shrLog(" # of Devices Available = %u\n", uiNumDevices);
if(shrGetCmdLineArgumentu(argc, (const char**)argv, "device", &uiTargetDevice)== shrTRUE)
{
uiTargetDevice = CLAMP(uiTargetDevice, 0, (uiNumDevices - 1));
}
shrLog(" Using Device %u: ", uiTargetDevice);
oclPrintDevName(LOGBOTH, cdDevices[uiTargetDevice]);
ciErrNum = clGetDeviceInfo(cdDevices[uiTargetDevice], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(uiNumComputeUnits), &uiNumComputeUnits, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("\n # of Compute Units = %u\n", uiNumComputeUnits);
// get command line arg for quick test, if provided
bNoPrompt = shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
// start logs
cExecutableName = argv[0];
shrSetLogFileName ("oclDotProduct.txt");
shrLog("%s Starting...\n\n# of float elements per Array \t= %u\n", argv[0], iNumElements);
// set and log Global and Local work size dimensions
szLocalWorkSize = 256;
szGlobalWorkSize = shrRoundUp((int)szLocalWorkSize, iNumElements); // rounded up to the nearest multiple of the LocalWorkSize
shrLog("Global Work Size \t\t= %u\nLocal Work Size \t\t= %u\n# of Work Groups \t\t= %u\n\n",
szGlobalWorkSize, szLocalWorkSize, (szGlobalWorkSize % szLocalWorkSize + szGlobalWorkSize/szLocalWorkSize));
// Allocate and initialize host arrays
shrLog( "Allocate and Init Host Mem...\n");
srcA = (void *)malloc(sizeof(cl_float4) * szGlobalWorkSize);
srcB = (void *)malloc(sizeof(cl_float4) * szGlobalWorkSize);
dst = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
Golden = (void *)malloc(sizeof(cl_float) * iNumElements);
shrFillArray((float*)srcA, 4 * iNumElements);
shrFillArray((float*)srcB, 4 * iNumElements);
// Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Get a GPU device
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevices[uiTargetDevice], NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Create the context
cxGPUContext = clCreateContext(0, 1, &cdDevices[uiTargetDevice], NULL, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Create a command-queue
shrLog("clCreateCommandQueue...\n");
cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevices[uiTargetDevice], 0, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Allocate the OpenCL buffer memory objects for source and result on the device GMEM
shrLog("clCreateBuffer (SrcA, SrcB and Dst in Device GMEM)...\n");
cmDevSrcA = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize * 4, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cmDevSrcB = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize * 4, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cmDevDst = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Read the OpenCL kernel in from source file
shrLog("oclLoadProgSource (%s)...\n", cSourceFile);
cPathAndName = shrFindFilePath(cSourceFile, argv[0]);
oclCheckErrorEX(cPathAndName != NULL, shrTRUE, pCleanup);
cSourceCL = oclLoadProgSource(cPathAndName, "", &szKernelLength);
oclCheckErrorEX(cSourceCL != NULL, shrTRUE, pCleanup);
// Create the program
shrLog("clCreateProgramWithSource...\n");
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErrNum);
// Build the program with 'mad' Optimization option
#ifdef MAC
char* flags = "-cl-fast-relaxed-math -DMAC";
#else
char* flags = "-cl-fast-relaxed-math";
#endif
shrLog("clBuildProgram...\n");
ciErrNum = clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);
if (ciErrNum != CL_SUCCESS)
{
// write out standard error, Build Log and PTX, then cleanup and exit
shrLogEx(LOGBOTH | ERRORMSG, ciErrNum, STDERROR);
oclLogBuildInfo(cpProgram, oclGetFirstDev(cxGPUContext));
oclLogPtx(cpProgram, oclGetFirstDev(cxGPUContext), "oclDotProduct.ptx");
Cleanup(EXIT_FAILURE);
}
// Create the kernel
shrLog("clCreateKernel (DotProduct)...\n");
ckKernel = clCreateKernel(cpProgram, "DotProduct", &ciErrNum);
// Set the Argument values
shrLog("clSetKernelArg 0 - 3...\n\n");
ciErrNum = clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevSrcA);
ciErrNum |= clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void*)&cmDevSrcB);
ciErrNum |= clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void*)&cmDevDst);
ciErrNum |= clSetKernelArg(ckKernel, 3, sizeof(cl_int), (void*)&iNumElements);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// --------------------------------------------------------
// Core sequence... copy input data to GPU, compute, copy results back
// Asynchronous write of data to GPU device
shrLog("clEnqueueWriteBuffer (SrcA and SrcB)...\n");
ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cmDevSrcA, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize * 4, srcA, 0, NULL, NULL);
ciErrNum |= clEnqueueWriteBuffer(cqCommandQueue, cmDevSrcB, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize * 4, srcB, 0, NULL, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Launch kernel
shrLog("clEnqueueNDRangeKernel (DotProduct)...\n");
ciErrNum = clEnqueueNDRangeKernel(cqCommandQueue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Read back results and check accumulated errors
shrLog("clEnqueueReadBuffer (Dst)...\n\n");
ciErrNum = clEnqueueReadBuffer(cqCommandQueue, cmDevDst, CL_TRUE, 0, sizeof(cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Compute and compare results for golden-host and report errors and pass/fail
shrLog("Comparing against Host/C++ computation...\n\n");
DotProductHost ((const float*)srcA, (const float*)srcB, (float*)Golden, iNumElements);
shrBOOL bMatch = shrComparefet((const float*)Golden, (const float*)dst, (unsigned int)iNumElements, 0.0f, 0);
// Cleanup and leave
Cleanup (EXIT_SUCCESS);
}
// Main function
// *********************************************************************
int ymain(int argc, char **argv)
{
shrQAStart(argc, argv);
// get command line arg for quick test, if provided
bNoPrompt = shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
// start logs
cExecutableName = argv[0];
shrSetLogFileName ("oclVectorAdd2.txt");
shrLog("%s Starting...\n\n# of float elements per Array \t= %i\n", argv[0], iNumElements);
// set and log Global and Local work size dimensions
szLocalWorkSize = 256;
szGlobalWorkSize = shrRoundUp((int)szLocalWorkSize, iNumElements); // rounded up to the nearest multiple of the LocalWorkSize
shrLog("Global Work Size \t\t= %u\nLocal Work Size \t\t= %u\n# of Work Groups \t\t= %u\n\n",
szGlobalWorkSize, szLocalWorkSize, (szGlobalWorkSize % szLocalWorkSize + szGlobalWorkSize/szLocalWorkSize));
// Allocate and initialize host arrays
shrLog( "Allocate and Init Host Mem...\n");
srcA = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
srcB = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
dst = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
Golden = (void *)malloc(sizeof(cl_float) * iNumElements);
shrFillArray((float*)srcA, iNumElements);
shrFillArray((float*)srcB, iNumElements);
//Get an OpenCL platform
ciErr1 = clGetPlatformIDs(1, &cpPlatform, NULL);
shrLog("clGetPlatformID...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clGetPlatformID, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
//Get the devices
ciErr1 = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL);
shrLog("clGetDeviceIDs...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clGetDeviceIDs, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
//Create the context
cxGPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErr1);
shrLog("clCreateContext...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateContext, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Create a command-queue
cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevice, 0, &ciErr1);
shrLog("clCreateCommandQueue...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateCommandQueue, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Allocate the OpenCL buffer memory objects for source and result on the device GMEM
cmDevSrcA = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr1);
cmDevSrcB = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr2);
ciErr1 |= ciErr2;
cmDevDst = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr2);
ciErr1 |= ciErr2;
shrLog("clCreateBuffer...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Read the OpenCL kernel in from source file
shrLog("oclLoadProgSource (%s)...\n", cSourceFile);
cPathAndName = shrFindFilePath(cSourceFile, argv[0]);
shrLog("Looking for: %s in Path: %s\n", cSourceFile, argv[0]);
cSourceCL = oclLoadProgSource(cPathAndName, "", &szKernelLength);
// Create the program
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErr1);
shrLog("clCreateProgramWithSource...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateProgramWithSource, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Build the program with 'mad' Optimization option
#ifdef MAC
char* flags = "-cl-fast-relaxed-math -DMAC";
#else
char* flags = "-cl-fast-relaxed-math";
#endif
ciErr1 = clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);
shrLog("clBuildProgram...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clBuildProgram, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Create the kernel
ckKernel = clCreateKernel(cpProgram, "VectorAdd", &ciErr1);
shrLog("clCreateKernel (VectorAdd)...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateKernel, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Set the Argument values
ciErr1 = clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevSrcA);
ciErr1 |= clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void*)&cmDevSrcB);
ciErr1 |= clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void*)&cmDevDst);
ciErr1 |= clSetKernelArg(ckKernel, 3, sizeof(cl_int), (void*)&iNumElements);
shrLog("clSetKernelArg 0 - 3...\n\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clSetKernelArg, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// --------------------------------------------------------
// Start Core sequence... copy input data to GPU, compute, copy results back
// Asynchronous write of data to GPU device
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, cmDevSrcA, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize, srcA, 0, NULL, NULL);
ciErr1 |= clEnqueueWriteBuffer(cqCommandQueue, cmDevSrcB, CL_FALSE, 0, sizeof(cl_float) * szGlobalWorkSize, srcB, 0, NULL, NULL);
shrLog("clEnqueueWriteBuffer (SrcA and SrcB)...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clEnqueueWriteBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Launch kernel
ciErr1 = clEnqueueNDRangeKernel(cqCommandQueue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);
shrLog("clEnqueueNDRangeKernel (VectorAdd)...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clEnqueueNDRangeKernel, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Synchronous/blocking read of results, and check accumulated errors
ciErr1 = clEnqueueReadBuffer(cqCommandQueue, cmDevDst, CL_TRUE, 0, sizeof(cl_float) * szGlobalWorkSize, dst, 0, NULL, NULL);
shrLog("clEnqueueReadBuffer (Dst)...\n\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clEnqueueReadBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
//--------------------------------------------------------
// Compute and compare results for golden-host and report errors and pass/fail
shrLog("Comparing against Host/C++ computation...\n\n");
VectorAddHost ((const float*)srcA, (const float*)srcB, (float*)Golden, iNumElements);
shrBOOL bMatch = shrComparefet((const float*)Golden, (const float*)dst, (unsigned int)iNumElements, 0.0f, 0);
// Cleanup and leave
Cleanup (argc, argv, (bMatch == shrTRUE) ? EXIT_SUCCESS : EXIT_FAILURE);
}
// Main function
// *********************************************************************
int main(int argc, char** argv)
{
shrQAStart(argc, argv);
// get command line arg for quick test, if provided
bNoPrompt = shrCheckCmdLineFlag(argc, (const char **)argv, "noprompt");
// start logs
cExecutableName = argv[0];
shrSetLogFileName ("oclMatVecMul.txt");
shrLog("%s Starting...\n\n", argv[0]);
// calculate matrix height given GPU memory
shrLog("Determining Matrix height from available GPU mem...\n");
memsize_t memsize;
getTargetDeviceGlobalMemSize(&memsize, argc, (const char **)argv);
height = memsize/width/16;
if (height > MAX_HEIGHT)
height = MAX_HEIGHT;
shrLog(" Matrix width\t= %u\n Matrix height\t= %u\n\n", width, height);
// Allocate and initialize host arrays
shrLog("Allocate and Init Host Mem...\n\n");
unsigned int size = width * height;
unsigned int mem_size_M = size * sizeof(float);
M = (float*)malloc(mem_size_M);
unsigned int mem_size_V = width * sizeof(float);
V = (float*)malloc(mem_size_V);
unsigned int mem_size_W = height * sizeof(float);
W = (float*)malloc(mem_size_W);
shrFillArray(M, size);
shrFillArray(V, width);
Golden = (float*)malloc(mem_size_W);
MatVecMulHost(M, V, width, height, Golden);
//Get the NVIDIA platform
shrLog("Get the Platform ID...\n\n");
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
//Get all the devices
shrLog("Get the Device info and select Device...\n");
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 0, NULL, &uiNumDevices);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cdDevices = (cl_device_id *)malloc(uiNumDevices * sizeof(cl_device_id) );
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, uiNumDevices, cdDevices, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Set target device and Query number of compute units on targetDevice
shrLog(" # of Devices Available = %u\n", uiNumDevices);
if(shrGetCmdLineArgumentu(argc, (const char **)argv, "device", &targetDevice)== shrTRUE)
{
targetDevice = CLAMP(targetDevice, 0, (uiNumDevices - 1));
}
shrLog(" Using Device %u: ", targetDevice);
oclPrintDevName(LOGBOTH, cdDevices[targetDevice]);
cl_uint num_compute_units;
clGetDeviceInfo(cdDevices[targetDevice], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(num_compute_units), &num_compute_units, NULL);
shrLog("\n # of Compute Units = %u\n\n", num_compute_units);
//Create the context
shrLog("clCreateContext...\n");
cxGPUContext = clCreateContext(0, uiNumDevsUsed, &cdDevices[targetDevice], NULL, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Create a command-queue
shrLog("clCreateCommandQueue...\n");
cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevices[targetDevice], CL_QUEUE_PROFILING_ENABLE, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Allocate the OpenCL buffer memory objects for source and result on the device GMEM
shrLog("clCreateBuffer (M, V and W in device global memory, mem_size_m = %u)...\n", mem_size_M);
cmM = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, mem_size_M, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cmV = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, mem_size_V, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cmW = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, mem_size_W, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Read the OpenCL kernel in from source file
shrLog("oclLoadProgSource (%s)...\n", cSourceFile);
cPathAndName = shrFindFilePath(cSourceFile, argv[0]);
oclCheckErrorEX(cPathAndName != NULL, shrTRUE, pCleanup);
cSourceCL = oclLoadProgSource(cPathAndName, "", &szKernelLength);
oclCheckErrorEX(cSourceCL != NULL, shrTRUE, pCleanup);
// Create the program
shrLog("clCreateProgramWithSource...\n");
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErrNum);
// Build the program
shrLog("clBuildProgram...\n");
ciErrNum = clBuildProgram(cpProgram, uiNumDevsUsed, &cdDevices[targetDevice], "-cl-fast-relaxed-math", NULL, NULL);
if (ciErrNum != CL_SUCCESS)
{
// write out standard error, Build Log and PTX, then cleanup and exit
shrLogEx(LOGBOTH | ERRORMSG, ciErrNum, STDERROR);
oclLogBuildInfo(cpProgram, oclGetFirstDev(cxGPUContext));
oclLogPtx(cpProgram, oclGetFirstDev(cxGPUContext), "oclMatVecMul.ptx");
shrQAFinish(argc, (const char **)argv, QA_FAILED);
Cleanup(EXIT_FAILURE);
}
// --------------------------------------------------------
// Core sequence... copy input data to GPU, compute, copy results back
// Asynchronous write of data to GPU device
shrLog("clEnqueueWriteBuffer (M and V)...\n\n");
ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cmM, CL_FALSE, 0, mem_size_M, M, 0, NULL, NULL);
ciErrNum |= clEnqueueWriteBuffer(cqCommandQueue, cmV, CL_FALSE, 0, mem_size_V, V, 0, NULL, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Kernels
const char* kernels[] = {
"MatVecMulUncoalesced0",
"MatVecMulUncoalesced1",
"MatVecMulCoalesced0",
"MatVecMulCoalesced1",
"MatVecMulCoalesced2",
"MatVecMulCoalesced3" };
for (int k = 0; k < (int)(sizeof(kernels)/sizeof(char*)); ++k) {
shrLog("Running with Kernel %s...\n\n", kernels[k]);
// Clear result
shrLog(" Clear result with clEnqueueWriteBuffer (W)...\n");
memset(W, 0, mem_size_W);
ciErrNum = clEnqueueWriteBuffer(cqCommandQueue, cmW, CL_FALSE, 0, mem_size_W, W, 0, NULL, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Create the kernel
shrLog(" clCreateKernel...\n");
if (ckKernel) {
clReleaseKernel(ckKernel);
ckKernel = 0;
}
ckKernel = clCreateKernel(cpProgram, kernels[k], &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Set and log Global and Local work size dimensions
szLocalWorkSize = 256;
if (k == 0)
szGlobalWorkSize = shrRoundUp((int)szLocalWorkSize, height); // rounded up to the nearest multiple of the LocalWorkSize
else
// Some experiments should be done here for determining the best global work size for a given device
// We will assume here that we can run 2 work-groups per compute unit
szGlobalWorkSize = 2 * num_compute_units * szLocalWorkSize;
shrLog(" Global Work Size \t\t= %u\n Local Work Size \t\t= %u\n # of Work Groups \t\t= %u\n",
szGlobalWorkSize, szLocalWorkSize, (szGlobalWorkSize % szLocalWorkSize + szGlobalWorkSize/szLocalWorkSize));
// Set the Argument values
shrLog(" clSetKernelArg...\n\n");
int n = 0;
ciErrNum = clSetKernelArg(ckKernel, n++, sizeof(cl_mem), (void*)&cmM);
ciErrNum |= clSetKernelArg(ckKernel, n++, sizeof(cl_mem), (void*)&cmV);
ciErrNum |= clSetKernelArg(ckKernel, n++, sizeof(cl_int), (void*)&width);
ciErrNum |= clSetKernelArg(ckKernel, n++, sizeof(cl_int), (void*)&height);
ciErrNum |= clSetKernelArg(ckKernel, n++, sizeof(cl_mem), (void*)&cmW);
if (k > 1)
ciErrNum |= clSetKernelArg(ckKernel, n++, szLocalWorkSize * sizeof(float), 0);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Launch kernel
shrLog(" clEnqueueNDRangeKernel (%s)...\n", kernels[k]);
ciErrNum = clEnqueueNDRangeKernel(cqCommandQueue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, &ceEvent);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Read back results and check accumulated errors
shrLog(" clEnqueueReadBuffer (W)...\n");
ciErrNum = clEnqueueReadBuffer(cqCommandQueue, cmW, CL_TRUE, 0, mem_size_W, W, 0, NULL, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
#ifdef GPU_PROFILING
// Execution time
ciErrNum = clWaitForEvents(1, &ceEvent);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cl_ulong start, end;
ciErrNum = clGetEventProfilingInfo(ceEvent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
ciErrNum |= clGetEventProfilingInfo(ceEvent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
double dSeconds = 1.0e-9 * (double)(end - start);
shrLog(" Kernel execution time: %.5f s\n\n", dSeconds);
#endif
// Compare results for golden-host and report errors and pass/fail
shrLog(" Comparing against Host/C++ computation...\n\n");
shrBOOL res = shrCompareL2fe(Golden, W, height, 1e-6f);
shrLog(" GPU Result %s CPU Result within allowable tolerance\n\n", (res == shrTRUE) ? "MATCHES" : "DOESN'T MATCH");
bPassFlag &= (res == shrTRUE);
// Release event
ciErrNum = clReleaseEvent(ceEvent);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
ceEvent = 0;
}
// Master status Pass/Fail (all tests)
shrQAFinish(argc, (const char **)argv, (bPassFlag ? QA_PASSED : QA_FAILED) );
// Cleanup and leave
Cleanup (EXIT_SUCCESS);
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int
main( int argc, char** argv)
{
CUdevice dev;
int major = 0, minor = 0;
int deviceCount = 0;
char deviceName[256];
shrQAStart(argc, argv);
// note your project will need to link with cuda.lib files on windows
printf("CUDA Device Query (Driver API) statically linked version \n");
CUresult error_id = cuInit(0);
if (error_id != CUDA_SUCCESS) {
printf("cuInit(0) returned %d\n-> %s\n", error_id, getCudaDrvErrorString(error_id));
shrQAFinishExit(argc, (const char **)argv, QA_FAILED);
}
error_id = cuDeviceGetCount(&deviceCount);
if (error_id != CUDA_SUCCESS) {
shrLog( "cuDeviceGetCount returned %d\n-> %s\n", (int)error_id, getCudaDrvErrorString(error_id) );
shrQAFinishExit(argc, (const char **)argv, QA_FAILED);
}
// This function call returns 0 if there are no CUDA capable devices.
if (deviceCount == 0)
printf("There are no available device(s) that support CUDA\n");
else if (deviceCount == 1)
printf("There is 1 device supporting CUDA\n");
else
printf("There are %d devices supporting CUDA\n", deviceCount);
for (dev = 0; dev < deviceCount; ++dev) {
error_id = cuDeviceComputeCapability(&major, &minor, dev);
if (error_id != CUDA_SUCCESS) {
shrLog( "cuDeviceComputeCapability returned %d\n-> %s\n", (int)error_id, getCudaDrvErrorString(error_id) );
shrQAFinishExit(argc, (const char **)argv, QA_FAILED);
}
error_id = cuDeviceGetName(deviceName, 256, dev);
if (error_id != CUDA_SUCCESS) {
shrLog( "cuDeviceGetName returned %d\n-> %s\n", (int)error_id, getCudaDrvErrorString(error_id) );
shrQAFinishExit(argc, (const char **)argv, QA_FAILED);
}
printf("\nDevice %d: \"%s\"\n", dev, deviceName);
#if CUDA_VERSION >= 2020
int driverVersion = 0;
cuDriverGetVersion(&driverVersion);
printf(" CUDA Driver Version: %d.%d\n", driverVersion/1000, (driverVersion%100)/10);
#endif
shrLog(" CUDA Capability Major/Minor version number: %d.%d\n", major, minor);
size_t totalGlobalMem;
error_id = cuDeviceTotalMem(&totalGlobalMem, dev);
if (error_id != CUDA_SUCCESS) {
shrLog( "cuDeviceTotalMem returned %d\n-> %s\n", (int)error_id, getCudaDrvErrorString(error_id) );
shrQAFinishExit(argc, (const char **)argv, QA_FAILED);
}
char msg[256];
sprintf(msg, " Total amount of global memory: %.0f MBytes (%llu bytes)\n",
(float)totalGlobalMem/1048576.0f, (unsigned long long) totalGlobalMem);
shrLog(msg);
#if CUDA_VERSION >= 2000
int multiProcessorCount;
getCudaAttribute<int>(&multiProcessorCount, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, dev);
shrLog(" (%2d) Multiprocessors x (%3d) CUDA Cores/MP: %d CUDA Cores\n",
multiProcessorCount, ConvertSMVer2Cores(major, minor),
ConvertSMVer2Cores(major, minor) * multiProcessorCount);
#endif
int clockRate;
getCudaAttribute<int>(&clockRate, CU_DEVICE_ATTRIBUTE_CLOCK_RATE, dev);
printf(" GPU Clock rate: %.0f MHz (%0.2f GHz)\n", clockRate * 1e-3f, clockRate * 1e-6f);
#if CUDA_VERSION >= 4000
int memoryClock;
getCudaAttribute<int>( &memoryClock, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, dev );
shrLog(" Memory Clock rate: %.0f Mhz\n", memoryClock * 1e-3f);
int memBusWidth;
getCudaAttribute<int>( &memBusWidth, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, dev );
shrLog(" Memory Bus Width: %d-bit\n", memBusWidth);
int L2CacheSize;
getCudaAttribute<int>( &L2CacheSize, CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE, dev );
if (L2CacheSize) {
shrLog(" L2 Cache Size: %d bytes\n", L2CacheSize);
}
int maxTex1D, maxTex2D[2], maxTex3D[3];
getCudaAttribute<int>( &maxTex1D, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_WIDTH, dev );
getCudaAttribute<int>( &maxTex2D[0], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_WIDTH, dev );
getCudaAttribute<int>( &maxTex2D[1], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_HEIGHT, dev );
getCudaAttribute<int>( &maxTex3D[0], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_WIDTH, dev );
getCudaAttribute<int>( &maxTex3D[1], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_HEIGHT, dev );
getCudaAttribute<int>( &maxTex3D[2], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE3D_DEPTH, dev );
shrLog(" Max Texture Dimension Sizes 1D=(%d) 2D=(%d,%d) 3D=(%d,%d,%d)\n",
maxTex1D, maxTex2D[0], maxTex2D[1], maxTex3D[0], maxTex3D[1], maxTex3D[2]);
int maxTex2DLayered[3];
getCudaAttribute<int>( &maxTex2DLayered[0], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH, dev );
getCudaAttribute<int>( &maxTex2DLayered[1], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT, dev );
getCudaAttribute<int>( &maxTex2DLayered[2], CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS, dev );
shrLog(" Max Layered Texture Size (dim) x layers 1D=(%d) x %d, 2D=(%d,%d) x %d\n",
maxTex2DLayered[0], maxTex2DLayered[2],
maxTex2DLayered[0], maxTex2DLayered[1], maxTex2DLayered[2]);
#endif
int totalConstantMemory;
getCudaAttribute<int>( &totalConstantMemory, CU_DEVICE_ATTRIBUTE_TOTAL_CONSTANT_MEMORY, dev );
printf(" Total amount of constant memory: %u bytes\n", totalConstantMemory);
int sharedMemPerBlock;
getCudaAttribute<int>( &sharedMemPerBlock, CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK, dev );
printf(" Total amount of shared memory per block: %u bytes\n", sharedMemPerBlock);
int regsPerBlock;
getCudaAttribute<int>( ®sPerBlock, CU_DEVICE_ATTRIBUTE_MAX_REGISTERS_PER_BLOCK, dev );
printf(" Total number of registers available per block: %d\n", regsPerBlock);
int warpSize;
getCudaAttribute<int>( &warpSize, CU_DEVICE_ATTRIBUTE_WARP_SIZE, dev );
printf(" Warp size: %d\n", warpSize);
int maxThreadsPerMultiProcessor;
getCudaAttribute<int>( &maxThreadsPerMultiProcessor, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_MULTIPROCESSOR, dev );
printf(" Maximum number of threads per multiprocessor: %d\n", maxThreadsPerMultiProcessor);
int maxThreadsPerBlock;
getCudaAttribute<int>( &maxThreadsPerBlock, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, dev );
printf(" Maximum number of threads per block: %d\n", maxThreadsPerBlock);
int blockDim[3];
getCudaAttribute<int>( &blockDim[0], CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X, dev );
getCudaAttribute<int>( &blockDim[1], CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y, dev );
getCudaAttribute<int>( &blockDim[2], CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z, dev );
printf(" Maximum sizes of each dimension of a block: %d x %d x %d\n", blockDim[0], blockDim[1], blockDim[2]);
int gridDim[3];
getCudaAttribute<int>( &gridDim[0], CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X, dev );
getCudaAttribute<int>( &gridDim[1], CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y, dev );
getCudaAttribute<int>( &gridDim[2], CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z, dev );
printf(" Maximum sizes of each dimension of a grid: %d x %d x %d\n", gridDim[0], gridDim[1], gridDim[2]);
int textureAlign;
getCudaAttribute<int>( &textureAlign, CU_DEVICE_ATTRIBUTE_TEXTURE_ALIGNMENT, dev );
printf(" Texture alignment: %u bytes\n", textureAlign);
int memPitch;
getCudaAttribute<int>( &memPitch, CU_DEVICE_ATTRIBUTE_MAX_PITCH, dev );
printf(" Maximum memory pitch: %u bytes\n", memPitch);
#if CUDA_VERSION >= 2000
int gpuOverlap;
getCudaAttribute<int>( &gpuOverlap, CU_DEVICE_ATTRIBUTE_GPU_OVERLAP, dev );
#endif
#if CUDA_VERSION >= 4000
int asyncEngineCount;
getCudaAttribute<int>( &asyncEngineCount, CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT, dev );
printf(" Concurrent copy and execution: %s with %d copy engine(s)\n", (gpuOverlap ? "Yes" : "No"), asyncEngineCount);
#else
printf(" Concurrent copy and execution: %s\n",gpuOverlap ? "Yes" : "No");
#endif
#if CUDA_VERSION >= 2020
int kernelExecTimeoutEnabled;
getCudaAttribute<int>( &kernelExecTimeoutEnabled, CU_DEVICE_ATTRIBUTE_KERNEL_EXEC_TIMEOUT, dev );
printf(" Run time limit on kernels: %s\n", kernelExecTimeoutEnabled ? "Yes" : "No");
int integrated;
getCudaAttribute<int>( &integrated, CU_DEVICE_ATTRIBUTE_INTEGRATED, dev );
printf(" Integrated GPU sharing Host Memory: %s\n", integrated ? "Yes" : "No");
int canMapHostMemory;
getCudaAttribute<int>( &canMapHostMemory, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, dev );
printf(" Support host page-locked memory mapping: %s\n", canMapHostMemory ? "Yes" : "No");
#endif
#if CUDA_VERSION >= 3000
int concurrentKernels;
getCudaAttribute<int>( &concurrentKernels, CU_DEVICE_ATTRIBUTE_CONCURRENT_KERNELS, dev );
printf(" Concurrent kernel execution: %s\n", concurrentKernels ? "Yes" : "No");
int surfaceAlignment;
getCudaAttribute<int>( &surfaceAlignment, CU_DEVICE_ATTRIBUTE_SURFACE_ALIGNMENT, dev );
printf(" Alignment requirement for Surfaces: %s\n", surfaceAlignment ? "Yes" : "No");
int eccEnabled;
getCudaAttribute<int>( &eccEnabled, CU_DEVICE_ATTRIBUTE_ECC_ENABLED, dev );
printf(" Device has ECC support enabled: %s\n", eccEnabled ? "Yes" : "No");
#endif
#if CUDA_VERSION >= 3020
int tccDriver ;
getCudaAttribute<int>( &tccDriver , CU_DEVICE_ATTRIBUTE_TCC_DRIVER, dev );
printf(" Device is using TCC driver mode: %s\n", tccDriver ? "Yes" : "No");
#endif
#if CUDA_VERSION >= 4000
int unifiedAddressing;
getCudaAttribute<int>( &unifiedAddressing, CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, dev );
printf(" Device supports Unified Addressing (UVA): %s\n", unifiedAddressing ? "Yes" : "No");
int pciBusID, pciDeviceID;
getCudaAttribute<int>( &pciBusID, CU_DEVICE_ATTRIBUTE_PCI_BUS_ID, dev );
getCudaAttribute<int>( &pciDeviceID, CU_DEVICE_ATTRIBUTE_PCI_DEVICE_ID, dev );
printf(" Device PCI Bus ID / PCI location ID: %d / %d\n", pciBusID, pciDeviceID );
const char *sComputeMode[] = {
"Default (multiple host threads can use ::cudaSetDevice() with device simultaneously)",
"Exclusive (only one host thread in one process is able to use ::cudaSetDevice() with this device)",
"Prohibited (no host thread can use ::cudaSetDevice() with this device)",
"Exclusive Process (many threads in one process is able to use ::cudaSetDevice() with this device)",
"Unknown",
NULL
};
int computeMode;
getCudaAttribute<int>( &computeMode, CU_DEVICE_ATTRIBUTE_COMPUTE_MODE, dev );
printf(" Compute Mode:\n");
printf(" < %s >\n", sComputeMode[computeMode]);
#endif
}
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
// Main Program
//*****************************************************************************
int main(int argc, char** argv)
{
pArgc = &argc;
pArgv = argv;
shrQAStart(argc, argv);
// start logs
cExecutableName = argv[0];
shrSetLogFileName ("oclPostProcessGL.txt");
shrLog("%s Starting...\n\n", argv[0]);
// process command line arguments
if (argc > 1)
{
bQATest = shrCheckCmdLineFlag(argc, (const char**)argv, "qatest");
bNoPrompt = shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
}
shrLog(" Image Width = %d, Image Height = %d, Blur Radius = %d\n\n", image_width, image_height, blur_radius);
// init GL
if(!bQATest)
{
InitGL(&argc, argv);
// create pbo
createPBO(&pbo_source);
createPBO(&pbo_dest);
// create texture for blitting onto the screen
createTexture(&tex_screen, image_width, image_height);
bGLinterop = shrTRUE;
}
// init CL
if( initCL(argc, (const char**)argv) != 0 )
{
return -1;
}
// init fps timer
shrDeltaT (1);
// Create buffers and textures,
// and then start main GLUT rendering loop for processing and rendering,
// or otherwise run No-GL Q/A test sequence
shrLog("\n%s...\n", bQATest ? "No-GL test sequence" : "Standard GL Loop");
printf("\n"
"\tControls\n"
"\t(right click mouse button for Menu)\n"
"\t[ ] : Toggle Post-Processing (blur filter) ON/OFF\n"
"\t[ p ] : Toggle Processing (between GPU or CPU)\n"
"\t[ a ] : Toggle OpenGL Animation (rotation) ON/OFF\n"
"\t[+/=] : Increase Blur Radius\n"
"\t[-/_] : Decrease Blur Radius\n"
"\t[Esc] - Quit\n\n"
);
if(!bQATest)
{
glutMainLoop();
}
else
{
TestNoGL();
}
Cleanup(EXIT_SUCCESS);
}
//////////////////////////////////////////////////////////////////////////////
// Program main
//////////////////////////////////////////////////////////////////////////////
int
main( int argc, char** argv)
{
bool bTestResults = true;
shrQAStart(argc, argv);
if( cutCheckCmdLineFlag(argc, (const char**)argv, "help") ) {
showHelp();
return 0;
}
shrLog("Run \"nbody -benchmark [-n=<numBodies>]\" to measure perfomance.\n");
shrLog("\t-fullscreen (run n-body simulation in fullscreen mode)\n");
shrLog("\t-fp64 (use double precision floating point values for simulation)\n");
shrLog("\t-numdevices=N (use first N CUDA devices for simulation)\n");
// shrLog("\t-hostmem (stores simulation data in host memory)\n");
// shrLog("\t-cpu (performs simulation on the host)\n");
shrLog("\n");
bFullscreen = (cutCheckCmdLineFlag(argc, (const char**) argv, "fullscreen") != 0);
if (bFullscreen)
bShowSliders = false;
benchmark = (cutCheckCmdLineFlag(argc, (const char**) argv, "benchmark") != 0);
compareToCPU = ((cutCheckCmdLineFlag(argc, (const char**) argv, "compare") != 0) ||
(cutCheckCmdLineFlag(argc, (const char**) argv, "qatest") != 0));
QATest = (cutCheckCmdLineFlag(argc, (const char**) argv, "qatest") != 0);
useHostMem = (cutCheckCmdLineFlag(argc, (const char**) argv, "hostmem") != 0);
fp64 = (cutCheckCmdLineFlag(argc, (const char**) argv, "fp64") != 0);
flopsPerInteraction = fp64 ? 30 : 20;
useCpu = (cutCheckCmdLineFlag(argc, (const char**) argv, "cpu") != 0);
cutGetCmdLineArgumenti(argc, (const char**) argv, "numdevices", &numDevsRequested);
// for multi-device we currently require using host memory -- the devices share
// data via the host
if (numDevsRequested > 1)
useHostMem = true;
int numDevsAvailable = 0;
bool customGPU = false;
cudaGetDeviceCount(&numDevsAvailable);
if (numDevsAvailable < numDevsRequested) {
shrLog("Error: only %d Devices available, %d requested. Exiting.\n", numDevsAvailable, numDevsRequested);
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
shrLog("> %s mode\n", bFullscreen ? "Fullscreen" : "Windowed");
shrLog("> Simulation data stored in %s memory\n", useHostMem ? "system" : "video" );
shrLog("> %s precision floating point simulation\n", fp64 ? "Double" : "Single");
shrLog("> %d Devices used for simulation\n", numDevsRequested);
int devID;
cudaDeviceProp props;
// Initialize GL and GLUT if necessary
if (!benchmark && !compareToCPU) {
initGL(&argc, argv);
initParameters();
}
if (useCpu) {
useHostMem = true;
compareToCPU = false;
bSupportDouble = true;
#ifdef OPENMP
shrLog("> Simulation with CPU using OpenMP\n");
#else
shrLog("> Simulation with CPU\n");
#endif
}
else
{
// Now choose the CUDA Device
// Either without GL interop:
if (benchmark || compareToCPU || useHostMem)
{
// Note if we are using host memory for the body system, we
// don't use CUDA-GL interop.
if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") )
{
devID = cutilDeviceInit(argc, argv);
if (devID < 0) {
printf("exiting...\n");
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
customGPU = true;
}
else
{
devID = cutGetMaxGflopsDeviceId();
cudaSetDevice( devID );
}
}
else // or with GL interop:
{
if( cutCheckCmdLineFlag(argc, (const char**)argv, "device") ) {
cutilGLDeviceInit(argc, argv);
customGPU = true;
} else {
devID = cutGetMaxGflopsDeviceId();
cudaGLSetGLDevice( devID );
}
}
cutilSafeCall(cudaGetDevice(&devID));
cutilSafeCall(cudaGetDeviceProperties(&props, devID));
bSupportDouble = true;
#if CUDART_VERSION < 4000
if (numDevsRequested > 1)
{
shrLog("MultiGPU n-body requires CUDA 4.0 or later\n");
cutilDeviceReset();
shrQAFinishExit(argc, (const char**)argv, QA_PASSED);
}
#endif
// Initialize devices
if(numDevsRequested > 1 && customGPU)
{
printf("You can't use --numdevices and --device at the same time.\n");
shrQAFinishExit(argc, (const char**)argv, QA_PASSED);
}
if(customGPU) {
cudaDeviceProp props;
cutilSafeCall(cudaGetDeviceProperties(&props, devID));
shrLog("> Compute %d.%d CUDA device: [%s]\n", props.major, props.minor, props.name);
}
else
{
for (int i = 0; i < numDevsRequested; i++)
{
cudaDeviceProp props;
cutilSafeCall(cudaGetDeviceProperties(&props, i));
shrLog("> Compute %d.%d CUDA device: [%s]\n", props.major, props.minor, props.name);
if (useHostMem)
{
#if CUDART_VERSION >= 2020
if(!props.canMapHostMemory)
{
fprintf(stderr, "Device %d cannot map host memory!\n", devID);
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
if (numDevsRequested > 1)
cutilSafeCall(cudaSetDevice(i));
cutilSafeCall(cudaSetDeviceFlags(cudaDeviceMapHost));
#else
fprintf(stderr, "This CUDART version does not support <cudaDeviceProp.canMapHostMemory> field\n");
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
#endif
}
}
// CC 1.2 and earlier do not support double precision
if (props.major*10 + props.minor <= 12)
bSupportDouble = false;
}
//if(numDevsRequested > 1)
// cutilSafeCall(cudaSetDevice(devID));
if (fp64 && !bSupportDouble) {
fprintf(stderr, "One or more of the requested devices does not support double precision floating-point\n");
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
}
numIterations = 0;
p = 0;
q = 1;
cutGetCmdLineArgumenti(argc, (const char**) argv, "i", &numIterations);
cutGetCmdLineArgumenti(argc, (const char**) argv, "p", &p);
cutGetCmdLineArgumenti(argc, (const char**) argv, "q", &q);
if (p == 0) // p not set on command line
{
p = 256;
if (q * p > 256)
{
p = 256 / q;
shrLog("Setting p=%d, q=%d to maintain %d threads per block\n", p, q, 256);
}
}
// default number of bodies is #SMs * 4 * CTA size
if (useCpu)
#ifdef OPENMP
numBodies = 8192;
#else
numBodies = 4096;
#endif
else if (numDevsRequested == 1)
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int
main( int argc, char** argv)
{
pArgc = &argc;
pArgv = argv;
shrQAStart(argc, argv);
shrSetLogFileName ("deviceQuery.txt");
shrLog("%s Starting...\n\n", argv[0]);
shrLog(" CUDA Device Query (Runtime API) version (CUDART static linking)\n\n");
int deviceCount = 0;
cudaError_t error_id = cudaGetDeviceCount(&deviceCount);
if (error_id != cudaSuccess) {
shrLog( "cudaGetDeviceCount returned %d\n-> %s\n", (int)error_id, cudaGetErrorString(error_id) );
shrQAFinishExit(*pArgc, (const char **)pArgv, QA_FAILED);
}
// This function call returns 0 if there are no CUDA capable devices.
if (deviceCount == 0)
shrLog("There is no device supporting CUDA\n");
else
shrLog("Found %d CUDA Capable device(s)\n", deviceCount);
int dev, driverVersion = 0, runtimeVersion = 0;
for (dev = 0; dev < deviceCount; ++dev) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
shrLog("\nDevice %d: \"%s\"\n", dev, deviceProp.name);
#if CUDART_VERSION >= 2020
// Console log
cudaDriverGetVersion(&driverVersion);
cudaRuntimeGetVersion(&runtimeVersion);
shrLog(" CUDA Driver Version / Runtime Version %d.%d / %d.%d\n", driverVersion/1000, (driverVersion%100)/10, runtimeVersion/1000, (runtimeVersion%100)/10);
#endif
shrLog(" CUDA Capability Major/Minor version number: %d.%d\n", deviceProp.major, deviceProp.minor);
char msg[256];
sprintf(msg, " Total amount of global memory: %.0f MBytes (%llu bytes)\n",
(float)deviceProp.totalGlobalMem/1048576.0f, (unsigned long long) deviceProp.totalGlobalMem);
shrLog(msg);
#if CUDART_VERSION >= 2000
shrLog(" (%2d) Multiprocessors x (%3d) CUDA Cores/MP: %d CUDA Cores\n",
deviceProp.multiProcessorCount,
ConvertSMVer2Cores(deviceProp.major, deviceProp.minor),
ConvertSMVer2Cores(deviceProp.major, deviceProp.minor) * deviceProp.multiProcessorCount);
#endif
shrLog(" GPU Clock rate: %.0f MHz (%0.2f GHz)\n", deviceProp.clockRate * 1e-3f, deviceProp.clockRate * 1e-6f);
#if CUDART_VERSION >= 4000
// This is not available in the CUDA Runtime API, so we make the necessary calls the driver API to support this for output
int memoryClock;
getCudaAttribute<int>( &memoryClock, CU_DEVICE_ATTRIBUTE_MEMORY_CLOCK_RATE, dev );
shrLog(" Memory Clock rate: %.0f Mhz\n", memoryClock * 1e-3f);
int memBusWidth;
getCudaAttribute<int>( &memBusWidth, CU_DEVICE_ATTRIBUTE_GLOBAL_MEMORY_BUS_WIDTH, dev );
shrLog(" Memory Bus Width: %d-bit\n", memBusWidth);
int L2CacheSize;
getCudaAttribute<int>( &L2CacheSize, CU_DEVICE_ATTRIBUTE_L2_CACHE_SIZE, dev );
if (L2CacheSize) {
shrLog(" L2 Cache Size: %d bytes\n", L2CacheSize);
}
shrLog(" Max Texture Dimension Size (x,y,z) 1D=(%d), 2D=(%d,%d), 3D=(%d,%d,%d)\n",
deviceProp.maxTexture1D, deviceProp.maxTexture2D[0], deviceProp.maxTexture2D[1],
deviceProp.maxTexture3D[0], deviceProp.maxTexture3D[1], deviceProp.maxTexture3D[2]);
shrLog(" Max Layered Texture Size (dim) x layers 1D=(%d) x %d, 2D=(%d,%d) x %d\n",
deviceProp.maxTexture1DLayered[0], deviceProp.maxTexture1DLayered[1],
deviceProp.maxTexture2DLayered[0], deviceProp.maxTexture2DLayered[1], deviceProp.maxTexture2DLayered[2]);
#endif
shrLog(" Total amount of constant memory: %u bytes\n", deviceProp.totalConstMem);
shrLog(" Total amount of shared memory per block: %u bytes\n", deviceProp.sharedMemPerBlock);
shrLog(" Total number of registers available per block: %d\n", deviceProp.regsPerBlock);
shrLog(" Warp size: %d\n", deviceProp.warpSize);
shrLog(" Maximum number of threads per multiprocessor: %d\n", deviceProp.maxThreadsPerMultiProcessor);
shrLog(" Maximum number of threads per block: %d\n", deviceProp.maxThreadsPerBlock);
shrLog(" Maximum sizes of each dimension of a block: %d x %d x %d\n",
deviceProp.maxThreadsDim[0],
deviceProp.maxThreadsDim[1],
deviceProp.maxThreadsDim[2]);
shrLog(" Maximum sizes of each dimension of a grid: %d x %d x %d\n",
deviceProp.maxGridSize[0],
deviceProp.maxGridSize[1],
deviceProp.maxGridSize[2]);
shrLog(" Maximum memory pitch: %u bytes\n", deviceProp.memPitch);
shrLog(" Texture alignment: %u bytes\n", deviceProp.textureAlignment);
#if CUDART_VERSION >= 4000
shrLog(" Concurrent copy and execution: %s with %d copy engine(s)\n", (deviceProp.deviceOverlap ? "Yes" : "No"), deviceProp.asyncEngineCount);
#else
shrLog(" Concurrent copy and execution: %s\n", deviceProp.deviceOverlap ? "Yes" : "No");
#endif
#if CUDART_VERSION >= 2020
shrLog(" Run time limit on kernels: %s\n", deviceProp.kernelExecTimeoutEnabled ? "Yes" : "No");
shrLog(" Integrated GPU sharing Host Memory: %s\n", deviceProp.integrated ? "Yes" : "No");
shrLog(" Support host page-locked memory mapping: %s\n", deviceProp.canMapHostMemory ? "Yes" : "No");
#endif
#if CUDART_VERSION >= 3000
shrLog(" Concurrent kernel execution: %s\n", deviceProp.concurrentKernels ? "Yes" : "No");
shrLog(" Alignment requirement for Surfaces: %s\n", deviceProp.surfaceAlignment ? "Yes" : "No");
#endif
#if CUDART_VERSION >= 3010
shrLog(" Device has ECC support enabled: %s\n", deviceProp.ECCEnabled ? "Yes" : "No");
#endif
#if CUDART_VERSION >= 3020
shrLog(" Device is using TCC driver mode: %s\n", deviceProp.tccDriver ? "Yes" : "No");
#endif
#if CUDART_VERSION >= 4000
shrLog(" Device supports Unified Addressing (UVA): %s\n", deviceProp.unifiedAddressing ? "Yes" : "No");
shrLog(" Device PCI Bus ID / PCI location ID: %d / %d\n", deviceProp.pciBusID, deviceProp.pciDeviceID );
#endif
#if CUDART_VERSION >= 2020
const char *sComputeMode[] = {
"Default (multiple host threads can use ::cudaSetDevice() with device simultaneously)",
"Exclusive (only one host thread in one process is able to use ::cudaSetDevice() with this device)",
"Prohibited (no host thread can use ::cudaSetDevice() with this device)",
"Exclusive Process (many threads in one process is able to use ::cudaSetDevice() with this device)",
"Unknown",
NULL
};
shrLog(" Compute Mode:\n");
shrLog(" < %s >\n", sComputeMode[deviceProp.computeMode]);
#endif
}
// csv masterlog info
// *****************************
// exe and CUDA driver name
shrLog("\n");
std::string sProfileString = "deviceQuery, CUDA Driver = CUDART";
char cTemp[10];
// driver version
sProfileString += ", CUDA Driver Version = ";
#ifdef WIN32
sprintf_s(cTemp, 10, "%d.%d", driverVersion/1000, (driverVersion%100)/10);
#else
sprintf(cTemp, "%d.%d", driverVersion/1000, (driverVersion%100)/10);
#endif
sProfileString += cTemp;
// Runtime version
sProfileString += ", CUDA Runtime Version = ";
#ifdef WIN32
sprintf_s(cTemp, 10, "%d.%d", runtimeVersion/1000, (runtimeVersion%100)/10);
#else
sprintf(cTemp, "%d.%d", runtimeVersion/1000, (runtimeVersion%100)/10);
#endif
sProfileString += cTemp;
// Device count
sProfileString += ", NumDevs = ";
#ifdef WIN32
sprintf_s(cTemp, 10, "%d", deviceCount);
#else
sprintf(cTemp, "%d", deviceCount);
#endif
sProfileString += cTemp;
// First 2 device names, if any
for (dev = 0; dev < ((deviceCount > 2) ? 2 : deviceCount); ++dev)
{
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
sProfileString += ", Device = ";
sProfileString += deviceProp.name;
}
sProfileString += "\n";
shrLogEx(LOGBOTH | MASTER, 0, sProfileString.c_str());
// finish
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
////////////////////////////////////////////////////////////////////////////////
// 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));
}
// Main function
// *********************************************************************
int main(int argc, char **argv)
{
shrQAStart(argc, argv);
int NUM_BLOCKS = 10;
shrSetLogFileName ("Barrier_Centralized.txt");
while(NUM_BLOCKS<=120)
{
int iNumElements = NUM_BLOCKS* NUM_THREADS; // total num of threads
// BARRIER GOAL
int goal_val = NUM_BLOCKS;
// get command line arg for quick test, if provided
bNoPrompt = shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
// start logs
cExecutableName = argv[0];
shrSetLogFileName ("Barrier.txt");
shrLog("%s Starting...\n\n# of THREADS \t= %i\n", argv[0], iNumElements);
// set and log Global and Local work size dimensions
szLocalWorkSize = NUM_THREADS ;
szGlobalWorkSize = shrRoundUp((int)szLocalWorkSize, iNumElements); // rounded up to the nearest multiple of the LocalWorkSize
shrLog("Global Work Size \t\t= %u\nLocal Work Size \t\t= %u\n# of Work Groups \t\t= %u\n\n",
szGlobalWorkSize, szLocalWorkSize, (szGlobalWorkSize % szLocalWorkSize + szGlobalWorkSize/szLocalWorkSize));
//Get an OpenCL platform
ciErr1 = clGetPlatformIDs(1, &cpPlatform, NULL);
shrLog("clGetPlatformID...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clGetPlatformID, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
//Get the devices
ciErr1 = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL);
shrLog("clGetDeviceIDs...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clGetDeviceIDs, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
//Create the context
cxGPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErr1);
shrLog("clCreateContext...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateContext, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Create a command-queue
cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevice, CL_QUEUE_PROFILING_ENABLE, &ciErr1);
shrLog("clCreateCommandQueue...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateCommandQueue, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Read the OpenCL kernel in from source file
shrLog("oclLoadProgSource (%s)...\n", cSourceFile);
cPathAndName = shrFindFilePath(cSourceFile, argv[0]);
cSourceCL = oclLoadProgSource(cPathAndName, "", &szKernelLength);
// Create the program
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErr1);
shrLog("clCreateProgramWithSource...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateProgramWithSource, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Build the program with 'mad' Optimization option
#ifdef MAC
char* flags = "-cl-fast-relaxed-math -DMAC";
#else
char* flags = "-cl-fast-relaxed-math";
#endif
ciErr1 = clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL);
shrLog("clBuildProgram...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clBuildProgram, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Create the kernel
ckKernel = clCreateKernel(cpProgram, "Barrier", &ciErr1);
shrLog("clCreateKernel (Barrier)...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateKernel, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Allocate and initialize host arrays
shrLog( "Allocate and Init Host Mem...\n");
input = (int *)malloc(sizeof(int) * NUM_BLOCKS);
for(int i =0; i<=NUM_BLOCKS; i++)
{
input[i]=0;
}
// Allocate the OpenCL buffer memory objects for source and result on the device GMEM
array_in = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, sizeof(int)* NUM_BLOCKS, NULL, &ciErr1);
array_out = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, sizeof(int)* NUM_BLOCKS, NULL, &ciErr1);
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clCreateBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Set the Argument values
ciErr1 = clSetKernelArg(ckKernel, 0, sizeof(cl_int), (void*)&goal_val);
ciErr1 |= clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void*)&array_in);
ciErr1 |= clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void*)&array_out);
// ciErr1 |= clSetKernelArg(ckKernel, 1, sizeof(cl_int), (void*)&iNumElements);
shrLog("clSetKernelArg 0 - 2...\n\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clSetKernelArg, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// --------------------------------------------------------
// Start Core sequence... copy input data to GPU, compute, copy results back
ciErr1 = clEnqueueWriteBuffer(cqCommandQueue, array_in, CL_FALSE, 0, sizeof(int) * NUM_BLOCKS,(void*) input, 0, NULL, NULL);
shrLog("clEnqueueWriteBuffer (SrcA and SrcB)...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clEnqueueWriteBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
// Launch kernel
ciErr1 = clEnqueueNDRangeKernel(cqCommandQueue, ckKernel, 1, NULL, &szGlobalWorkSize, &szLocalWorkSize, 0, NULL, &ceEvent);
shrLog("clEnqueueNDRangeKernel (Barrier)...\n");
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clEnqueueNDRangeKernel, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}
/*ciErr1 = clEnqueueReadBuffer(cqCommandQueue, global_mutex, CL_TRUE, 0, sizeof(cl_int), &original_goal, 0, NULL, NULL);
shrLog("clEnqueueReadBuffer (Dst)...%d \n\n", original_goal);
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error in clEnqueueReadBuffer, Line %u in file %s !!!\n\n", __LINE__, __FILE__);
Cleanup(argc, argv, EXIT_FAILURE);
}*/
//GPU_PROFILING
ciErr1=clWaitForEvents(1, &ceEvent);
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error 1 !\n\n");
Cleanup(argc, argv, EXIT_FAILURE);
}
cl_ulong start, end;
ciErr1 = clGetEventProfilingInfo(ceEvent, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
ciErr1 |= clGetEventProfilingInfo(ceEvent, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
if (ciErr1 != CL_SUCCESS)
{
shrLog("Error 2 !\n\n");
Cleanup(argc, argv, EXIT_FAILURE);
}
double dSeconds = 1.0e-9 * (double)(end - start);
shrLog("Done! time taken %ul \n",end - start );
// shrLog("Done! Kernel execution time: %.5f s\n\n", dSeconds);
// Release event
clReleaseEvent(ceEvent);
ceEvent = 0;
Cleanup (argc, argv, EXIT_SUCCESS);
NUM_BLOCKS = NUM_BLOCKS+10;
}
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
int main(int argc, char **argv)
{
shrQAStart(argc, argv);
if (argc > 1) {
if (cutCheckCmdLineFlag(argc, (const char **)argv, "qatest") ||
cutCheckCmdLineFlag(argc, (const char **)argv, "noprompt"))
{
g_bQAReadback = true;
fpsLimit = frameCheckNumber;
}
if (cutCheckCmdLineFlag(argc, (const char **)argv, "glverify"))
{
g_bOpenGLQA = true;
g_bFBODisplay = false;
fpsLimit = frameCheckNumber;
}
if (cutCheckCmdLineFlag(argc, (const char **)argv, "fbo")) {
g_bFBODisplay = true;
fpsLimit = frameCheckNumber;
}
}
if (g_bQAReadback) {
runAutoTest(argc, argv);
} else {
printf("[%s] ", sSDKsample);
if (g_bFBODisplay) printf("[FBO Display] ");
if (g_bOpenGLQA) printf("[OpenGL Readback Comparisons] ");
printf("\n");
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
if ( cutCheckCmdLineFlag(argc, (const char **)argv, "device")) {
printf("[%s]\n", argv[0]);
printf(" Does not explicitly support -device=n in OpenGL mode\n");
printf(" To use -device=n, the sample must be running w/o OpenGL\n\n");
printf(" > %s -device=n -qatest\n", argv[0]);
printf("exiting...\n");
exit(0);
}
// First load the image, so we know what the size of the image (imageW and imageH)
printf("Allocating host and CUDA memory and loading image file...\n");
const char *image_path = cutFindFilePath("portrait_noise.bmp", argv[0]);
if (image_path == NULL) {
printf( "imageDenoisingGL was unable to find and load image file <portrait_noise.bmp>.\nExiting...\n");
shrQAFinishExit(argc, (const char **)argv, QA_FAILED);
}
LoadBMPFile(&h_Src, &imageW, &imageH, image_path);
printf("Data init done.\n");
// First initialize OpenGL context, so we can properly set the GL for CUDA.
// This is necessary in order to achieve optimal performance with OpenGL/CUDA interop.
initGL( &argc, argv );
cudaGLSetGLDevice( cutGetMaxGflopsDeviceId() );
cutilSafeCall( CUDA_MallocArray(&h_Src, imageW, imageH) );
initOpenGLBuffers();
// Creating the Auto-Validation Code
if (g_bOpenGLQA) {
if (g_bFBODisplay) {
g_CheckRender = new CheckFBO(imageW, imageH, 4);
} else {
g_CheckRender = new CheckBackBuffer(imageW, imageH, 4);
}
g_CheckRender->setPixelFormat(GL_RGBA);
g_CheckRender->setExecPath(argv[0]);
g_CheckRender->EnableQAReadback(g_bOpenGLQA);
}
}
printf("Starting GLUT main loop...\n");
printf("Press [1] to view noisy image\n");
printf("Press [2] to view image restored with knn filter\n");
printf("Press [3] to view image restored with nlm filter\n");
printf("Press [4] to view image restored with modified nlm filter\n");
printf("Press [ ] to view smooth/edgy areas [RED/BLUE] Ct's\n");
printf("Press [f] to print frame rate\n");
printf("Press [?] to print Noise and Lerp Ct's\n");
printf("Press [q] to exit\n");
glutDisplayFunc(displayFunc);
glutKeyboardFunc(shutDown);
cutilCheckError( cutCreateTimer(&hTimer) );
cutilCheckError( cutStartTimer(hTimer) );
glutTimerFunc(REFRESH_DELAY, timerEvent,0);
glutMainLoop();
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
// Main program
//*****************************************************************************
int main(int argc, char** argv)
{
pArgc = &argc;
pArgv = argv;
shrQAStart(argc, argv);
// Start logs
cExecutableName = argv[0];
shrSetLogFileName ("oclSobelFilter.txt");
shrLog("%s Starting (Using %s)...\n\n", argv[0], clSourcefile);
// Get command line args for quick test or QA test, if provided
bNoPrompt = (bool)shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
bQATest = (bool)shrCheckCmdLineFlag(argc, (const char**)argv, "qatest");
// Menu items
if (!(bQATest))
{
ShowMenuItems();
}
// Find the path from the exe to the image file
cPathAndName = shrFindFilePath(cImageFile, argv[0]);
oclCheckErrorEX(cPathAndName != NULL, shrTRUE, pCleanup);
shrLog("Image File\t = %s\nImage Dimensions = %u w x %u h x %u bpp\n\n", cPathAndName, uiImageWidth, uiImageHeight, sizeof(unsigned int)<<3);
// Initialize OpenGL items (if not No-GL QA test)
shrLog("%sInitGL...\n\n", bQATest ? "Skipping " : "Calling ");
if (!(bQATest))
{
InitGL(&argc, argv);
}
//Get the NVIDIA platform if available, otherwise use default
char cBuffer[1024];
bool bNV = false;
shrLog("Get Platform ID... ");
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
ciErrNum = clGetPlatformInfo (cpPlatform, CL_PLATFORM_NAME, sizeof(cBuffer), cBuffer, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("%s\n\n", cBuffer);
bNV = (strstr(cBuffer, "NVIDIA") != NULL);
//Get the devices
shrLog("Get Device Info...\n");
cl_uint uiNumAllDevs = 0;
GpuDevMngr = new DeviceManager(cpPlatform, &uiNumAllDevs, pCleanup);
// Get selected device if specified, otherwise examine avaiable ones and choose by perf
cl_int iSelectedDevice = 0;
if((shrGetCmdLineArgumenti(argc, (const char**)argv, "device", &iSelectedDevice)) || (uiNumAllDevs == 1))
{
// Use 1 selected device
GpuDevMngr->uiUsefulDevCt = 1;
iSelectedDevice = CLAMP((cl_uint)iSelectedDevice, 0, (uiNumAllDevs - 1));
GpuDevMngr->uiUsefulDevs[0] = iSelectedDevice;
GpuDevMngr->fLoadProportions[0] = 1.0f;
shrLog(" Using 1 Selected Device for Sobel Filter Computation...\n");
}
else
{
// Use available useful devices and Compute the device load proportions
ciErrNum = GpuDevMngr->GetDevLoadProportions(bNV);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
if (GpuDevMngr->uiUsefulDevCt == 1)
{
iSelectedDevice = GpuDevMngr->uiUsefulDevs[0];
}
shrLog(" Using %u Device(s) for Sobel Filter Computation\n", GpuDevMngr->uiUsefulDevCt);
}
//Create the context
shrLog("\nclCreateContext...\n\n");
cxGPUContext = clCreateContext(0, uiNumAllDevs, GpuDevMngr->cdDevices, NULL, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
// Allocate per-device OpenCL objects for useful devices
cqCommandQueue = new cl_command_queue[GpuDevMngr->uiUsefulDevCt];
ckSobel = new cl_kernel[GpuDevMngr->uiUsefulDevCt];
cmDevBufIn = new cl_mem[GpuDevMngr->uiUsefulDevCt];
cmDevBufOut = new cl_mem[GpuDevMngr->uiUsefulDevCt];
szAllocDevBytes = new size_t[GpuDevMngr->uiUsefulDevCt];
uiInHostPixOffsets = new cl_uint[GpuDevMngr->uiUsefulDevCt];
uiOutHostPixOffsets = new cl_uint[GpuDevMngr->uiUsefulDevCt];
uiDevImageHeight = new cl_uint[GpuDevMngr->uiUsefulDevCt];
// Create command queue(s) for device(s)
shrLog("clCreateCommandQueue...\n");
for (cl_uint i = 0; i < GpuDevMngr->uiUsefulDevCt; i++)
{
cqCommandQueue[i] = clCreateCommandQueue(cxGPUContext, GpuDevMngr->cdDevices[GpuDevMngr->uiUsefulDevs[i]], 0, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog(" CommandQueue %u, Device %u, Device Load Proportion = %.2f, ", i, GpuDevMngr->uiUsefulDevs[i], GpuDevMngr->fLoadProportions[i]);
oclPrintDevName(LOGBOTH, GpuDevMngr->cdDevices[GpuDevMngr->uiUsefulDevs[i]]);
shrLog("\n");
}
// Allocate pinned input and output host image buffers: mem copy operations to/from pinned memory is much faster than paged memory
szBuffBytes = uiImageWidth * uiImageHeight * sizeof (unsigned int);
cmPinnedBufIn = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, szBuffBytes, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cmPinnedBufOut = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR, szBuffBytes, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("\nclCreateBuffer (Input and Output Pinned Host buffers)...\n");
// Get mapped pointers for writing to pinned input and output host image pointers
uiInput = (cl_uint*)clEnqueueMapBuffer(cqCommandQueue[0], cmPinnedBufIn, CL_TRUE, CL_MAP_WRITE, 0, szBuffBytes, 0, NULL, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
uiOutput = (cl_uint*)clEnqueueMapBuffer(cqCommandQueue[0], cmPinnedBufOut, CL_TRUE, CL_MAP_READ, 0, szBuffBytes, 0, NULL, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("clEnqueueMapBuffer (Pointer to Input and Output pinned host buffers)...\n");
// Load image data from file to pinned input host buffer
ciErrNum = shrLoadPPM4ub(cPathAndName, (unsigned char **)&uiInput, &uiImageWidth, &uiImageHeight);
oclCheckErrorEX(ciErrNum, shrTRUE, pCleanup);
shrLog("Load Input Image to Input pinned host buffer...\n");
// Read the kernel in from file
free(cPathAndName);
cPathAndName = shrFindFilePath(clSourcefile, argv[0]);
oclCheckErrorEX(cPathAndName != NULL, shrTRUE, pCleanup);
cSourceCL = oclLoadProgSource(cPathAndName, "// My comment\n", &szKernelLength);
oclCheckErrorEX(cSourceCL != NULL, shrTRUE, pCleanup);
shrLog("Load OpenCL Prog Source from File...\n");
// Create the program object
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&cSourceCL, &szKernelLength, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("clCreateProgramWithSource...\n");
// Build the program with 'mad' Optimization option
#ifdef MAC
char *flags = "-cl-fast-relaxed-math -DMAC";
#else
char *flags = "-cl-fast-relaxed-math";
#endif
ciErrNum = clBuildProgram(cpProgram, 0, NULL, flags, NULL, NULL);
if (ciErrNum != CL_SUCCESS)
{
// On error: write out standard error, Build Log and PTX, then cleanup and exit
shrLogEx(LOGBOTH | ERRORMSG, ciErrNum, STDERROR);
oclLogBuildInfo(cpProgram, oclGetFirstDev(cxGPUContext));
oclLogPtx(cpProgram, oclGetFirstDev(cxGPUContext), "oclSobelFilter.ptx");
Cleanup(EXIT_FAILURE);
}
shrLog("clBuildProgram...\n\n");
// Determine, the size/shape of the image portions for each dev and create the device buffers
unsigned uiSumHeight = 0;
for (cl_uint i = 0; i < GpuDevMngr->uiUsefulDevCt; i++)
{
// Create kernel instance
ckSobel[i] = clCreateKernel(cpProgram, "ckSobel", &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("clCreateKernel (ckSobel), Device %u...\n", i);
// Allocations and offsets for the portion of the image worked on by each device
if (GpuDevMngr->uiUsefulDevCt == 1)
{
// One device processes the whole image with no offset
uiDevImageHeight[i] = uiImageHeight;
uiInHostPixOffsets[i] = 0;
uiOutHostPixOffsets[i] = 0;
szAllocDevBytes[i] = uiDevImageHeight[i] * uiImageWidth * sizeof(cl_uint);
}
else if (i == 0)
{
// Multiple devices, top stripe zone including topmost row of image:
// Over-allocate on device by 1 row
// Set offset and size to copy extra 1 padding row H2D (below bottom of stripe)
// Won't return the last row (dark/garbage row) D2H
uiInHostPixOffsets[i] = 0;
uiOutHostPixOffsets[i] = 0;
uiDevImageHeight[i] = (cl_uint)(GpuDevMngr->fLoadProportions[GpuDevMngr->uiUsefulDevs[i]] * (float)uiImageHeight); // height is proportional to dev perf
uiSumHeight += uiDevImageHeight[i];
uiDevImageHeight[i] += 1;
szAllocDevBytes[i] = uiDevImageHeight[i] * uiImageWidth * sizeof(cl_uint);
}
else if (i < (GpuDevMngr->uiUsefulDevCt - 1))
{
// Multiple devices, middle stripe zone:
// Over-allocate on device by 2 rows
// Set offset and size to copy extra 2 padding rows H2D (above top and below bottom of stripe)
// Won't return the first and last rows (dark/garbage rows) D2H
uiInHostPixOffsets[i] = (uiSumHeight - 1) * uiImageWidth;
uiOutHostPixOffsets[i] = uiInHostPixOffsets[i] + uiImageWidth;
uiDevImageHeight[i] = (cl_uint)(GpuDevMngr->fLoadProportions[GpuDevMngr->uiUsefulDevs[i]] * (float)uiImageHeight); // height is proportional to dev perf
uiSumHeight += uiDevImageHeight[i];
uiDevImageHeight[i] += 2;
szAllocDevBytes[i] = uiDevImageHeight[i] * uiImageWidth * sizeof(cl_uint);
}
else
{
// Multiple devices, last boundary tile:
// Over-allocate on device by 1 row
// Set offset and size to copy extra 1 padding row H2D (above top of stripe)
// Won't return the first row (dark/garbage rows D2H
uiInHostPixOffsets[i] = (uiSumHeight - 1) * uiImageWidth;
uiOutHostPixOffsets[i] = uiInHostPixOffsets[i] + uiImageWidth;
uiDevImageHeight[i] = uiImageHeight - uiSumHeight; // "leftover" rows
uiSumHeight += uiDevImageHeight[i];
uiDevImageHeight[i] += 1;
szAllocDevBytes[i] = uiDevImageHeight[i] * uiImageWidth * sizeof(cl_uint);
}
shrLog("Image Height (rows) for Device %u = %u...\n", i, uiDevImageHeight[i]);
// Create the device buffers in GMEM on each device
cmDevBufIn[i] = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, szAllocDevBytes[i], NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
cmDevBufOut[i] = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, szAllocDevBytes[i], NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("clCreateBuffer (Input and Output GMEM buffers, Device %u)...\n", i);
// Set the common argument values for the Median kernel instance for each device
int iLocalPixPitch = iBlockDimX + 2;
ciErrNum = clSetKernelArg(ckSobel[i], 0, sizeof(cl_mem), (void*)&cmDevBufIn[i]);
ciErrNum |= clSetKernelArg(ckSobel[i], 1, sizeof(cl_mem), (void*)&cmDevBufOut[i]);
ciErrNum |= clSetKernelArg(ckSobel[i], 2, (iLocalPixPitch * (iBlockDimY + 2) * sizeof(cl_uchar4)), NULL);
ciErrNum |= clSetKernelArg(ckSobel[i], 3, sizeof(cl_int), (void*)&iLocalPixPitch);
ciErrNum |= clSetKernelArg(ckSobel[i], 4, sizeof(cl_uint), (void*)&uiImageWidth);
ciErrNum |= clSetKernelArg(ckSobel[i], 6, sizeof(cl_float), (void*)&fThresh);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
shrLog("clSetKernelArg (0-4), Device %u...\n\n", i);
}
// Set common global and local work sizes for Median kernel
szLocalWorkSize[0] = iBlockDimX;
szLocalWorkSize[1] = iBlockDimY;
szGlobalWorkSize[0] = shrRoundUp((int)szLocalWorkSize[0], uiImageWidth);
// init running timers
shrDeltaT(0); // timer 0 used for computation timing
shrDeltaT(1); // timer 1 used for fps computation
// Start main GLUT rendering loop for processing and rendering,
// or otherwise run No-GL Q/A test sequence
if (!(bQATest))
{
glutMainLoop();
}
else
{
TestNoGL();
}
Cleanup(EXIT_SUCCESS);
}
// Main function
// *********************************************************************
int main(int argc, char** argv)
{
shrQAStart(argc, argv);
int use_gpu = 0;
for(int i = 0; i < argc && argv; i++)
{
if(!argv[i])
continue;
if(strstr(argv[i], "cpu"))
use_gpu = 0;
else if(strstr(argv[i], "gpu"))
use_gpu = 1;
}
// start logs
shrSetLogFileName ("oclDXTCompression.txt");
shrLog("%s Starting...\n\n", argv[0]);
cl_platform_id cpPlatform = NULL;
cl_uint uiNumDevices = 0;
cl_device_id *cdDevices = NULL;
cl_context cxGPUContext;
cl_command_queue cqCommandQueue;
cl_program cpProgram;
cl_kernel ckKernel;
cl_mem cmMemObjs[3];
cl_mem cmAlphaTable4, cmProds4;
cl_mem cmAlphaTable3, cmProds3;
size_t szGlobalWorkSize[1];
size_t szLocalWorkSize[1];
cl_int ciErrNum;
// Get the path of the filename
char *filename;
if (shrGetCmdLineArgumentstr(argc, (const char **)argv, "image", &filename)) {
image_filename = filename;
}
// load image
const char* image_path = shrFindFilePath(image_filename, argv[0]);
oclCheckError(image_path != NULL, shrTRUE);
shrLoadPPM4ub(image_path, (unsigned char **)&h_img, &width, &height);
oclCheckError(h_img != NULL, shrTRUE);
shrLog("Loaded '%s', %d x %d pixels\n\n", image_path, width, height);
// Convert linear image to block linear.
const uint memSize = width * height * sizeof(cl_uint);
uint* block_image = (uint*)malloc(memSize);
// Convert linear image to block linear.
for(uint by = 0; by < height/4; by++) {
for(uint bx = 0; bx < width/4; bx++) {
for (int i = 0; i < 16; i++) {
const int x = i & 3;
const int y = i / 4;
block_image[(by * width/4 + bx) * 16 + i] =
((uint *)h_img)[(by * 4 + y) * 4 * (width/4) + bx * 4 + x];
}
}
}
// Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckError(ciErrNum, CL_SUCCESS);
// Get the platform's GPU devices
ciErrNum = clGetDeviceIDs(cpPlatform, use_gpu?CL_DEVICE_TYPE_GPU:CL_DEVICE_TYPE_CPU, 0, NULL, &uiNumDevices);
oclCheckError(ciErrNum, CL_SUCCESS);
cdDevices = (cl_device_id *)malloc(uiNumDevices * sizeof(cl_device_id) );
ciErrNum = clGetDeviceIDs(cpPlatform, use_gpu?CL_DEVICE_TYPE_GPU:CL_DEVICE_TYPE_CPU, uiNumDevices, cdDevices, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
// Create the context
cxGPUContext = clCreateContext(0, uiNumDevices, cdDevices, NULL, NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
// get and log device
cl_device_id device;
if( shrCheckCmdLineFlag(argc, (const char **)argv, "device") ) {
int device_nr = 0;
shrGetCmdLineArgumenti(argc, (const char **)argv, "device", &device_nr);
device = oclGetDev(cxGPUContext, device_nr);
if( device == (cl_device_id)-1 ) {
shrLog(" Invalid GPU Device: devID=%d. %d valid GPU devices detected\n\n", device_nr, uiNumDevices);
shrLog(" exiting...\n");
return -1;
}
} else {
device = oclGetMaxFlopsDev(cxGPUContext);
}
oclPrintDevName(LOGBOTH, device);
shrLog("\n");
// create a command-queue
cqCommandQueue = clCreateCommandQueue(cxGPUContext, device, 0, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
// Memory Setup
// Constants
cmAlphaTable4 = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 4 * sizeof(cl_float), (void*)&alphaTable4[0], &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
cmProds4 = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 4 * sizeof(cl_int), (void*)&prods4[0], &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
cmAlphaTable3 = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 4 * sizeof(cl_float), (void*)&alphaTable3[0], &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
cmProds3 = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 4 * sizeof(cl_int), (void*)&prods3[0], &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
// Compute permutations.
cl_uint permutations[1024];
computePermutations(permutations);
// Upload permutations.
cmMemObjs[0] = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(cl_uint) * 1024, permutations, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
// Image
cmMemObjs[1] = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, memSize, NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
// Result
const uint compressedSize = (width / 4) * (height / 4) * 8;
cmMemObjs[2] = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, compressedSize, NULL , &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
unsigned int * h_result = (uint*)malloc(compressedSize);
// Program Setup
size_t program_length;
const char* source_path = shrFindFilePath("DXTCompression.cl", argv[0]);
oclCheckError(source_path != NULL, shrTRUE);
char *source = oclLoadProgSource(source_path, "", &program_length);
oclCheckError(source != NULL, shrTRUE);
// create the program
cpProgram = clCreateProgramWithSource(cxGPUContext, 1,
(const char **) &source, &program_length, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
// build the program
ciErrNum = clBuildProgram(cpProgram, 0, NULL, "-cl-fast-relaxed-math", NULL, NULL);
if (ciErrNum != CL_SUCCESS)
{
// write out standard error, Build Log and PTX, then cleanup and exit
shrLogEx(LOGBOTH | ERRORMSG, ciErrNum, STDERROR);
oclLogBuildInfo(cpProgram, oclGetFirstDev(cxGPUContext));
oclLogPtx(cpProgram, oclGetFirstDev(cxGPUContext), "oclDXTCompression.ptx");
oclCheckError(ciErrNum, CL_SUCCESS);
}
// create the kernel
ckKernel = clCreateKernel(cpProgram, "compress", &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
// set the args values
ciErrNum = clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void *) &cmMemObjs[0]);
ciErrNum |= clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void *) &cmMemObjs[1]);
ciErrNum |= clSetKernelArg(ckKernel, 2, sizeof(cl_mem), (void *) &cmMemObjs[2]);
ciErrNum |= clSetKernelArg(ckKernel, 3, sizeof(cl_mem), (void*)&cmAlphaTable4);
ciErrNum |= clSetKernelArg(ckKernel, 4, sizeof(cl_mem), (void*)&cmProds4);
ciErrNum |= clSetKernelArg(ckKernel, 5, sizeof(cl_mem), (void*)&cmAlphaTable3);
ciErrNum |= clSetKernelArg(ckKernel, 6, sizeof(cl_mem), (void*)&cmProds3);
oclCheckError(ciErrNum, CL_SUCCESS);
// Copy input data host to device
clEnqueueWriteBuffer(cqCommandQueue, cmMemObjs[1], CL_FALSE, 0, sizeof(cl_uint) * width * height, block_image, 0,0,0);
// Determine launch configuration and run timed computation numIterations times
int blocks = ((width + 3) / 4) * ((height + 3) / 4); // rounds up by 1 block in each dim if %4 != 0
// Restrict the numbers of blocks to launch on low end GPUs to avoid kernel timeout
cl_uint compute_units;
clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(compute_units), &compute_units, NULL);
int blocksPerLaunch = MIN(blocks, 768 * (int)compute_units);
// set work-item dimensions
szGlobalWorkSize[0] = blocksPerLaunch * NUM_THREADS;
szLocalWorkSize[0]= NUM_THREADS;
#ifdef GPU_PROFILING
shrLog("\nRunning DXT Compression on %u x %u image...\n", width, height);
shrLog("\n%u Workgroups, %u Work Items per Workgroup, %u Work Items in NDRange...\n\n",
blocks, NUM_THREADS, blocks * NUM_THREADS);
int numIterations = 50;
for (int i = -1; i < numIterations; ++i) {
if (i == 0) { // start timing only after the first warmup iteration
clFinish(cqCommandQueue); // flush command queue
shrDeltaT(0); // start timer
}
#endif
// execute kernel
for( int j=0; j<blocks; j+= blocksPerLaunch ) {
clSetKernelArg(ckKernel, 7, sizeof(int), &j);
szGlobalWorkSize[0] = MIN( blocksPerLaunch, blocks-j ) * NUM_THREADS;
ciErrNum = clEnqueueNDRangeKernel(cqCommandQueue, ckKernel, 1, NULL,
szGlobalWorkSize, szLocalWorkSize,
0, NULL, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
}
#ifdef GPU_PROFILING
}
clFinish(cqCommandQueue);
double dAvgTime = shrDeltaT(0) / (double)numIterations;
shrLogEx(LOGBOTH | MASTER, 0, "oclDXTCompression, Throughput = %.4f MPixels/s, Time = %.5f s, Size = %u Pixels, NumDevsUsed = %i, Workgroup = %d\n",
(1.0e-6 * (double)(width * height)/ dAvgTime), dAvgTime, (width * height), 1, szLocalWorkSize[0]);
#endif
// blocking read output
ciErrNum = clEnqueueReadBuffer(cqCommandQueue, cmMemObjs[2], CL_TRUE, 0,
compressedSize, h_result, 0, NULL, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
// Write DDS file.
FILE* fp = NULL;
char output_filename[1024];
#ifdef WIN32
strcpy_s(output_filename, 1024, image_path);
strcpy_s(output_filename + strlen(image_path) - 3, 1024 - strlen(image_path) + 3, "dds");
fopen_s(&fp, output_filename, "wb");
#else
strcpy(output_filename, image_path);
strcpy(output_filename + strlen(image_path) - 3, "dds");
fp = fopen(output_filename, "wb");
#endif
oclCheckError(fp != NULL, shrTRUE);
DDSHeader header;
header.fourcc = FOURCC_DDS;
header.size = 124;
header.flags = (DDSD_WIDTH|DDSD_HEIGHT|DDSD_CAPS|DDSD_PIXELFORMAT|DDSD_LINEARSIZE);
header.height = height;
header.width = width;
header.pitch = compressedSize;
header.depth = 0;
header.mipmapcount = 0;
memset(header.reserved, 0, sizeof(header.reserved));
header.pf.size = 32;
header.pf.flags = DDPF_FOURCC;
header.pf.fourcc = FOURCC_DXT1;
header.pf.bitcount = 0;
header.pf.rmask = 0;
header.pf.gmask = 0;
header.pf.bmask = 0;
header.pf.amask = 0;
header.caps.caps1 = DDSCAPS_TEXTURE;
header.caps.caps2 = 0;
header.caps.caps3 = 0;
header.caps.caps4 = 0;
header.notused = 0;
fwrite(&header, sizeof(DDSHeader), 1, fp);
fwrite(h_result, compressedSize, 1, fp);
fclose(fp);
// Make sure the generated image matches the reference image (regression check)
shrLog("\nComparing against Host/C++ computation...\n");
const char* reference_image_path = shrFindFilePath(refimage_filename, argv[0]);
oclCheckError(reference_image_path != NULL, shrTRUE);
// read in the reference image from file
#ifdef WIN32
fopen_s(&fp, reference_image_path, "rb");
#else
fp = fopen(reference_image_path, "rb");
#endif
oclCheckError(fp != NULL, shrTRUE);
fseek(fp, sizeof(DDSHeader), SEEK_SET);
uint referenceSize = (width / 4) * (height / 4) * 8;
uint * reference = (uint *)malloc(referenceSize);
fread(reference, referenceSize, 1, fp);
fclose(fp);
// compare the reference image data to the sample/generated image
float rms = 0;
for (uint y = 0; y < height; y += 4)
{
for (uint x = 0; x < width; x += 4)
{
// binary comparison of data
uint referenceBlockIdx = ((y/4) * (width/4) + (x/4));
uint resultBlockIdx = ((y/4) * (width/4) + (x/4));
int cmp = compareBlock(((BlockDXT1 *)h_result) + resultBlockIdx, ((BlockDXT1 *)reference) + referenceBlockIdx);
// log deviations, if any
if (cmp != 0.0f)
{
compareBlock(((BlockDXT1 *)h_result) + resultBlockIdx, ((BlockDXT1 *)reference) + referenceBlockIdx);
shrLog("Deviation at (%d, %d):\t%f rms\n", x/4, y/4, float(cmp)/16/3);
}
rms += cmp;
}
}
rms /= width * height * 3;
shrLog("RMS(reference, result) = %f\n\n", rms);
// Free OpenCL resources
oclDeleteMemObjs(cmMemObjs, 3);
clReleaseMemObject(cmAlphaTable4);
clReleaseMemObject(cmProds4);
clReleaseMemObject(cmAlphaTable3);
clReleaseMemObject(cmProds3);
clReleaseKernel(ckKernel);
clReleaseProgram(cpProgram);
clReleaseCommandQueue(cqCommandQueue);
clReleaseContext(cxGPUContext);
// Free host memory
free(source);
free(h_img);
// finish
shrQAFinishExit(argc, (const char **)argv, (rms <= ERROR_THRESHOLD) ? QA_PASSED : QA_FAILED);
}
////////////////////////////////////////////////////////////////////////////////
// Main program
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
cl_platform_id cpPlatform; //OpenCL platform
cl_device_id cdDevice; //OpenCL device
cl_context cxGPUContext; //OpenCL context
cl_command_queue cqCommandQueue; //OpenCL command que
cl_mem d_Input, d_Output; //OpenCL memory buffer objects
cl_int ciErrNum;
float *h_Input, *h_OutputCPU, *h_OutputGPU;
const uint
imageW = 2048,
imageH = 2048,
stride = 2048;
const int dir = DCT_FORWARD;
shrQAStart(argc, argv);
int use_gpu = 0;
for(int i = 0; i < argc && argv; i++)
{
if(!argv[i])
continue;
if(strstr(argv[i], "cpu"))
use_gpu = 0;
else if(strstr(argv[i], "gpu"))
use_gpu = 1;
}
// set logfile name and start logs
shrSetLogFileName ("oclDCT8x8.txt");
shrLog("%s Starting...\n\n", argv[0]);
shrLog("Allocating and initializing host memory...\n");
h_Input = (float *)malloc(imageH * stride * sizeof(float));
h_OutputCPU = (float *)malloc(imageH * stride * sizeof(float));
h_OutputGPU = (float *)malloc(imageH * stride * sizeof(float));
srand(2009);
for(uint i = 0; i < imageH; i++)
for(uint j = 0; j < imageW; j++)
h_Input[i * stride + j] = (float)rand() / (float)RAND_MAX;
shrLog("Initializing OpenCL...\n");
//Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckError(ciErrNum, CL_SUCCESS);
//Get a GPU device
ciErrNum = clGetDeviceIDs(cpPlatform, use_gpu?CL_DEVICE_TYPE_GPU:CL_DEVICE_TYPE_CPU, 1, &cdDevice, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
//Create the context
cxGPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
//Create a command-queue
cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevice, CL_QUEUE_PROFILING_ENABLE, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("Initializing OpenCL DCT 8x8...\n");
initDCT8x8(cxGPUContext, cqCommandQueue, (const char **)argv);
shrLog("Creating OpenCL memory objects...\n");
d_Input = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, imageH * stride * sizeof(cl_float), h_Input, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
d_Output = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, imageH * stride * sizeof(cl_float), NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("Performing DCT8x8 of %u x %u image...\n\n", imageH, imageW);
//Just a single iteration or a warmup iteration
DCT8x8(
cqCommandQueue,
d_Output,
d_Input,
stride,
imageH,
imageW,
dir
);
#ifdef GPU_PROFILING
const int numIterations = 16;
cl_event startMark, endMark;
ciErrNum = clEnqueueMarker(cqCommandQueue, &startMark);
ciErrNum |= clFinish(cqCommandQueue);
shrCheckError(ciErrNum, CL_SUCCESS);
shrDeltaT(0);
for(int iter = 0; iter < numIterations; iter++)
DCT8x8(
NULL,
d_Output,
d_Input,
stride,
imageH,
imageW,
dir
);
ciErrNum = clEnqueueMarker(cqCommandQueue, &endMark);
ciErrNum |= clFinish(cqCommandQueue);
shrCheckError(ciErrNum, CL_SUCCESS);
//Calculate performance metrics by wallclock time
double gpuTime = shrDeltaT(0) / (double)numIterations;
shrLogEx(LOGBOTH | MASTER, 0, "oclDCT8x8, Throughput = %.4f MPixels/s, Time = %.5f s, Size = %u Pixels, NumDevsUsed = %i, Workgroup = %u\n",
(1.0e-6 * (double)(imageW * imageH)/ gpuTime), gpuTime, (imageW * imageH), 1, 0);
//Get profiler time
cl_ulong startTime = 0, endTime = 0;
ciErrNum = clGetEventProfilingInfo(startMark, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &startTime, NULL);
ciErrNum |= clGetEventProfilingInfo(endMark, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime, NULL);
shrCheckError(ciErrNum, CL_SUCCESS);
shrLog("\nOpenCL time: %.5f s\n\n", 1.0e-9 * ((double)endTime - (double)startTime) / (double)numIterations);
#endif
shrLog("Reading back OpenCL results...\n");
ciErrNum = clEnqueueReadBuffer(cqCommandQueue, d_Output, CL_TRUE, 0, imageH * stride * sizeof(cl_float), h_OutputGPU, 0, NULL, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("Comparing against Host/C++ computation...\n");
DCT8x8CPU(h_OutputCPU, h_Input, stride, imageH, imageW, dir);
double sum = 0, delta = 0;
double L2norm;
for(uint i = 0; i < imageH; i++)
for(uint j = 0; j < imageW; j++){
sum += h_OutputCPU[i * stride + j] * h_OutputCPU[i * stride + j];
delta += (h_OutputGPU[i * stride + j] - h_OutputCPU[i * stride + j]) * (h_OutputGPU[i * stride + j] - h_OutputCPU[i * stride + j]);
}
L2norm = sqrt(delta / sum);
shrLog("Relative L2 norm: %.3e\n\n", L2norm);
shrLog("Shutting down...\n");
//Release kernels and program
closeDCT8x8();
//Release other OpenCL objects
ciErrNum = clReleaseMemObject(d_Output);
ciErrNum |= clReleaseMemObject(d_Input);
ciErrNum |= clReleaseCommandQueue(cqCommandQueue);
ciErrNum |= clReleaseContext(cxGPUContext);
oclCheckError(ciErrNum, CL_SUCCESS);
//Release host buffers
free(h_OutputGPU);
free(h_OutputCPU);
free(h_Input);
//Finish
shrQAFinishExit(argc, (const char **)argv, (L2norm < 1E-6) ? QA_PASSED : QA_FAILED);
}
int main(int argc, char **argv)
{
uchar *h_Data;
uint *h_HistogramCPU, *h_HistogramGPU;
uchar *d_Data;
uint *d_Histogram;
uint hTimer;
int PassFailFlag = 1;
uint byteCount = 64 * 1048576;
uint uiSizeMult = 1;
cudaDeviceProp deviceProp;
deviceProp.major = 0;
deviceProp.minor = 0;
int dev;
shrQAStart(argc, argv);
// set logfile name and start logs
shrSetLogFileName ("histogram.txt");
//Use command-line specified CUDA device, otherwise use device with highest Gflops/s
if( shrCheckCmdLineFlag(argc, (const char**)argv, "device") ) {
dev = cutilDeviceInit(argc, argv);
if (dev < 0) {
printf("No CUDA Capable Devices found, exiting...\n");
shrQAFinishExit(argc, (const char **)argv, QA_WAIVED);
}
} else {
cudaSetDevice( dev = cutGetMaxGflopsDeviceId() );
cutilSafeCall( cudaChooseDevice(&dev, &deviceProp) );
}
cutilSafeCall( cudaGetDeviceProperties(&deviceProp, dev) );
printf("CUDA device [%s] has %d Multi-Processors, Compute %d.%d\n",
deviceProp.name, deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
int version = deviceProp.major * 0x10 + deviceProp.minor;
if(version < 0x11)
{
printf("There is no device supporting a minimum of CUDA compute capability 1.1 for this SDK sample\n");
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_WAIVED);
}
cutilCheckError(cutCreateTimer(&hTimer));
// Optional Command-line multiplier to increase size of array to histogram
if (shrGetCmdLineArgumentu(argc, (const char**)argv, "sizemult", &uiSizeMult))
{
uiSizeMult = CLAMP(uiSizeMult, 1, 10);
byteCount *= uiSizeMult;
}
shrLog("Initializing data...\n");
shrLog("...allocating CPU memory.\n");
h_Data = (uchar *)malloc(byteCount);
h_HistogramCPU = (uint *)malloc(HISTOGRAM256_BIN_COUNT * sizeof(uint));
h_HistogramGPU = (uint *)malloc(HISTOGRAM256_BIN_COUNT * sizeof(uint));
shrLog("...generating input data\n");
srand(2009);
for(uint i = 0; i < byteCount; i++)
h_Data[i] = rand() % 256;
shrLog("...allocating GPU memory and copying input data\n\n");
cutilSafeCall( cudaMalloc((void **)&d_Data, byteCount ) );
cutilSafeCall( cudaMalloc((void **)&d_Histogram, HISTOGRAM256_BIN_COUNT * sizeof(uint) ) );
cutilSafeCall( cudaMemcpy(d_Data, h_Data, byteCount, cudaMemcpyHostToDevice) );
{
shrLog("Starting up 64-bin histogram...\n\n");
initHistogram64();
shrLog("Running 64-bin GPU histogram for %u bytes (%u runs)...\n\n", byteCount, numRuns);
for(int iter = -1; iter < numRuns; iter++){
//iter == -1 -- warmup iteration
if(iter == 0){
cutilSafeCall( cutilDeviceSynchronize() );
cutilCheckError( cutResetTimer(hTimer) );
cutilCheckError( cutStartTimer(hTimer) );
}
histogram64(d_Histogram, d_Data, byteCount);
}
cutilSafeCall( cutilDeviceSynchronize() );
cutilCheckError( cutStopTimer(hTimer));
double dAvgSecs = 1.0e-3 * (double)cutGetTimerValue(hTimer) / (double)numRuns;
shrLog("histogram64() time (average) : %.5f sec, %.4f MB/sec\n\n", dAvgSecs, ((double)byteCount * 1.0e-6) / dAvgSecs);
shrLogEx(LOGBOTH | MASTER, 0, "histogram64, Throughput = %.4f MB/s, Time = %.5f s, Size = %u Bytes, NumDevsUsed = %u, Workgroup = %u\n",
(1.0e-6 * (double)byteCount / dAvgSecs), dAvgSecs, byteCount, 1, HISTOGRAM64_THREADBLOCK_SIZE);
shrLog("\nValidating GPU results...\n");
shrLog(" ...reading back GPU results\n");
cutilSafeCall( cudaMemcpy(h_HistogramGPU, d_Histogram, HISTOGRAM64_BIN_COUNT * sizeof(uint), cudaMemcpyDeviceToHost) );
shrLog(" ...histogram64CPU()\n");
histogram64CPU(
h_HistogramCPU,
h_Data,
byteCount
);
shrLog(" ...comparing the results...\n");
for(uint i = 0; i < HISTOGRAM64_BIN_COUNT; i++)
if(h_HistogramGPU[i] != h_HistogramCPU[i]) PassFailFlag = 0;
shrLog(PassFailFlag ? " ...64-bin histograms match\n\n" : " ***64-bin histograms do not match!!!***\n\n" );
shrLog("Shutting down 64-bin histogram...\n\n\n");
closeHistogram64();
}
{
shrLog("Initializing 256-bin histogram...\n");
initHistogram256();
shrLog("Running 256-bin GPU histogram for %u bytes (%u runs)...\n\n", byteCount, numRuns);
for(int iter = -1; iter < numRuns; iter++){
//iter == -1 -- warmup iteration
if(iter == 0){
cutilSafeCall( cutilDeviceSynchronize() );
cutilCheckError( cutResetTimer(hTimer) );
cutilCheckError( cutStartTimer(hTimer) );
}
histogram256(d_Histogram, d_Data, byteCount);
}
cutilSafeCall( cutilDeviceSynchronize() );
cutilCheckError( cutStopTimer(hTimer));
double dAvgSecs = 1.0e-3 * (double)cutGetTimerValue(hTimer) / (double)numRuns;
shrLog("histogram256() time (average) : %.5f sec, %.4f MB/sec\n\n", dAvgSecs, ((double)byteCount * 1.0e-6) / dAvgSecs);
shrLogEx(LOGBOTH | MASTER, 0, "histogram256, Throughput = %.4f MB/s, Time = %.5f s, Size = %u Bytes, NumDevsUsed = %u, Workgroup = %u\n",
(1.0e-6 * (double)byteCount / dAvgSecs), dAvgSecs, byteCount, 1, HISTOGRAM256_THREADBLOCK_SIZE);
shrLog("\nValidating GPU results...\n");
shrLog(" ...reading back GPU results\n");
cutilSafeCall( cudaMemcpy(h_HistogramGPU, d_Histogram, HISTOGRAM256_BIN_COUNT * sizeof(uint), cudaMemcpyDeviceToHost) );
shrLog(" ...histogram256CPU()\n");
histogram256CPU(
h_HistogramCPU,
h_Data,
byteCount
);
shrLog(" ...comparing the results\n");
for(uint i = 0; i < HISTOGRAM256_BIN_COUNT; i++)
if(h_HistogramGPU[i] != h_HistogramCPU[i]) PassFailFlag = 0;
shrLog(PassFailFlag ? " ...256-bin histograms match\n\n" : " ***256-bin histograms do not match!!!***\n\n" );
shrLog("Shutting down 256-bin histogram...\n\n\n");
closeHistogram256();
}
shrLog("Shutting down...\n");
cutilCheckError(cutDeleteTimer(hTimer));
cutilSafeCall( cudaFree(d_Histogram) );
cutilSafeCall( cudaFree(d_Data) );
free(h_HistogramGPU);
free(h_HistogramCPU);
free(h_Data);
cutilDeviceReset();
shrLog("%s - Test Summary\n", sSDKsample);
// pass or fail (for both 64 bit and 256 bit histograms)
shrQAFinishExit(argc, (const char **)argv, (PassFailFlag ? QA_PASSED : QA_FAILED));
}
int main(int argc, char** argv)
{
pArgc = &argc;
pArgv = argv;
shrQAStart(argc, argv);
if (argc > 1) {
if (cutCheckCmdLineFlag(argc, (const char **)argv, "help")) {
printHelp();
}
if (cutCheckCmdLineFlag(argc, (const char **)argv, "qatest") ||
cutCheckCmdLineFlag(argc, (const char **)argv, "noprompt"))
{
g_bQAReadback = true;
fpsLimit = frameCheckNumber;
}
if (cutCheckCmdLineFlag(argc, (const char **)argv, "glverify"))
{
g_bOpenGLQA = true;
fpsLimit = frameCheckNumber;
}
if (cutCheckCmdLineFlag(argc, (const char **)argv, "fbo")) {
g_bFBODisplay = true;
fpsLimit = frameCheckNumber;
}
}
if (g_bQAReadback)
{
runAutoTest(argc, argv);
}
else
{
if ( cutCheckCmdLineFlag(argc, (const char **)argv, "device")) {
printf(" This SDK does not explicitly support -device=n when running with OpenGL.\n");
printf(" When specifying -device=n (n=0,1,2,....) the sample must not use OpenGL.\n");
printf(" See details below to run without OpenGL:\n\n");
printf(" > %s -device=n -qatest\n\n", argv[0]);
printf("exiting...\n");
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
// First initialize OpenGL context, so we can properly set the GL for CUDA.
// This is necessary in order to achieve optimal performance with OpenGL/CUDA interop.
initGL( &argc, argv );
//cudaGLSetGLDevice (cutGetMaxGflopsDeviceId() );
int dev = findCapableDevice(argc, argv);
if( dev != -1 ) {
cudaGLSetGLDevice( dev );
} else {
shrQAFinishExit2(g_bQAReadback, *pArgc, (const char **)pArgv, QA_PASSED);
}
cutilCheckError(cutCreateTimer(&timer));
cutilCheckError(cutResetTimer(timer));
glutDisplayFunc(display);
glutKeyboardFunc(keyboard);
glutReshapeFunc(reshape);
if (g_bOpenGLQA) {
loadDefaultImage( argc, argv );
}
if (argc > 1) {
char *filename;
if (cutGetCmdLineArgumentstr(argc, (const char **)argv, "file", &filename)) {
initializeData(filename, argc, argv);
}
} else {
loadDefaultImage( argc, argv );
}
// If code is not printing the USage, then we execute this path.
if (!bQuit) {
if (g_bOpenGLQA) {
g_CheckRender = new CheckBackBuffer(wWidth, wHeight, 4);
g_CheckRender->setPixelFormat(GL_BGRA);
g_CheckRender->setExecPath(argv[0]);
g_CheckRender->EnableQAReadback(true);
}
printf("I: display Image (no filtering)\n");
printf("T: display Sobel Edge Detection (Using Texture)\n");
printf("S: display Sobel Edge Detection (Using SMEM+Texture)\n");
printf("Use the '-' and '=' keys to change the brightness.\n");
printf("b: switch block filter operation (mean/Sobel)\n");
printf("p: switch point filter operation (threshold on/off)\n");
fflush(stdout);
atexit(cleanup);
glutTimerFunc(REFRESH_DELAY, timerEvent,0);
glutMainLoop();
}
}
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_PASSED);
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
shrQAStart(argc, argv);
// start logs
shrSetLogFileName ("oclSimpleMultiGPU.txt");
shrLog("%s Starting, Array = %u float values...\n\n", argv[0], DATA_N);
// OpenCL
cl_platform_id cpPlatform;
cl_uint ciDeviceCount;
cl_device_id* cdDevices;
cl_context cxGPUContext;
cl_device_id cdDevice; // GPU device
int deviceNr[MAX_GPU_COUNT];
cl_command_queue commandQueue[MAX_GPU_COUNT];
cl_mem d_Data[MAX_GPU_COUNT];
cl_mem d_Result[MAX_GPU_COUNT];
cl_program cpProgram;
cl_kernel reduceKernel[MAX_GPU_COUNT];
cl_event GPUDone[MAX_GPU_COUNT];
cl_event GPUExecution[MAX_GPU_COUNT];
size_t programLength;
cl_int ciErrNum;
char cDeviceName [256];
cl_mem h_DataBuffer;
// Vars for reduction results
float h_SumGPU[MAX_GPU_COUNT * ACCUM_N];
float *h_Data;
double sumGPU;
double sumCPU, dRelError;
// allocate and init host buffer with with some random generated input data
h_Data = (float *)malloc(DATA_N * sizeof(float));
shrFillArray(h_Data, DATA_N);
// start timer & logs
shrLog("Setting up OpenCL on the Host...\n\n");
shrDeltaT(1);
// Annotate profiling state
#ifdef GPU_PROFILING
shrLog("OpenCL Profiling is enabled...\n\n");
#endif
//Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clGetPlatformID...\n");
//Get the devices
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 0, NULL, &ciDeviceCount);
oclCheckError(ciErrNum, CL_SUCCESS);
cdDevices = (cl_device_id *)malloc(ciDeviceCount * sizeof(cl_device_id) );
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, ciDeviceCount, cdDevices, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clGetDeviceIDs...\n");
//Create the context
cxGPUContext = clCreateContext(0, ciDeviceCount, cdDevices, NULL, NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateContext...\n");
// Set up command queue(s) for GPU's specified on the command line or all GPU's
if(shrCheckCmdLineFlag(argc, (const char **)argv, "device"))
{
// User specified GPUs
int ciMaxDeviceID = ciDeviceCount-1;
ciDeviceCount = 0;
char* deviceList;
char* deviceStr;
char* next_token;
shrGetCmdLineArgumentstr(argc, (const char **)argv, "device", &deviceList);
#ifdef WIN32
deviceStr = strtok_s (deviceList," ,.-", &next_token);
#else
deviceStr = strtok (deviceList," ,.-");
#endif
// Create command queues for all Requested GPU's
while(deviceStr != NULL)
{
// get & log device index # and name
deviceNr[ciDeviceCount] = atoi(deviceStr);
if( deviceNr[ciDeviceCount] > ciMaxDeviceID ) {
shrLog(" Invalid user specified device ID: %d\n", deviceNr[ciDeviceCount]);
return 1;
}
cdDevice = oclGetDev(cxGPUContext, deviceNr[ciDeviceCount]);
ciErrNum = clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(cDeviceName), cDeviceName, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog(" Device %i: %s\n\n", deviceNr[ciDeviceCount], cDeviceName);
// create a command que
commandQueue[ciDeviceCount] = clCreateCommandQueue(cxGPUContext, cdDevice, CL_QUEUE_PROFILING_ENABLE, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateCommandQueue\n");
++ciDeviceCount;
#ifdef WIN32
deviceStr = strtok_s (NULL," ,.-", &next_token);
#else
deviceStr = strtok (NULL," ,.-");
#endif
}
free(deviceList);
}
else
{
// Find out how many GPU's to compute on all available GPUs
size_t nDeviceBytes;
ciErrNum = clGetContextInfo(cxGPUContext, CL_CONTEXT_DEVICES, 0, NULL, &nDeviceBytes);
oclCheckError(ciErrNum, CL_SUCCESS);
ciDeviceCount = (cl_uint)nDeviceBytes/sizeof(cl_device_id);
for(unsigned int i = 0; i < ciDeviceCount; ++i )
{
// get & log device index # and name
deviceNr[i] = i;
cdDevice = oclGetDev(cxGPUContext, i);
ciErrNum = clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(cDeviceName), cDeviceName, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog(" Device %i: %s\n", i, cDeviceName);
// create a command que
commandQueue[i] = clCreateCommandQueue(cxGPUContext, cdDevice, CL_QUEUE_PROFILING_ENABLE, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateCommandQueue\n\n");
}
}
// Load the OpenCL source code from the .cl file
const char* source_path = shrFindFilePath("simpleMultiGPU.cl", argv[0]);
char *source = oclLoadProgSource(source_path, "", &programLength);
oclCheckError(source != NULL, shrTRUE);
shrLog("oclLoadProgSource\n");
// Create the program for all GPUs in the context
cpProgram = clCreateProgramWithSource(cxGPUContext, 1, (const char **)&source, &programLength, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateProgramWithSource\n");
// build the program
ciErrNum = clBuildProgram(cpProgram, 0, NULL, "-cl-fast-relaxed-math", NULL, NULL);
if (ciErrNum != CL_SUCCESS)
{
// write out standard error, Build Log and PTX, then cleanup and exit
shrLogEx(LOGBOTH | ERRORMSG, ciErrNum, STDERROR);
oclLogBuildInfo(cpProgram, oclGetFirstDev(cxGPUContext));
oclLogPtx(cpProgram, oclGetFirstDev(cxGPUContext), "oclSimpleMultiGPU.ptx");
oclCheckError(ciErrNum, CL_SUCCESS);
}
shrLog("clBuildProgram\n");
// Create host buffer with page-locked memory
h_DataBuffer = clCreateBuffer(cxGPUContext, CL_MEM_COPY_HOST_PTR | CL_MEM_ALLOC_HOST_PTR,
DATA_N * sizeof(float), h_Data, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateBuffer (Page-locked Host)\n\n");
// Create buffers for each GPU, with data divided evenly among GPU's
int sizePerGPU = DATA_N / ciDeviceCount;
int workOffset[MAX_GPU_COUNT];
int workSize[MAX_GPU_COUNT];
workOffset[0] = 0;
for(unsigned int i = 0; i < ciDeviceCount; ++i )
{
workSize[i] = (i != (ciDeviceCount - 1)) ? sizePerGPU : (DATA_N - workOffset[i]);
// Input buffer
d_Data[i] = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, workSize[i] * sizeof(float), NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateBuffer (Input)\t\tDev %i\n", i);
// Copy data from host to device
ciErrNum = clEnqueueCopyBuffer(commandQueue[i], h_DataBuffer, d_Data[i], workOffset[i] * sizeof(float),
0, workSize[i] * sizeof(float), 0, NULL, NULL);
shrLog("clEnqueueCopyBuffer (Input)\tDev %i\n", i);
// Output buffer
d_Result[i] = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, ACCUM_N * sizeof(float), NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateBuffer (Output)\t\tDev %i\n", i);
// Create kernel
reduceKernel[i] = clCreateKernel(cpProgram, "reduce", &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clCreateKernel\t\t\tDev %i\n", i);
// Set the args values and check for errors
ciErrNum |= clSetKernelArg(reduceKernel[i], 0, sizeof(cl_mem), &d_Result[i]);
ciErrNum |= clSetKernelArg(reduceKernel[i], 1, sizeof(cl_mem), &d_Data[i]);
ciErrNum |= clSetKernelArg(reduceKernel[i], 2, sizeof(int), &workSize[i]);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("clSetKernelArg\t\t\tDev %i\n\n", i);
workOffset[i + 1] = workOffset[i] + workSize[i];
}
// Set # of work items in work group and total in 1 dimensional range
size_t localWorkSize[] = {THREAD_N};
size_t globalWorkSize[] = {ACCUM_N};
// Start timer and launch reduction kernel on each GPU, with data split between them
shrLog("Launching Kernels on GPU(s)...\n\n");
for(unsigned int i = 0; i < ciDeviceCount; i++)
{
ciErrNum = clEnqueueNDRangeKernel(commandQueue[i], reduceKernel[i], 1, 0, globalWorkSize, localWorkSize,
0, NULL, &GPUExecution[i]);
oclCheckError(ciErrNum, CL_SUCCESS);
}
// Copy result from device to host for each device
for(unsigned int i = 0; i < ciDeviceCount; i++)
{
ciErrNum = clEnqueueReadBuffer(commandQueue[i], d_Result[i], CL_FALSE, 0, ACCUM_N * sizeof(float),
h_SumGPU + i * ACCUM_N, 0, NULL, &GPUDone[i]);
oclCheckError(ciErrNum, CL_SUCCESS);
}
// Synchronize with the GPUs and do accumulated error check
clWaitForEvents(ciDeviceCount, GPUDone);
shrLog("clWaitForEvents complete...\n\n");
// Aggregate results for multiple GPU's and stop/log processing time
sumGPU = 0;
for(unsigned int i = 0; i < ciDeviceCount * ACCUM_N; i++)
{
sumGPU += h_SumGPU[i];
}
// Print Execution Times for each GPU
#ifdef GPU_PROFILING
shrLog("Profiling Information for GPU Processing:\n\n");
for(unsigned int i = 0; i < ciDeviceCount; i++)
{
cdDevice = oclGetDev(cxGPUContext, deviceNr[i]);
clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(cDeviceName), cDeviceName, NULL);
shrLog("Device %i : %s\n", deviceNr[i], cDeviceName);
shrLog(" Reduce Kernel : %.5f s\n", executionTime(GPUExecution[i]));
shrLog(" Copy Device->Host : %.5f s\n\n\n", executionTime(GPUDone[i]));
}
#endif
// Run the computation on the Host CPU and log processing time
shrLog("Launching Host/CPU C++ Computation...\n\n");
sumCPU = 0;
for(unsigned int i = 0; i < DATA_N; i++)
{
sumCPU += h_Data[i];
}
// Check GPU result against CPU result
dRelError = 100.0 * fabs(sumCPU - sumGPU) / fabs(sumCPU);
shrLog("Comparing against Host/C++ computation...\n");
shrLog(" GPU sum: %f\n CPU sum: %f\n", sumGPU, sumCPU);
shrLog(" Relative Error (100.0 * Error / Golden) = %f \n\n", dRelError);
// cleanup
free(source);
free(h_Data);
for(unsigned int i = 0; i < ciDeviceCount; ++i )
{
clReleaseKernel(reduceKernel[i]);
clReleaseCommandQueue(commandQueue[i]);
}
clReleaseProgram(cpProgram);
clReleaseContext(cxGPUContext);
// finish
shrQAFinishExit(argc, (const char **)argv, (dRelError < 1e-4) ? QA_PASSED : QA_FAILED);
}
// main function
//*****************************************************************************
int main(int argc, const char **argv)
{
cl_platform_id cpPlatform; // OpenCL platform
cl_uint nDevice; // OpenCL device count
cl_device_id* cdDevices; // OpenCL device list
cl_context cxGPUContext; // OpenCL context
cl_command_queue cqCommandQue[MAX_GPU_COUNT]; // OpenCL command que
cl_int ciErrNum = 1; // Error code var
shrQAStart(argc, (char **)argv);
// Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("clGetPlatformID...\n");
//Get all the devices
cl_uint uiNumDevices = 0; // Number of devices available
cl_uint uiTargetDevice = 0; // Default Device to compute on
cl_uint uiNumComputeUnits; // Number of compute units (SM's on NV GPU)
shrLog("Get the Device info and select Device...\n");
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 0, NULL, &uiNumDevices);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
cdDevices = (cl_device_id *)malloc(uiNumDevices * sizeof(cl_device_id) );
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, uiNumDevices, cdDevices, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
// Get command line device options and config accordingly
shrLog(" # of Devices Available = %u\n", uiNumDevices);
if(shrGetCmdLineArgumentu(argc, (const char**)argv, "device", &uiTargetDevice)== shrTRUE)
{
uiTargetDevice = CLAMP(uiTargetDevice, 0, (uiNumDevices - 1));
}
shrLog(" Using Device %u: ", uiTargetDevice);
oclPrintDevName(LOGBOTH, cdDevices[uiTargetDevice]);
ciErrNum = clGetDeviceInfo(cdDevices[uiTargetDevice], CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(uiNumComputeUnits), &uiNumComputeUnits, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("\n # of Compute Units = %u\n", uiNumComputeUnits);
shrSetLogFileName ("oclHiddenMarkovModel.txt");
shrLog("%s Starting...\n\n", argv[0]);
shrLog("Get platform...\n");
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("Get devices...\n");
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 0, NULL, &nDevice);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
cdDevices = (cl_device_id *)malloc(nDevice * sizeof(cl_device_id) );
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, nDevice, cdDevices, NULL);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("clCreateContext\n");
cxGPUContext = clCreateContext(0, nDevice, cdDevices, NULL, NULL, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
shrLog("clCreateCommandQueue\n");
int id_device;
if(shrGetCmdLineArgumenti(argc, argv, "device", &id_device)) // Set up command queue(s) for GPU specified on the command line
{
// create a command que
cqCommandQue[0] = clCreateCommandQueue(cxGPUContext, cdDevices[id_device], 0, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
oclPrintDevInfo(LOGBOTH, cdDevices[id_device]);
nDevice = 1;
}
else
{ // create command queues for all available devices
for (cl_uint i = 0; i < nDevice; i++)
{
cqCommandQue[i] = clCreateCommandQueue(cxGPUContext, cdDevices[i], CL_QUEUE_PROFILING_ENABLE, &ciErrNum);
oclCheckErrorEX(ciErrNum, CL_SUCCESS, NULL);
}
for (cl_uint i = 0; i < nDevice; i++) oclPrintDevInfo(LOGBOTH, cdDevices[i]);
}
shrLog("\nUsing %d GPU(s)...\n\n", nDevice);
int wgSize;
if (!shrGetCmdLineArgumenti(argc, argv, "work-group-size", &wgSize))
{
wgSize = 256;
}
shrLog("Init Hidden Markov Model parameters\n");
int nState = 256*16; // number of states, must be a multiple of 256
int nEmit = 128; // number of possible observations
float *initProb = (float*)malloc(sizeof(float)*nState); // initial probability
float *mtState = (float*)malloc(sizeof(float)*nState*nState); // state transition matrix
float *mtEmit = (float*)malloc(sizeof(float)*nEmit*nState); // emission matrix
initHMM(initProb, mtState, mtEmit, nState, nEmit);
// define observational sequence
int nObs = 100; // size of observational sequence
int **obs = (int**)malloc(nDevice*sizeof(int*));
int **viterbiPathCPU = (int**)malloc(nDevice*sizeof(int*));
int **viterbiPathGPU = (int**)malloc(nDevice*sizeof(int*));
float *viterbiProbCPU = (float*)malloc(nDevice*sizeof(float));
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
obs[iDevice] = (int*)malloc(sizeof(int)*nObs);
for (int i = 0; i < nObs; i++)
obs[iDevice][i] = i % 15;
viterbiPathCPU[iDevice] = (int*)malloc(sizeof(int)*nObs);
viterbiPathGPU[iDevice] = (int*)malloc(sizeof(int)*nObs);
}
shrLog("# of states = %d\n# of possible observations = %d \nSize of observational sequence = %d\n\n",
nState, nEmit, nObs);
shrLog("Compute Viterbi path on GPU\n\n");
HMM **oHmm = (HMM**)malloc(nDevice*sizeof(HMM*));
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
oHmm[iDevice] = new HMM(cxGPUContext, cqCommandQue[iDevice], initProb, mtState, mtEmit, nState, nEmit, nObs, argv[0], wgSize);
}
cl_mem *vProb = (cl_mem*)malloc(sizeof(cl_mem)*nDevice);
cl_mem *vPath = (cl_mem*)malloc(sizeof(cl_mem)*nDevice);
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
vProb[iDevice] = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, sizeof(float), NULL, &ciErrNum);
vPath[iDevice] = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, sizeof(int)*nObs, NULL, &ciErrNum);
}
#ifdef GPU_PROFILING
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
clFinish(cqCommandQue[iDevice]);;
}
shrDeltaT(1);
#endif
size_t szWorkGroup;
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
szWorkGroup = oHmm[iDevice]->ViterbiSearch(vProb[iDevice], vPath[iDevice], obs[iDevice]);
}
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
clFinish(cqCommandQue[iDevice]);
}
#ifdef GPU_PROFILING
double dElapsedTime = shrDeltaT(1);
shrLogEx(LOGBOTH | MASTER, 0, "oclHiddenMarkovModel, Throughput = %.4f GB/s, Time = %.5f s, Size = %u items, NumDevsUsed = %u, Workgroup = %u\n",
(1.0e-9 * 2.0 * sizeof(float) * nDevice * nState * nState * (nObs-1))/dElapsedTime, dElapsedTime, (nDevice * nState * nObs), nDevice, szWorkGroup);
#endif
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
ciErrNum = clEnqueueReadBuffer(cqCommandQue[iDevice], vPath[iDevice], CL_TRUE, 0, sizeof(int)*nObs, viterbiPathGPU[iDevice], 0, NULL, NULL);
}
shrLog("\nCompute Viterbi path on CPU\n");
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
ciErrNum = ViterbiCPU(viterbiProbCPU[iDevice], viterbiPathCPU[iDevice], obs[iDevice], nObs, initProb, mtState, nState, mtEmit);
}
if (!ciErrNum)
{
shrEXIT(argc, argv);
}
bool pass = true;
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
for (int i = 0; i < nObs; i++)
{
if (viterbiPathCPU[iDevice][i] != viterbiPathGPU[iDevice][i])
{
pass = false;
break;
}
}
}
// NOTE: Most properly this should be done at any of the exit points above, but it is omitted elsewhere for clarity.
shrLog("Release CPU buffers and OpenCL objects...\n");
free(initProb);
free(mtState);
free(mtEmit);
for (cl_uint iDevice = 0; iDevice < nDevice; iDevice++)
{
free(obs[iDevice]);
free(viterbiPathCPU[iDevice]);
free(viterbiPathGPU[iDevice]);
delete oHmm[iDevice];
clReleaseCommandQueue(cqCommandQue[iDevice]);
}
free(obs);
free(viterbiPathCPU);
free(viterbiPathGPU);
free(viterbiProbCPU);
free(cdDevices);
free(oHmm);
clReleaseContext(cxGPUContext);
// finish
shrQAFinishExit(argc, (const char **)argv, pass ? QA_PASSED : QA_FAILED);
shrEXIT(argc, argv);
}
////////////////////////////////////////////////////////////////////////////////
// Test driver
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
cl_platform_id cpPlatform;
cl_device_id cdDevice;
cl_context cxGPUContext; //OpenCL context
cl_command_queue cqCommandQueue; //OpenCL command queue
cl_mem c_Kernel, d_Input, d_Buffer, d_Output; //OpenCL memory buffer objects
cl_float *h_Kernel, *h_Input, *h_Buffer, *h_OutputCPU, *h_OutputGPU;
cl_int ciErrNum;
const unsigned int imageW = 3072;
const unsigned int imageH = 3072;
shrQAStart(argc, argv);
// set logfile name and start logs
shrSetLogFileName ("oclConvolutionSeparable.txt");
shrLog("%s Starting...\n\n", argv[0]);
shrLog("Allocating and initializing host memory...\n");
h_Kernel = (cl_float *)malloc(KERNEL_LENGTH * sizeof(cl_float));
h_Input = (cl_float *)malloc(imageW * imageH * sizeof(cl_float));
h_Buffer = (cl_float *)malloc(imageW * imageH * sizeof(cl_float));
h_OutputCPU = (cl_float *)malloc(imageW * imageH * sizeof(cl_float));
h_OutputGPU = (cl_float *)malloc(imageW * imageH * sizeof(cl_float));
srand(2009);
for(unsigned int i = 0; i < KERNEL_LENGTH; i++)
h_Kernel[i] = (cl_float)(rand() % 16);
for(unsigned int i = 0; i < imageW * imageH; i++)
h_Input[i] = (cl_float)(rand() % 16);
shrLog("Initializing OpenCL...\n");
//Get the NVIDIA platform
ciErrNum = oclGetPlatformID(&cpPlatform);
oclCheckError(ciErrNum, CL_SUCCESS);
//Get the devices
ciErrNum = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU, 1, &cdDevice, NULL);
//Create the context
cxGPUContext = clCreateContext(0, 1, &cdDevice, NULL, NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
//Create a command-queue
cqCommandQueue = clCreateCommandQueue(cxGPUContext, cdDevice, CL_QUEUE_PROFILING_ENABLE, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("Initializing OpenCL separable convolution...\n");
initConvolutionSeparable(cxGPUContext, cqCommandQueue, (const char **)argv);
shrLog("Creating OpenCL memory objects...\n");
c_Kernel = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, KERNEL_LENGTH * sizeof(cl_float), h_Kernel, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
d_Input = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, imageW * imageH * sizeof(cl_float), h_Input, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
d_Buffer = clCreateBuffer(cxGPUContext, CL_MEM_READ_WRITE, imageW * imageH * sizeof(cl_float), NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
d_Output = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, imageW * imageH * sizeof(cl_float), NULL, &ciErrNum);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("Applying separable convolution to %u x %u image...\n\n", imageW, imageH);
//Just a single run or a warmup iteration
convolutionRows(
NULL,
d_Buffer,
d_Input,
c_Kernel,
imageW,
imageH
);
convolutionColumns(
NULL,
d_Output,
d_Buffer,
c_Kernel,
imageW,
imageH
);
#ifdef GPU_PROFILING
const int numIterations = 16;
cl_event startMark, endMark;
ciErrNum = clEnqueueMarker(cqCommandQueue, &startMark);
ciErrNum |= clFinish(cqCommandQueue);
shrCheckError(ciErrNum, CL_SUCCESS);
shrDeltaT(0);
for(int iter = 0; iter < numIterations; iter++){
convolutionRows(
cqCommandQueue,
d_Buffer,
d_Input,
c_Kernel,
imageW,
imageH
);
convolutionColumns(
cqCommandQueue,
d_Output,
d_Buffer,
c_Kernel,
imageW,
imageH
);
}
ciErrNum = clEnqueueMarker(cqCommandQueue, &endMark);
ciErrNum |= clFinish(cqCommandQueue);
shrCheckError(ciErrNum, CL_SUCCESS);
//Calculate performance metrics by wallclock time
double gpuTime = shrDeltaT(0) / (double)numIterations;
shrLogEx(LOGBOTH | MASTER, 0, "oclConvolutionSeparable, Throughput = %.4f MPixels/s, Time = %.5f s, Size = %u Pixels, NumDevsUsed = %i, Workgroup = %u\n",
(1.0e-6 * (double)(imageW * imageH)/ gpuTime), gpuTime, (imageW * imageH), 1, 0);
//Get OpenCL profiler info
cl_ulong startTime = 0, endTime = 0;
ciErrNum = clGetEventProfilingInfo(startMark, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &startTime, NULL);
ciErrNum |= clGetEventProfilingInfo(endMark, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &endTime, NULL);
shrCheckError(ciErrNum, CL_SUCCESS);
shrLog("\nOpenCL time: %.5f s\n\n", 1.0e-9 * ((double)endTime - (double)startTime)/ (double)numIterations);
#endif
shrLog("Reading back OpenCL results...\n\n");
ciErrNum = clEnqueueReadBuffer(cqCommandQueue, d_Output, CL_TRUE, 0, imageW * imageH * sizeof(cl_float), h_OutputGPU, 0, NULL, NULL);
oclCheckError(ciErrNum, CL_SUCCESS);
shrLog("Comparing against Host/C++ computation...\n");
convolutionRowHost(h_Buffer, h_Input, h_Kernel, imageW, imageH, KERNEL_RADIUS);
convolutionColumnHost(h_OutputCPU, h_Buffer, h_Kernel, imageW, imageH, KERNEL_RADIUS);
double sum = 0, delta = 0;
double L2norm;
for(unsigned int i = 0; i < imageW * imageH; i++){
delta += (h_OutputCPU[i] - h_OutputGPU[i]) * (h_OutputCPU[i] - h_OutputGPU[i]);
sum += h_OutputCPU[i] * h_OutputCPU[i];
}
L2norm = sqrt(delta / sum);
shrLog("Relative L2 norm: %.3e\n\n", L2norm);
// cleanup
closeConvolutionSeparable();
ciErrNum = clReleaseMemObject(d_Output);
ciErrNum |= clReleaseMemObject(d_Buffer);
ciErrNum |= clReleaseMemObject(d_Input);
ciErrNum |= clReleaseMemObject(c_Kernel);
ciErrNum |= clReleaseCommandQueue(cqCommandQueue);
ciErrNum |= clReleaseContext(cxGPUContext);
oclCheckError(ciErrNum, CL_SUCCESS);
free(h_OutputGPU);
free(h_OutputCPU);
free(h_Buffer);
free(h_Input);
free(h_Kernel);
// finish
shrQAFinishExit(argc, (const char **)argv, (L2norm < 1e-6) ? QA_PASSED : QA_FAILED);
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int main( int argc, char** argv)
{
shrQAStart(argc, argv);
// start logs
shrSetLogFileName ("bilateralFilter.txt");
shrLog("%s Starting...\n\n", argv[0]);
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
cutGetCmdLineArgumenti( argc, (const char**) argv, "threads", &nthreads );
cutGetCmdLineArgumenti( argc, (const char**) argv, "radius", &filter_radius);
// load image to process
loadImageData(argc, argv);
if (cutCheckCmdLineFlag(argc, (const char **)argv, "qatest") ||
cutCheckCmdLineFlag(argc, (const char **)argv, "noprompt"))
{
// Running CUDA kernel (bilateralFilter) without visualization (QA Testing/Verification)
runAutoTest(argc, argv);
shrQAFinishExit(argc, (const char **)argv, (g_TotalErrors == 0 ? QA_PASSED : QA_FAILED));
}
else if (cutCheckCmdLineFlag(argc, (const char **)argv, "benchmark"))
{
// Running CUDA kernel (bilateralFilter) in Benchmarking Mode
runBenchmark(argc, argv);
shrQAFinishExit(argc, (const char **)argv, (g_TotalErrors == 0 ? QA_PASSED : QA_FAILED));
}
else
{
// Running CUDA kernel (bilateralFilter) in CUDA + OpenGL Visualization Mode
if ( cutCheckCmdLineFlag(argc, (const char **)argv, "device")) {
printf("[%s]\n", argv[0]);
printf(" Does not explicitly support -device=n in OpenGL mode\n");
printf(" To use -device=n, the sample must be running w/o OpenGL\n\n");
printf(" > %s -device=n -qatest\n", argv[0]);
printf("exiting...\n");
exit(0);
}
// First initialize OpenGL context, so we can properly set the GL for CUDA.
// This is necessary in order to achieve optimal performance with OpenGL/CUDA interop.
initGL( argc, argv );
if ( cutCheckCmdLineFlag(argc, (const char **)argv, "device")) {
cutilGLDeviceInit(argc, argv);
} else {
cudaGLSetGLDevice (cutGetMaxGflopsDeviceId() );
}
initCuda();
initOpenGL();
}
atexit(cleanup);
printf("Running Standard Demonstration with GLUT loop...\n\n");
printf("Press '+' and '-' to change number of iterations\n"
"Press LEFT and RIGHT change euclidean delta\n"
"Press UP and DOWN to change gaussian delta\n"
"Press '1' to show original image\n"
"Press '2' to show result\n\n");
glutMainLoop();
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, (g_TotalErrors == 0 ? QA_PASSED : QA_FAILED));
}
//-----------------------------------------------------------------------------
// Program main
//-----------------------------------------------------------------------------
int main(int argc, char** argv)
{
pArgc = &argc;
pArgv = argv;
// start logs
shrQAStart(argc, argv);
shrSetLogFileName ("oclSimpleD3D9Texture.txt");
shrLog("%s Starting...\n\n", argv[0]);
// process command line arguments
if (argc > 1)
{
bQATest = shrCheckCmdLineFlag(argc, (const char **)argv, "qatest");
bNoPrompt = shrCheckCmdLineFlag(argc, (const char **)argv, "noprompt");
}
//
// create window
//
// Register the window class
WNDCLASSEX wc = { sizeof(WNDCLASSEX), CS_CLASSDC, MsgProc, 0L, 0L,
GetModuleHandle(NULL), NULL, NULL, NULL, NULL,
"OpenCL/D3D9 Texture InterOP", NULL };
RegisterClassEx( &wc );
int xBorder = ::GetSystemMetrics(SM_CXSIZEFRAME);
int yMenu = ::GetSystemMetrics(SM_CYMENU);
int yBorder = ::GetSystemMetrics(SM_CYSIZEFRAME);
// Create the application's window (padding by window border for uniform BB sizes across OSs)
HWND hWnd = CreateWindow( wc.lpszClassName, "OpenCL/D3D9 Texture InterOP",
WS_OVERLAPPEDWINDOW, 0, 0, g_WindowWidth + 2*xBorder, g_WindowHeight+ 2*yBorder+yMenu,
NULL, NULL, wc.hInstance, NULL );
ShowWindow(hWnd, SW_SHOWDEFAULT);
UpdateWindow(hWnd);
// init fps timer
shrDeltaT (1);
// Initialize Direct3D
if( SUCCEEDED( InitD3D9(hWnd) ) &&
SUCCEEDED( InitCL(argc, (const char **)argv) ) &&
SUCCEEDED( InitTextures() ) )
{
if (!g_bDeviceLost)
{
RegisterD3D9ResourceWithCL();
}
}
//
// the main loop
//
while(false == g_bDone)
{
RunCL();
DrawScene();
//
// handle I/O
//
MSG msg;
ZeroMemory( &msg, sizeof(msg) );
while( msg.message!=WM_QUIT )
{
if( PeekMessage( &msg, NULL, 0U, 0U, PM_REMOVE ) )
{
TranslateMessage( &msg );
DispatchMessage( &msg );
}
else
{
RunCL();
DrawScene();
if(bQATest)
{
for(int count=0;count<g_iFrameToCompare;count++)
{
RunCL();
DrawScene();
}
const char *ref_image_path = "ref_oclSimpleD3D9Texture.ppm";
const char *cur_image_path = "oclSimpleD3D9Texture.ppm";
// Save a reference of our current test run image
CheckRenderD3D9::BackbufferToPPM(g_pD3DDevice,cur_image_path);
// compare to offical reference image, printing PASS or FAIL.
g_bPassed = CheckRenderD3D9::PPMvsPPM(cur_image_path,ref_image_path,argv[0],MAX_EPSILON, 0.15f);
PostQuitMessage(0);
g_bDone = true;
}
}
}
};
// Unregister windows class
UnregisterClass( wc.lpszClassName, wc.hInstance );
// Cleanup and leave
Cleanup (g_bPassed ? EXIT_SUCCESS : EXIT_FAILURE);
}
////////////////////////////////////////////////////////////////////////////////
// Test driver
////////////////////////////////////////////////////////////////////////////////
int main(int argc, char** argv){
uint *h_SrcKey, *h_SrcVal, *h_DstKey, *h_DstVal;
uint *d_SrcKey, *d_SrcVal, *d_BufKey, *d_BufVal, *d_DstKey, *d_DstVal;
uint hTimer;
const uint N = 4 * 1048576;
const uint DIR = 1;
const uint numValues = 65536;
shrQAStart(argc, argv);
printf("Allocating and initializing host arrays...\n\n");
cutCreateTimer(&hTimer);
h_SrcKey = (uint *)malloc(N * sizeof(uint));
h_SrcVal = (uint *)malloc(N * sizeof(uint));
h_DstKey = (uint *)malloc(N * sizeof(uint));
h_DstVal = (uint *)malloc(N * sizeof(uint));
srand(2009);
for(uint i = 0; i < N; i++)
h_SrcKey[i] = rand() % numValues;
fillValues(h_SrcVal, N);
printf("Allocating and initializing CUDA arrays...\n\n");
cutilSafeCall( cudaMalloc((void **)&d_DstKey, N * sizeof(uint)) );
cutilSafeCall( cudaMalloc((void **)&d_DstVal, N * sizeof(uint)) );
cutilSafeCall( cudaMalloc((void **)&d_BufKey, N * sizeof(uint)) );
cutilSafeCall( cudaMalloc((void **)&d_BufVal, N * sizeof(uint)) );
cutilSafeCall( cudaMalloc((void **)&d_SrcKey, N * sizeof(uint)) );
cutilSafeCall( cudaMalloc((void **)&d_SrcVal, N * sizeof(uint)) );
cutilSafeCall( cudaMemcpy(d_SrcKey, h_SrcKey, N * sizeof(uint), cudaMemcpyHostToDevice) );
cutilSafeCall( cudaMemcpy(d_SrcVal, h_SrcVal, N * sizeof(uint), cudaMemcpyHostToDevice) );
printf("Initializing GPU merge sort...\n");
initMergeSort();
printf("Running GPU merge sort...\n");
cutilSafeCall( cutilDeviceSynchronize() );
cutResetTimer(hTimer);
cutStartTimer(hTimer);
mergeSort(
d_DstKey,
d_DstVal,
d_BufKey,
d_BufVal,
d_SrcKey,
d_SrcVal,
N,
DIR
);
cutilSafeCall( cutilDeviceSynchronize() );
cutStopTimer(hTimer);
printf("Time: %f ms\n", cutGetTimerValue(hTimer));
printf("Reading back GPU merge sort results...\n");
cutilSafeCall( cudaMemcpy(h_DstKey, d_DstKey, N * sizeof(uint), cudaMemcpyDeviceToHost) );
cutilSafeCall( cudaMemcpy(h_DstVal, d_DstVal, N * sizeof(uint), cudaMemcpyDeviceToHost) );
printf("Inspecting the results...\n");
uint keysFlag = validateSortedKeys(
h_DstKey,
h_SrcKey,
1,
N,
numValues,
DIR
);
uint valuesFlag = validateSortedValues(
h_DstKey,
h_DstVal,
h_SrcKey,
1,
N
);
printf("Shutting down...\n");
closeMergeSort();
cutilCheckError( cutDeleteTimer(hTimer) );
cutilSafeCall( cudaFree(d_SrcVal) );
cutilSafeCall( cudaFree(d_SrcKey) );
cutilSafeCall( cudaFree(d_BufVal) );
cutilSafeCall( cudaFree(d_BufKey) );
cutilSafeCall( cudaFree(d_DstVal) );
cutilSafeCall( cudaFree(d_DstKey) );
free(h_DstVal);
free(h_DstKey);
free(h_SrcVal);
free(h_SrcKey);
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, (keysFlag && valuesFlag) ? QA_PASSED : QA_FAILED);
}
// Main program
//*****************************************************************************
int main(int argc, char** argv)
{
numParticles = NUM_PARTICLES;
uint gridDim = GRID_SIZE;
pArgc = &argc;
pArgv = argv;
shrQAStart(argc, argv);
// Start logs and timers
cExecutableName = argv[0];
shrSetLogFileName ("oclParticles.txt");
shrLog("%s Starting...\n\n", argv[0]);
// check command line flags and parameters
if (argc > 1)
{
shrGetCmdLineArgumenti(argc, (const char**)argv, "n", (int*)&numParticles);
shrGetCmdLineArgumenti(argc, (const char**)argv, "grid", (int*)&gridDim);
bQATest = shrCheckCmdLineFlag(argc, (const char**)argv, "qatest");
bNoPrompt = shrCheckCmdLineFlag(argc, (const char**)argv, "noprompt");
}
// Set and log grid size and particle count, after checking optional command-line inputs
gridSize.x = gridSize.y = gridSize.z = gridDim;
shrLog(" grid: %d x %d x %d = %d cells\n", gridSize.x, gridSize.y, gridSize.z, gridSize.x * gridSize.y * gridSize.z);
shrLog(" particles: %d\n\n", numParticles);
// initialize GLUT and GLEW
if(!bQATest)
{
InitGL(&argc, argv);
}
// initialize OpenCL
startupOpenCL(argc,(const char**)argv);
// init simulation parameters and objects
initParticleSystem(numParticles, gridSize);
initParams();
// Init timers
shrDeltaT(0); // timer 0 is for processing time measurements
shrDeltaT(1); // timer 1 is for fps measurement
// Start main GLUT rendering loop for processing and rendering,
// or otherwise run No-GL Q/A test sequence
if(!bQATest)
{
glutMainLoop();
}
else
{
TestNoGL();
}
// Normally unused return path
shrQAFinish(argc, (const char **)argv, QA_PASSED);
Cleanup(EXIT_SUCCESS);
}
///////////////////////////////////////////////////////////////////////////////
// Main program
///////////////////////////////////////////////////////////////////////////////
int main(int argc, char **argv)
{
// Start logs
shrQAStart(argc, argv);
// initialize the GPU, either identified by --device
// or by picking the device with highest flop rate.
int devID = findCudaDevice(argc, (const char **)argv);
// parsing the number of random numbers to generate
int rand_n = DEFAULT_RAND_N;
if( checkCmdLineFlag(argc, (const char**) argv, "count") )
{
rand_n = getCmdLineArgumentInt(argc, (const char**) argv, "count");
}
printf("Allocating data for %i samples...\n", rand_n);
// parsing the seed
int seed = DEFAULT_SEED;
if( checkCmdLineFlag(argc, (const char**) argv, "seed") )
{
seed = getCmdLineArgumentInt(argc, (const char**) argv, "seed");
}
printf("Seeding with %i ...\n", seed);
float *d_Rand;
checkCudaErrors( cudaMalloc((void **)&d_Rand, rand_n * sizeof(float)) );
curandGenerator_t prngGPU;
checkCurandErrors( curandCreateGenerator(&prngGPU, CURAND_RNG_PSEUDO_MTGP32) );
checkCurandErrors( curandSetPseudoRandomGeneratorSeed(prngGPU, seed) );
curandGenerator_t prngCPU;
checkCurandErrors( curandCreateGeneratorHost(&prngCPU, CURAND_RNG_PSEUDO_MTGP32) );
checkCurandErrors( curandSetPseudoRandomGeneratorSeed(prngCPU, seed) );
//
// Example 1: Compare random numbers generated on GPU and CPU
float *h_RandGPU = (float *)malloc(rand_n * sizeof(float));
printf("Generating random numbers on GPU...\n\n");
checkCurandErrors( curandGenerateUniform(prngGPU, (float*) d_Rand, rand_n) );
printf("\nReading back the results...\n");
checkCudaErrors( cudaMemcpy(h_RandGPU, d_Rand, rand_n * sizeof(float), cudaMemcpyDeviceToHost) );
float *h_RandCPU = (float *)malloc(rand_n * sizeof(float));
printf("Generating random numbers on CPU...\n\n");
checkCurandErrors( curandGenerateUniform(prngCPU, (float*) h_RandCPU, rand_n) );
printf("Comparing CPU/GPU random numbers...\n\n");
float L1norm = compareResults(rand_n, h_RandGPU, h_RandCPU);
//
// Example 2: Timing of random number generation on GPU
const int numIterations = 10;
int i;
StopWatchInterface *hTimer;
checkCudaErrors( cudaDeviceSynchronize() );
sdkCreateTimer(&hTimer);
sdkResetTimer(&hTimer);
sdkStartTimer(&hTimer);
for (i = 0; i < numIterations; i++)
{
checkCurandErrors( curandGenerateUniform(prngGPU, (float*) d_Rand, rand_n) );
}
checkCudaErrors( cudaDeviceSynchronize() );
sdkStopTimer(&hTimer);
double gpuTime = 1.0e-3 * sdkGetTimerValue(&hTimer)/(double)numIterations;
printf("MersenneTwister, Throughput = %.4f GNumbers/s, Time = %.5f s, Size = %u Numbers\n",
1.0e-9 * rand_n / gpuTime, gpuTime, rand_n);
printf("Shutting down...\n");
checkCurandErrors( curandDestroyGenerator(prngGPU) );
checkCurandErrors( curandDestroyGenerator(prngCPU) );
checkCudaErrors( cudaFree(d_Rand) );
sdkDeleteTimer( &hTimer);
free(h_RandGPU);
free(h_RandCPU);
cudaDeviceReset();
shrQAFinishExit(argc, (const char**)argv, (L1norm < 1e-6) ? QA_PASSED : QA_FAILED);
}
int main(int argc, char* argv[])
{
shrQAStart(argc, argv);
try
{
std::string sFilename;
char *filePath = findFilePath("Lena.pgm", argv[0]);
if (filePath) {
sFilename = filePath;
} else {
printf("Error unable to find Lena.pgm\n");
shrQAFinishExit(argc, (const char **)argv, QA_FAILED);
}
// Parse the command line arguments for proper configuration
parseCommandLineArguments(argc, argv);
printfNPPinfo(argc, argv);
if (g_bQATest == false && (g_nDevice == -1) && argc > 1) {
sFilename = argv[1];
}
// if we specify the filename at the command line, then we only test sFilename[0].
int file_errors = 0;
std::ifstream infile(sFilename.data(), std::ifstream::in);
if (infile.good()) {
std::cout << "boxFilterNPP opened: <" << sFilename.data() << "> successfully!" << std::endl;
file_errors = 0;
infile.close();
} else {
std::cout << "boxFilterNPP unable to open: <" << sFilename.data() << ">" << std::endl;
file_errors++;
infile.close();
}
if (file_errors > 0) {
shrQAFinish(argc, (const char **)argv, QA_FAILED);
exit(EXIT_FAILURE);
}
std::string sResultFilename = sFilename;
std::string::size_type dot = sResultFilename.rfind('.');
if (dot != std::string::npos) sResultFilename = sResultFilename.substr(0, dot);
sResultFilename += "_boxFilter.pgm";
if (argc >= 3 && !g_bQATest)
sResultFilename = argv[2];
// declare a host image object for an 8-bit grayscale image
npp::ImageCPU_8u_C1 oHostSrc;
// load gray-scale image from disk
npp::loadImage(sFilename, oHostSrc);
// declare a device image and copy construct from the host image,
// i.e. upload host to device
npp::ImageNPP_8u_C1 oDeviceSrc(oHostSrc);
// create struct with box-filter mask size
NppiSize oMaskSize = {5, 5};
// create struct with ROI size given the current mask
NppiSize oSizeROI = {oDeviceSrc.width() - oMaskSize.width + 1, oDeviceSrc.height() - oMaskSize.height + 1};
// allocate device image of appropriatedly reduced size
npp::ImageNPP_8u_C1 oDeviceDst(oSizeROI.width, oSizeROI.height);
// set anchor point inside the mask to (0, 0)
NppiPoint oAnchor = {0, 0};
// run box filter
NppStatus eStatusNPP;
eStatusNPP = nppiFilterBox_8u_C1R(oDeviceSrc.data(), oDeviceSrc.pitch(),
oDeviceDst.data(), oDeviceDst.pitch(),
oSizeROI, oMaskSize, oAnchor);
NPP_ASSERT(NPP_NO_ERROR == eStatusNPP);
// declare a host image for the result
npp::ImageCPU_8u_C1 oHostDst(oDeviceDst.size());
// and copy the device result data into it
oDeviceDst.copyTo(oHostDst.data(), oHostDst.pitch());
saveImage(sResultFilename, oHostDst);
std::cout << "Saved image: " << sResultFilename << std::endl;
shrQAFinish(argc, (const char **)argv, QA_PASSED);
exit(EXIT_SUCCESS);
}
catch (npp::Exception & rException)
{
std::cerr << "Program error! The following exception occurred: \n";
std::cerr << rException << std::endl;
std::cerr << "Aborting." << std::endl;
shrQAFinish(argc, (const char **)argv, QA_FAILED);
exit(EXIT_FAILURE);
}
catch (...)
{
std::cerr << "Program error! An unknow type of exception occurred. \n";
std::cerr << "Aborting." << std::endl;
shrQAFinish(argc, (const char **)argv, QA_FAILED);
exit(EXIT_FAILURE);
return -1;
}
return 0;
}
////////////////////////////////////////////////////////////////////////////////
// Program main
////////////////////////////////////////////////////////////////////////////////
int
main( int argc, char** argv)
{
pArgc = &argc;
pArgv = argv;
shrQAStart(argc, argv);
// start logs
shrSetLogFileName ("boxFilter.txt");
shrLog("%s Starting...\n\n", argv[0]);
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
if (argc > 1) {
cutGetCmdLineArgumenti( argc, (const char**) argv, "threads", &nthreads );
cutGetCmdLineArgumenti( argc, (const char**) argv, "radius", &filter_radius);
if (cutCheckCmdLineFlag(argc, (const char **)argv, "glverify"))
{
g_bOpenGLQA = true;
fpsLimit = frameCheckNumber;
}
if (cutCheckCmdLineFlag(argc, (const char **)argv, "fbo")) {
g_bFBODisplay = true;
}
}
// load image to process
loadImageData(argc, argv);
if (cutCheckCmdLineFlag(argc, (const char **)argv, "qatest") ||
cutCheckCmdLineFlag(argc, (const char **)argv, "noprompt"))
{
// Running CUDA kernel (boxFilter) without visualization (QA Testing/Verification)
runAutoTest(argc, argv);
shrQAFinishExit(argc, (const char **)argv, (g_TotalErrors == 0 ? QA_PASSED : QA_FAILED));
}
else if (cutCheckCmdLineFlag(argc, (const char **)argv, "benchmark"))
{
// Running CUDA kernels (boxfilter) in Benchmarking mode
runBenchmark(argc, argv);
shrQAFinishExit(argc, (const char **)argv, (g_TotalErrors == 0 ? QA_PASSED : QA_FAILED));
}
else
{
// Running CUDA kernels (boxFilter) with OpenGL visualization
if (g_bFBODisplay) shrLog("[FBO Display] ");
if (g_bOpenGLQA) shrLog("[OpenGL Readback Comparisons] ");
shrLog("\n");
if ( cutCheckCmdLineFlag(argc, (const char **)argv, "device")) {
printf(" This SDK does not explicitly support -device=n when running with OpenGL.\n");
printf(" When specifying -device=n (n=0,1,2,....) the sample must not use OpenGL.\n");
printf(" See details below to run without OpenGL:\n\n");
printf(" > %s -device=n -qatest\n\n", argv[0]);
printf("exiting...\n");
shrQAFinishExit(argc, (const char **)argv, QA_WAIVED);
}
// First initialize OpenGL context, so we can properly set the GL for CUDA.
// This is necessary in order to achieve optimal performance with OpenGL/CUDA interop.
initGL( &argc, argv );
int dev = findCapableDevice(argc, argv);
if( dev != -1 ) {
cudaGLSetGLDevice( dev );
} else {
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, QA_WAIVED);
}
// Now we can create a CUDA context and bind it to the OpenGL context
initCuda();
initGLResources();
if (g_bOpenGLQA) {
if (g_bFBODisplay) {
g_CheckRender = new CheckFBO(width, height, 4, g_FrameBufferObject);
} else {
g_CheckRender = new CheckBackBuffer(width, height, 4);
}
g_CheckRender->setPixelFormat(GL_RGBA);
g_CheckRender->setExecPath(argv[0]);
g_CheckRender->EnableQAReadback(true);
}
}
// sets the callback function so it will call cleanup upon exit
atexit(cleanup);
shrLog("Running Standard Demonstration with GLUT loop...\n\n");
shrLog("Press '+' and '-' to change filter width\n"
"Press ']' and '[' to change number of iterations\n\n");
// Main OpenGL loop that will run visualization for every vsync
glutMainLoop();
cutilDeviceReset();
shrQAFinishExit(argc, (const char **)argv, (g_TotalErrors == 0 ? QA_PASSED : QA_FAILED));
}
