OpenCL 和 CUDA 之間的區別

OpenCL馬上就要發布了, 根據nvidia的官方文檔,對OpenCL和CUDA的異同做比較:

 

•         指標遍曆

OpenCL不支援CUDA那樣的指標遍曆方式, 你只能用下標方式間接實現指標遍曆. 例子代碼如下:
// CUDA

struct Node { Node* next; }
n = n->next;

 // OpenCL

struct Node { unsigned int next; }

n = bufBase + n;

• Kernel 程式異同

CUDA的代碼最終編譯成顯卡上的二進位格式，最後由cudart.dll(個人猜測)裝載到GPU並且執行。OpenCL中執行階段程式庫中包含編譯器，

使用虛擬碼，程式運行時即時編譯和裝載。這個類似JAVA, .net 程式，道理也一樣，為了支援跨平台的相容。kernel程式的文法也

有略微不同，如下：
__global__ void vectorAdd(const float * a, const float * b, float * c)<br />{ // CUDA<br /> int nIndex = blockIdx.x * blockDim.x + threadIdx.x;<br /> c[nIndex] = a[nIndex] + b[nIndex];<br />}
__kernel void vectorAdd(__global const float * a, __global const float * b, __global float * c)<br />{ // OpenCL<br /> int nIndex = get_global_id(0);<br /> c[nIndex] = a[nIndex] + b[nIndex];<br />} 

可以看出大部分都相同。只是細節有差異：

1）CUDA 的kernel函數使用“__global__”申明而OpenCL的kernel函數使用“__kernel”作為申明。

2）OpenCL的所有參數都有“__global”修飾符，代表這個參數所指地址是在全域記憶體。

3）眾所周知，CUDA採用threadIdx.{x|y|z}, blockIdx.{x|y|z}來獲得當前線程的索引號，而OpenCL

     通過一個特定的get_global_id()函數來獲得在kernel中的全域索引號。OpenCL中如果要獲得在當前工作

     組（對等於CUDA中的block）中的局部索引號，可以使用get_local_id()

• Host代碼的異同

把上面的kernel代碼編譯成“vectorAdd.cubin”，CUDA調用方法如下：

const unsigned int cnBlockSize = 512;<br />const unsigned int cnBlocks = 3;<br />const unsigned int cnDimension = cnBlocks * cnBlockSize;<br />CUdevice hDevice;<br />CUcontext hContext;<br />CUmodule hModule;<br />CUfunction hFunction;<br />// create CUDA device & context<br />cuInit(0);<br />cuDeviceGet(&hContext, 0); // pick first device<br />cuCtxCreate(&hContext, 0, hDevice));<br />cuModuleLoad(&hModule, “vectorAdd.cubin”);<br />cuModuleGetFunction(&hFunction, hModule, "vectorAdd");<br />// allocate host vectors<br />float * pA = new float[cnDimension];<br />float * pB = new float[cnDimension];<br />float * pC = new float[cnDimension];<br />// initialize host memory<br />randomInit(pA, cnDimension);<br />randomInit(pB, cnDimension);<br />// allocate memory on the device<br />CUdeviceptr pDeviceMemA, pDeviceMemB, pDeviceMemC;<br />cuMemAlloc(&pDeviceMemA, cnDimension * sizeof(float));<br />cuMemAlloc(&pDeviceMemB, cnDimension * sizeof(float));<br />cuMemAlloc(&pDeviceMemC, cnDimension * sizeof(float));<br />// copy host vectors to device<br />cuMemcpyHtoD(pDeviceMemA, pA, cnDimension * sizeof(float));<br />cuMemcpyHtoD(pDeviceMemB, pB, cnDimension * sizeof(float));<br />// setup parameter values<br />cuFuncSetBlockShape(cuFunction, cnBlockSize, 1, 1);<br />cuParamSeti(cuFunction, 0, pDeviceMemA);<br />cuParamSeti(cuFunction, 4, pDeviceMemB);<br />cuParamSeti(cuFunction, 8, pDeviceMemC);<br />cuParamSetSize(cuFunction, 12);<br />// execute kernel<br />cuLaunchGrid(cuFunction, cnBlocks, 1);<br />// copy the result from device back to host<br />cuMemcpyDtoH((void *) pC, pDeviceMemC, cnDimension * sizeof(float));<br />delete[] pA;<br />delete[] pB;<br />delete[] pC;<br />cuMemFree(pDeviceMemA);<br />cuMemFree(pDeviceMemB);<br />cuMemFree(pDeviceMemC);

OpenCL的代碼以文本方式存放在“sProgramSource”。 調用方式如下：

const unsigned int cnBlockSize = 512;<br />const unsigned int cnBlocks = 3;<br />const unsigned int cnDimension = cnBlocks * cnBlockSize;<br />// create OpenCL device & context<br />cl_context hContext;<br />hContext = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, 0, 0, 0);<br />// query all devices available to the context<br />size_t nContextDescriptorSize;<br />clGetContextInfo(hContext, CL_CONTEXT_DEVICES, 0, 0, &nContextDescriptorSize);<br />cl_device_id * aDevices = malloc(nContextDescriptorSize);<br />clGetContextInfo(hContext, CL_CONTEXT_DEVICES, nContextDescriptorSize, aDevices, 0);<br />// create a command queue for first device the context reported<br />cl_command_queue hCmdQueue;<br />hCmdQueue = clCreateCommandQueue(hContext, aDevices[0], 0, 0);<br />// create & compile program<br />cl_program hProgram;<br />hProgram = clCreateProgramWithSource(hContext, 1, sProgramSource, 0, 0);<br />clBuildProgram(hProgram, 0, 0, 0, 0, 0);// create kernel<br />cl_kernel hKernel;<br />hKernel = clCreateKernel(hProgram, “vectorAdd”, 0);<br />// allocate host vectors<br />float * pA = new float[cnDimension];<br />float * pB = new float[cnDimension];<br />float * pC = new float[cnDimension];<br />// initialize host memory<br />randomInit(pA, cnDimension);<br />randomInit(pB, cnDimension);<br />// allocate device memory<br />cl_mem hDeviceMemA, hDeviceMemB, hDeviceMemC;<br />hDeviceMemA = clCreateBuffer(hContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cnDimension * sizeof(cl_float), pA, 0);<br />hDeviceMemB = clCreateBuffer(hContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cnDimension * sizeof(cl_float), pA, 0);<br />hDeviceMemC = clCreateBuffer(hContext,<br />CL_MEM_WRITE_ONLY,<br />cnDimension * sizeof(cl_float), 0, 0);<br />// setup parameter values<br />clSetKernelArg(hKernel, 0, sizeof(cl_mem), (void *)&hDeviceMemA);<br />clSetKernelArg(hKernel, 1, sizeof(cl_mem), (void *)&hDeviceMemB);<br />clSetKernelArg(hKernel, 2, sizeof(cl_mem), (void *)&hDeviceMemC);<br />// execute kernel<br />clEnqueueNDRangeKernel(hCmdQueue, hKernel, 1, 0, &cnDimension, 0, 0, 0, 0);<br />// copy results from device back to host<br />clEnqueueReadBuffer(hContext, hDeviceMemC, CL_TRUE, 0, cnDimension * sizeof(cl_float),<br />pC, 0, 0, 0);<br />delete[] pA;<br />delete[] pB;<br />delete[] pC;<br />clReleaseMemObj(hDeviceMemA);<br />clReleaseMemObj(hDeviceMemB);<br />clReleaseMemObj(hDeviceMemC);

• 初始化部分的異同 

CUDA 在使用任何API之前必須調用cuInit(0)，然後是獲得當前系統的可用裝置並獲得Context。
cuInit(0);
cuDeviceGet(&hContext, 0);
cuCtxCreate(&hContext, 0, hDevice));
OpenCL不用全域的初始化，直接指定裝置獲得控制代碼就可以了
cl_context hContext;
hContext = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, 0, 0, 0);
裝置建立完畢後，可以通過下面的方法獲得裝置資訊和上下文：
size_t nContextDescriptorSize;
clGetContextInfo(hContext, CL_CONTEXT_DEVICES, 0, 0, &nContextDescriptorSize);
cl_device_id * aDevices = malloc(nContextDescriptorSize);

clGetContextInfo(hContext, CL_CONTEXT_DEVICES, nContextDescriptorSize, aDevices, 0);

OpenCL introduces an additional concept: Command Queues. Commands launching kernels and
reading or writing memory are always issued for a specific command queue. A command queue is
created on a specific device in a context. The following code creates a command queue for the
device and context created so far:
cl_command_queue hCmdQueue;
hCmdQueue = clCreateCommandQueue(hContext, aDevices[0], 0, 0);
With this the program has progressed to the point where data can be uploaded to the device’s
memory and processed by launching compute kernels on the device.

• Kernel Creation

CUDA kernel 以二進位格式存放與CUBIN檔案中間，其調用格式和DLL的用法比較類似，先裝載二進位庫，然後通過函數名尋找

函數地址，最後用將函數裝載到GPU運行。範例程式碼如下：
CUmodule hModule;
cuModuleLoad(&hModule, “vectorAdd.cubin”);
cuModuleGetFunction(&hFunction, hModule, "vectorAdd");
OpenCL 為了支援多平台，所以不使用編譯後的代碼，採用類似JAVA的方式，裝載文字格式設定的代碼檔案，然後即時編譯並運行。
需要注意的是，OpenCL也提供API訪問kernel的二進位程式，前提是這個kernel已經被編譯並且放在某個特定的緩衝中了。

// 裝載代碼，即時編譯
cl_program hProgram;
hProgram = clCreateProgramWithSource(hContext, 1, “vectorAdd.c", 0, 0);
clBuildProgram(hProgram, 0, 0, 0, 0, 0);
// 獲得kernel函數控制代碼
cl_kernel hKernel;
hKernel = clCreateKernel(hProgram, “vectorAdd”, 0);

 

• 裝置記憶體配置

記憶體配置沒有什麼大區別，OpenCL提供兩組特殊的標誌，CL_MEM_READ_ONLY  和 CL_MEM_WRITE_ONLY 用來控制記憶體

的讀寫權限。另外一個標誌比較有用：CL_MEM_COPY_HOST_PTR 表示這個記憶體在主機分配，但是GPU可以使用，運行時會自動

將主機記憶體內容拷貝到GPU，主機記憶體配置，裝置記憶體配置，主機拷貝資料到裝置，3個步驟一氣呵成。
// CUDA

CUdeviceptr pDeviceMemA, pDeviceMemB, pDeviceMemC;
cuMemAlloc(&pDeviceMemA, cnDimension * sizeof(float));
cuMemAlloc(&pDeviceMemB, cnDimension * sizeof(float));
cuMemAlloc(&pDeviceMemC, cnDimension * sizeof(float));
cuMemcpyHtoD(pDeviceMemA, pA, cnDimension * sizeof(float));
cuMemcpyHtoD(pDeviceMemB, pB, cnDimension * sizeof(float));
// OpenCL
hDeviceMemA = clCreateBuffer(hContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cnDimension * sizeof(cl_float), pA, 0);
hDeviceMemB = clCreateBuffer(hContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cnDimension * sizeof(cl_float), pA, 0);
hDeviceMemC = clCreateBuffer(hContext, CL_MEM_WRITE_ONLY, cnDimension * sizeof(cl_float), 0, 0);

• Kernel Parameter Specification

The next step in preparing the kernels for launch is to establish a mapping between the kernels’
parameters, essentially pointers to the three vectors A, B and C, to the three device memory regions,
which were allocated in the previous section.
Parameter setting in both APIs is a pretty low-level affair. It requires knowledge of the total number
, order, and types of a given kernel’s parameters. The order and types of the parameters are used to
determine a specific parameters offset inside the data block made up of all parameters. The offset in
bytes for the n-th parameter is essentially the sum of the sizes of all (n-1) preceding parameters.
Using the CUDA Driver API:
In CUDA device pointers are represented as unsigned int and the CUDA Driver API has a
dedicated method for setting that type. Here’s the code for setting the three parameters. Note how
the offset is incrementally computed as the sum of the previous parameters’ sizes.
cuParamSeti(cuFunction, 0, pDeviceMemA);
cuParamSeti(cuFunction, 4, pDeviceMemB);
cuParamSeti(cuFunction, 8, pDeviceMemC);
cuParamSetSize(cuFunction, 12);
Using OpenCL:
In OpenCL parameter setting is done via a single function that takes a pointer to the location of the
parameter to be set.
clSetKernelArg(hKernel, 0, sizeof(cl_mem), (void *)&hDeviceMemA);
clSetKernelArg(hKernel, 1, sizeof(cl_mem), (void *)&hDeviceMemB);
clSetKernelArg(hKernel, 2, sizeof(cl_mem), (void *)&hDeviceMemC);

• Kernel Launch

Launching a kernel requires the specification of the dimension and size of the “thread-grid”. The
CUDA Programming Guide and the OpenCL specification contain details about the structure of
those grids. For NVIDIA GPUs the permissible structures are the same for CUDA and OpenCL.
For the vectorAdd sample we need to start one thread per vector-element (of the output vector).
The number of elements in the vector is given in the cnDimension variable. It is defined to be
cnDimension = cnBlockSize * cnBlocks. This means that cnDimension threads
need to be executed. The threads are structured into cnBlocks one-dimensional thread blocks of
size cnBlockSize.
Using the CUDA Driver API:
A kernel’s block size is specified in a call separate from the actual kernel launch using
cuFunctSetBlockShape. The kernel launching function cuLaunchGrid then only
specifies the number of blocks to be launched.
cuFuncSetBlockShape(cuFunction, cnBlockSize, 1, 1);
cuLaunchGrid (cuFunction, cnBlocks, 1);
Using OpenCL:
The OpenCL equivalent of kernel launching is to “enqueue” a kernel for execution into a command
queue. The enqueue function takes parameters for both the work group size (work group is the
OpenCL equivalent of a CUDA thread-block), and the global work size, which is the size of the
global array of threads.
Note: Where in CUDA the global work size is specified in terms of number of thread
blocks, it is given in number of threads in OpenCL.
Both work group size and global work size are potentially one, two, or three dimensional arrays. The
function expects pointers of unsigned ints to be passed in the fourth and fifth parameters.
For the vectorAdd example, work groups and total work size is a one-dimensional grid of threads.
clEnqueueNDRangeKernel(hCmdQueue, hKernel, 1, 0,
&cnDimension, &cnBlockSize, 0, 0, 0);
The parameters of cnDimension and cnBlockSize must be pointers to unsigned int.
Work group sizes that are dimensions greater than 1, the parameters will be a pointer to arrays of
sizes.

• Result Data Retrieval

Both kernel launch functions (CUDA and OpenCL) are asynchronous, i.e. they return immediately
after scheduling the kernel to be executed on the GPU. In order for a copy operation that retrieves
the result vector C (copy from device to host) to produce correct results in synchronization with the
kernel completion needs to happen.
CUDA memcpy functions automatically synchronize and complete any outstanding kernel launches
proceeding. Both API’s also provide a set of asynchronous memory transfer functions which
allows a user to overlap memory transfers with computation to increase throughput.
Using the CUDA Driver API:
Use cuMemcpyDtoH() to copy results back to the host.
cuMemcpyDtoH((void *)pC, pDeviceMemC, cnDimension * sizeof(float));
Using OpenCL:
OpenCL’s clEnqueueReadBuffer() function allows the user to specify whether a read is to
be synchronous or asynchronous (third argument). For the simple vectorAdd sample a
synchronizing read is used, which results in the same behavior as the simple synchronous CUDA
memory copy above:
clEnqueueReadBuffer(hContext, hDeviceC, CL_TRUE, 0, cnDimension * sizeof(cl_float), pC, 0, 0, 0);
When used for asynchronous reads, OpenCL has an event mechanism that allows the host
application to query the status or wait for the completion of a given call.

 

The current version of OpenCL does not support stream offsets at the API/kernel invocation level.
Offsets must be passed in as a parameter to the kernel and the address of the memory computed
inside it. CUDA kernels may be started at offsets within buffers at the API/kernel invocation level.