從 CUDA 到 OpenCL


這篇文章是參考 nVidia CUDA 官方網站提供的《OpenCL JumpStart Guide》這份文件,來源是「OpenCL for NVIDIA」這個頁面;基本上這份文件主要的目的,就是透過由 CUDA Driver API 的寫法(Low Level),來和 OpenCL 做一個對比;如此一來,對於已經會 CUDA 的人來說,應該可以更簡單的學會如何使用 OpenCL 了~(但是現在會寫也沒得用啊…)

由此也看得出來,其實 OpenCL 算是和 CUDA 的 Driver API 比較接近、而非 Runtime API。但是,一般 Heresy 在玩 CUDA 的時候,都是用操作比較簡單的 high level runtime API,low level 的 driver API 到還沒真的用過了…

總之,接下來還是來看 CUDA 和 OpenCL 的差異吧~在文件中,是以 Vector Addition 得成是來當作基本的範例。

在 kernel Code 的部分,兩者的程式分別為:

C for CUDA
OpenCL
__global__ void
vectorAdd(const float * a,
          const float * b,
          float * c)
{
  // Vector element index</font>
  int nIndex = blockIdx.x * blockDim.x
             + threadIdx.x;
  c[nIndex] = a[nIndex] + b[nIndex];
}
__kernel void
vectorAdd(__global const float * a,
          __global const float * b,
          __global float * c)
{
  // Vector element index
  int nIndex = get_global_id(0);

  c[nIndex] = a[nIndex] + b[nIndex];
}

可以發現,在這個例子裡,兩邊的差異基本上只有 syntax 的部分~包括了:

  1. Kernel 的宣告:在 CUDA 是要用「__global__」來宣告 kernel function,而在 OpenCL 中,則是要用「__kernel」。

  2. 指標的宣告:OpenCL 會強制要求你去指定傳到 kernel function 的 pointer 是在哪一個記憶體空間;在這個例子則是要加上「__global」。

  3. Index 的計算:在 CUDA 裡,要計算出某個 thread 在整體的座標,是要透過 threadIdxblockIdxblockDimgridDim 這四個變數來做計算而得出。

    而在 OpenCL 裡,則是可以接透過 get_global_id() 這個函式來取得計算後的結果!當然,OpenCL 也有提供 get_local_id()get_work_dim()get_global_size() 這些函式,來取得細部的資料。

而在 Host code 的部分,兩者的程式分別是:

CUDA
OpenCL
const unsigned int cnBlockSize = 512;
const unsigned int cnBlocks = 3;
const unsigned int cnDimension
                   = cnBlocks * cnBlockSize;
const unsigned int cnBlockSize = 512;
const unsigned int cnBlocks = 3;
const unsigned int cnDimension
                   = cnBlocks * cnBlockSize;
CUdevice hDevice;
CUcontext hContext;

// create CUDA device & context
cuInit(0);
// pick first device
cuDeviceGet(&hContext, 0);
cuCtxCreate(&hContext, 0, hDevice));
// create OpenCL device & context
cl_context hContext;
hContext = clCreateContextFromType(  
       0, CL_DEVICE_TYPE_GPU, 0, 0, 0);
 
// query all devices available to the context
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);
 
// create a command queue
// for first device the context reported
cl_command_queue hCmdQueue;
hCmdQueue = clCreateCommandQueue(hContext,
                                 aDevices[0],
                                 0, 0);
CUmodule hModule;
CUfunction hFunction;

cuModuleLoad(&hModule, "vectorAdd.cubin");
cuModuleGetFunction(&hFunction, hModule,
                    "vectorAdd");
// create & compile program
cl_program hProgram;
hProgram = clCreateProgramWithSource(
             hContext, 1, sProgramSource, 0, 0);
clBuildProgram(hProgram, 0, 0, 0, 0, 0);
 
// create kernel
cl_kernel hKernel;
hKernel = clCreateKernel(hProgram, "vectorAdd", 0);
// allocate host vectors
float * pA = new float[cnDimension];
float * pB = new float[cnDimension];
float * pC = new float[cnDimension];
 
// initialize host memory
randomInit(pA, cnDimension);
randomInit(pB, cnDimension);
// allocate host vectors
float * pA = new float[cnDimension];
float * pB = new float[cnDimension];
float * pC = new float[cnDimension];
 
// initialize host memory
randomInit(pA, cnDimension);
randomInit(pB, cnDimension);
// allocate memory on the device
CUdeviceptr pDeviceMemA,
            pDeviceMemB,
            pDeviceMemC;
cuMemAlloc(&pDeviceMemA,
           cnDimension * sizeof(float));
cuMemAlloc(&pDeviceMemB,
            cnDimension * sizeof(float));
cuMemAlloc(&pDeviceMemC,
           cnDimension * sizeof(float));
 
// copy host vectors to device
cuMemcpyHtoD(pDeviceMemA,
             pA,
             cnDimension * sizeof(float));
cuMemcpyHtoD(pDeviceMemB,
             pB,
             cnDimension * sizeof(float));
// allocate device memory
cl_mem hDeviceMemA, hDeviceMemB, hDeviceMemC;
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);
// setup parameter values
cuFuncSetBlockShape(cuFunction,
                    cnBlockSize, 1, 1);
cuParamSeti(cuFunction, 0, pDeviceMemA);
cuParamSeti(cuFunction, 4, pDeviceMemB);
cuParamSeti(cuFunction, 8, pDeviceMemC);
cuParamSetSize(cuFunction, 12);
 
// execute kernel
cuLaunchGrid(cuFunction, cnBlocks, 1);
// setup parameter values
clSetKernelArg(hKernel, 0, sizeof(cl_mem),
               (void *)&hDeviceMemA);
clSetKernelArg(hKernel, 1, sizeof(cl_mem),
               (void *)&hDeviceMemB);
clSetKernelArg(hKernel, 2, sizeof(cl_mem),
                (void *)&hDeviceMemC);
 
// execute kernel
clEnqueueNDRangeKernel(hCmdQueue, hKernel,
                    1, 0, &cnDimension, 0, 0, 0, 0);
// copy the result from device back to host
cuMemcpyDtoH( (void *) pC,
              pDeviceMemC,
              cnDimension * sizeof(float));
// copy results from device back to host
clEnqueueReadBuffer(hContext,
                    hDeviceMemC,
                    CL_TRUE,
                    0,
                    cnDimension * sizeof(cl_float),
                    pC, 0, 0, 0);
delete[] pA;
delete[] pB;
delete[] pC;
 
cuMemFree(pDeviceMemA);
cuMemFree(pDeviceMemB);
cuMemFree(pDeviceMemC);
delete[] pA;
delete[] pB;
delete[] pC;
 
clReleaseMemObj(hDeviceMemA);
clReleaseMemObj(hDeviceMemB);
clReleaseMemObj(hDeviceMemC);

從上面這個例子可以發現,其實 CUDA 和 OpenCL 在概念層面算是差不多的!不過如果是用 CUDA 的 Runtime API 的話,倒是可以省掉不少麻煩(可以參考《簡單的 CUDA 程式:VectorAdd》)。

而在不少情況下,CUDA 和 OpenCL 的差異都只是所呼叫的函式不同罷了!

現階段,大概就先整理整裡而已吧…畢竟,在沒有編譯器可用的情況下,實在沒什麼興趣認真玩啊…等之後有可以玩的平台後,在來認真研究它的語法吧~(謎之聲:到時候有空嗎?)


對「從 CUDA 到 OpenCL」的想法

  1. ㄟ…這種和文章內容無關的,麻煩到訪客留言的地方貼吧。 @@http://heresy.spaces.live.com/guestbook/此外,不知道你的預算、用途?

  2. 大虾,我最近要装机,不知道你能否给我一些意见呢?比如a平台还是i平台?a卡还是n卡?用什么主机板?

發表迴響

在下方填入你的資料或按右方圖示以社群網站登入:

WordPress.com 標誌

您的留言將使用 WordPress.com 帳號。 登出 /  變更 )

Google photo

您的留言將使用 Google 帳號。 登出 /  變更 )

Twitter picture

您的留言將使用 Twitter 帳號。 登出 /  變更 )

Facebook照片

您的留言將使用 Facebook 帳號。 登出 /  變更 )

連結到 %s

This site uses Akismet to reduce spam. Learn how your comment data is processed.