GPGPU OpenCL 程式設計步驟與向量運算實作範例

OpenCL (Open Computing Language) 是一個為異構平台編寫程式的開放標準框架，該平台可由 CPU、GPU、DSP、FPGA 等不同類型的處理器組成。OpenCL 包含一門基於 C99 的核心語言（用於編寫在 OpenCL 裝置上執行的 kernels）以及一組用於定義與控制平台的 API。它提供任務並行與資料並行兩種並行機制，讓開發者能充分發揮異構運算硬體的效能。

OpenCL 與 CUDA 的比較

不同點： OpenCL 是通用的異構平台程式設計語言，為了相容不同裝置，API 設計較為繁瑣；CUDA 則是 NVIDIA 專為自家 GPGPU 設計的程式設計框架，API 相對簡單、入門門檻較低。

相同點： 兩者都基於任務並行與資料並行模型，核心程式設計概念一致。

OpenCL 程式設計 13 個標準步驟

1. 發現並初始化平台： 呼叫兩次 clGetPlatformIDs，第一次取得可用平台數量，第二次取得平台 ID。
2. 發現並初始化裝置： 呼叫兩次 clGetDeviceIDs，分別取得裝置數量與裝置 ID。
3. 建立內容： 呼叫 clCreateContext，內容可管理多個裝置。
4. 建立命令佇列： 呼叫 clCreateCommandQueue，每個裝置對應一個命令佇列，內容將命令發送至裝置對應的佇列執行。
5. 建立裝置緩衝區： 呼叫 clCreateBuffer，緩衝區儲存資料物件，由內容建立以便多裝置共享資料。
6. 寫入主機資料到裝置緩衝區： 呼叫 clEnqueueWriteBuffer。
7. 建立並編譯程式： 建立程式物件（原始碼或二進制），呼叫 clCreateProgramWithSource 與 clBuildProgram。
8. 建立核心： 呼叫 clCreateKernel，核心代表裝置程式的進入點。
9. 設定核心引數： 呼叫 clSetKernelArg 綁定緩衡區與參數。
10. 配置工作項目結構： 設定 work-item 組織形式（維度、工作群組大小等）。
11. 將核心加入佇列執行： 呼叫 clEnqueueNDRangeKernel。
12. 讀回輸出緩衝區到主機： 呼叫 clEnqueueReadBuffer。
13. 釋放 OpenCL 資源： 依序釋放核心、程式、記憶體物件、命令佇列、內容等。

程式碼架構說明

前三個步驟（平台、裝置、內容初始化）通常固定，可獨立封裝為工具模組。其餘步驟依應用需求調整，但基本流程不變。OpenCL 核心函數必須獨立放在 .cl 檔案中。

完整實作範例：向量運算

工具模組

tool.h - 宣告平台/裝置初始化與核心原始碼載入函式：

#ifndef TOOLH
#define TOOLH
#include <CL/cl.h>
#include <string>
#include <fstream>

int convertToString(const char *filename, std::string& s);
int getPlatform(cl_platform_id &platform);
cl_device_id *getCl_device_id(cl_platform_id &platform);
#endif

tool.cpp - 實作核心載入、平台選擇（取第一個可用）、GPU 裝置查詢：

int convertToString(const char *filename, std::string& s) {
    std::fstream f(filename, std::fstream::in | std::fstream::binary);
    if (f.is_open()) {
        size_t fileSize = (size_t)f.tellg();
        f.seekg(0, std::fstream::beg);
        char* str = new char[fileSize + 1];
        f.read(str, fileSize);
        str[fileSize] = '\0';
        s = str;
        delete[] str;
        return 0;
    }
    return -1;
}

int getPlatform(cl_platform_id &platform) {
    cl_uint numPlatforms;
    cl_int status = clGetPlatformIDs(0, NULL, &numPlatforms);
    if (status != CL_SUCCESS || numPlatforms == 0) return -1;
    cl_platform_id* platforms = (cl_platform_id*)malloc(numPlatforms * sizeof(cl_platform_id));
    status = clGetPlatformIDs(numPlatforms, platforms, NULL);
    platform = platforms[0];
    free(platforms);
    return 0;
}

cl_device_id *getCl_device_id(cl_platform_id &platform) {
    cl_uint numDevices = 0;
    cl_int status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &numDevices);
    if (numDevices > 0) {
        cl_device_id *devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));
        status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
        return devices;
    }
    return NULL;
}

核心程式

HelloWorld_Kernel.cl - 簡單向量運算核心：

__kernel void helloworld(__global double* in, __global double* out) {
    int num = get_global_id(0);
    out[num] = in[num] / 2.4 * (in[num] / 6);
}

主程式

HelloWorld.cpp - 完整執行流程：

int main() {
    cl_int status;
    cl_platform_id platform;
    getPlatform(platform);
    cl_device_id *devices = getCl_device_id(platform);
    
    cl_context context = clCreateContext(NULL, 1, devices, NULL, NULL, NULL);
    cl_command_queue commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
    
    // 建立並編譯程式
    std::string sourceStr;
    convertToString("HelloWorld_Kernel.cl", sourceStr);
    const char *source = sourceStr.c_str();
    size_t sourceSize[] = {strlen(source)};
    cl_program program = clCreateProgramWithSource(context, 1, &source, sourceSize, NULL);
    status = clBuildProgram(program, 1, devices, NULL, NULL, NULL);
    
    // 準備資料
    const int NUM = 512000;
    double* input = new double[NUM];
    for (int i = 0; i < NUM; i++) input[i] = i;
    double* output = new double[NUM];
    
    cl_mem inputBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
        NUM * sizeof(double), input, NULL);
    cl_mem outputBuffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
        NUM * sizeof(double), NULL, NULL);
    
    // 建立核心、設定引數、執行
    cl_kernel kernel = clCreateKernel(program, "helloworld", NULL);
    clSetKernelArg(kernel, 0, sizeof(cl_mem), &inputBuffer);
    clSetKernelArg(kernel, 1, sizeof(cl_mem), &outputBuffer);
    
    size_t global_work_size[1] = {NUM};
    cl_event event;
    clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, global_work_size, NULL, 0, NULL, &event);
    clWaitForEvents(1, &event);
    clReleaseEvent(event);
    
    // 讀回結果
    clEnqueueReadBuffer(commandQueue, outputBuffer, CL_TRUE, 0,
        NUM * sizeof(double), output, 0, NULL, NULL);
    std::cout << output[NUM-1] << std::endl;
    
    // 清理資源
    clReleaseKernel(kernel);
    clReleaseProgram(program);
    clReleaseMemObject(inputBuffer);
    clReleaseMemObject(outputBuffer);
    clReleaseCommandQueue(commandQueue);
    clReleaseContext(context);
    delete[] input;
    delete[] output;
    free(devices);
    return 0;
}

編譯與執行

g++ -I /opt/AMDAPP/include/ -o A *.cpp -lOpenCL
./A

這個範例展示了 OpenCL 從平台初始化到核心執行的完整流程，雖然步驟繁多，但掌握此標準流程後，即可將其應用於更複雜的並行運算任務。