{"id":1135,"date":"2012-02-28T16:49:18","date_gmt":"2012-02-28T21:49:18","guid":{"rendered":"http:\/\/codinggorilla.domemtech.com\/?p=1135"},"modified":"2012-03-03T10:09:28","modified_gmt":"2012-03-03T15:09:28","slug":"performance-comparison-of-cuda-opencl-and-c-amp","status":"publish","type":"post","link":"http:\/\/165.227.223.229\/index.php\/2012\/02\/28\/performance-comparison-of-cuda-opencl-and-c-amp\/","title":{"rendered":"Performance comparison of CUDA, OpenCL, and C++ AMP"},"content":{"rendered":"

Trying to get information of the underlying design of a GPGPU programming language environment and hardware can be difficult. \u00c2\u00a0Companies will not publish design information because they do not want you or other companies to copy the technology. \u00c2\u00a0But, sometimes you need to know details of a technology that are just not published in order to use it effectively. \u00c2\u00a0If they won’t tell you how the technology works, the only recourse to gain an understanding is experimentation [1, 2]. \u00c2\u00a0What is the performance of OpenCL, CUDA, and C++ AMP? \u00c2\u00a0What can we learn from this information?<\/p>\n

<\/p>\n

Introduction<\/h1>\n

Performance of OpenCL, CUDA, and C++ AMP is not easy to compare because there is no one-to-one correspondence between the APIs. \u00c2\u00a0For example, there is no direct analog of the OpenCL function clSetKernelArg within CUDA. \u00c2\u00a0In fact, clSetKernelArg can be called once and the kernel called multiple times. \u00c2\u00a0In CUDA, this is impossible because the kernel argument is set when the kernel is called: a kernel call is translated into internal functions cudaSetupArgument and cudaLaunch.<\/p>\n

Moreover, each GPGPU programming language environment performs initialization when a function of the API\u00c2\u00a0called, which is order dependent. For example, the first call to the CUDA API could be a cudaMalloc or cudaSetDevice depending on whether the programmer accepts the default device. \u00c2\u00a0However, when either function is called, a long wait follows because CUDA is initializing. \u00c2\u00a0But it is not hopeless if the developer implements the solution carefully.<\/p>\n

As it turns out, a block of code in a program written in one language can usually be translated into equivalent code in another GPGPU programming language. \u00c2\u00a0For example, code to find and select a device in CUDA has an equivalent block of code in OpenCL. \u00c2\u00a0These blocks of code are called phases. \u00c2\u00a0What are the phases of a program? \u00c2\u00a0This depends, but one possible definition follows:<\/p>\n

    \n
  1. Setup resources on the GPU for problem solving, e.g., pick the GPU, allocate memory on the GPU, and compile kernel code for the GPU.<\/li>\n
  2. Copy data structures from the CPU to the GPU.<\/li>\n
  3. Call kernel: Run the kernel code to solve the problem on the GPU.<\/li>\n
  4. Copy data structures\u00c2\u00a0from the GPU to the CPU.<\/li>\n
  5. Release resources on the GPU.<\/li>\n<\/ol>\n
    The functions of each language can be grouped into these phases:<\/div>\n\n\n\n\n\n\n\n\n
    <\/th>\nCUDA<\/th>\nOpenCL<\/th>\nC++ AMP<\/th>\n<\/tr>\n
    Setup<\/td>\ncudaGetDeviceCount
    \ncudaGetDeviceProperties
    \ncudaSetDevice
    \ncudaMalloc<\/td>\n    		clGetPlatformIDs
    \nclGetPlatformInfo
    \nclGetDeviceIDs
    \nclGetDeviceInfo
    \nclCreateContext
    \nclCreateProgramWithSource
    \nclBuildProgram
    \nclGetProgramBuildInfo
    \nclCreateKernel
    \nclCreateBuffer
    \nclSetKernelArg
    \nclCreateCommandQueue<\/td>\n    		allocate vector<
    \naccelerator>get_accelerators()
    \nallocate\u00c2\u00a0vector::iterator
    \naccelerators.begin()
    \naccelerators.end()
    \nget_description()
    \nallocate array<>
    \nallocate\u00c2\u00a0extent<>
    \nallocate\u00c2\u00a0grid<><\/td>\n<\/tr>\n
    Copy to GPU<\/td>\ncudaMemcpy<\/td>\nclEnqueueWriteBuffer<\/td>\ncopy<\/td>\n<\/tr>\n
    Call kernel<\/td>\nkernel<<<…>>>(…)
    \ncudaDeviceSynchronize<\/td>\n    		clEnqueueNDRangeKernel
    \nclWaitForEvents<\/td>\n    		parallel_for_each
    \nflush
    \nwait<\/td>\n<\/tr>\n
    Copy to CPU<\/td>\ncudaMemcpy<\/td>\nclEnqueueReadBuffer<\/td>\ncopy<\/td>\n<\/tr>\n
    Release<\/td>\ncudaFree<\/td>\nclReleaseMemObject
    \nclReleaseKernel
    \nclReleaseContext
    \nclReleaseProgram<\/td>\n    		allocate grid<>
    \nallocate\u00c2\u00a0extent<>
    \nallocate\u00c2\u00a0array<>
    \ndeallocate\u00c2\u00a0accelerator_view
    \ndeallocate\u00c2\u00a0vector<accelerator>
    \ndeallocate\u00c2\u00a0vector::iterator<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

    Gathering run-time data<\/h1>\n

    How do we collect the run-times for each phase? \u00c2\u00a0There are a few ways to do this:<\/p>\n