OpenCL使用CL_MEM_USE_HOST_PTR存儲器對象屬性與存儲器映射

隨着OpenCL的普及，現在有越來越多的移動設備以及平板、超級本等都支持OpenCL異構計算。而這些設備與桌面計算機、服務器相比而言性能不是占主要因素的，反而能耗更受人關注。因此，這些移動設備上的GPU與CPU基本都是在同一芯片上（SoC），或者GPU就已經成為了處理器的一部分，像Intel Ivy Bridge架構開始的處理器（Intel HD Graphics 4000開始支持OpenCL），AMD APU等。

因此，在這些設備上做OpenCL的異構並行計算的話，我們不需要像桌面端那些獨立GPU那樣，要把主存數據通過PCIe搬運到GPU端，然后等GPU計算結束后再搬回到主存。我們只需要將給GPU端分配的顯存映射到主機端即可。這樣，在主機端我們也能直接通過指針來操作這塊存儲數據。

下面編寫了一個比較簡單的例子來描述如何使用OpenCL的存儲器映射特性。這個例子在MacBook Air，OS X 10.9.2下完成，並通過Xcode 5.1，Apple LLVM 5.1的編譯與運行。 硬件環境為：Intel Core i7 4650U, Intel Graphics 5000, 8GB DDR3L, 128GB SSD

這是主機端代碼（C源文件）：

#include <stdio.h>
#include <string.h>
#include <stdlib.h>
#include <time.h>

#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif


int main(void)
{
    cl_int ret;
    
    cl_platform_id platform_id = NULL;
    cl_device_id device_id = NULL;
    cl_context context = NULL;
    cl_command_queue command_queue = NULL;
    cl_mem memObj = NULL;
    char *kernelSource = NULL;
    cl_program program = NULL;
    cl_kernel kernel = NULL;
    int *pHostBuffer = NULL;

    
    clGetPlatformIDs(1, &platform_id, NULL);
    if(platform_id == NULL)
    {
        puts("Get OpenCL platform failed!");
        goto FINISH;
    }
    
    clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
    if(device_id == NULL)
    {
        puts("No GPU available as a compute device!");
        goto FINISH;
    }
    
    context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
    if(context == NULL)
    {
        puts("Context not established!");
        goto FINISH;
    }
    
    command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
    if(command_queue == NULL)
    {
        puts("Command queue cannot be created!");
        goto FINISH;
    }
    
    // 指定內核源文件路徑
    const char *pFileName = "/Users/zennychen/Downloads/test.cl";
    
    FILE *fp = fopen(pFileName, "r");
    if (fp == NULL)
    {
        puts("The specified kernel source file cannot be opened!");
        goto FINISH;
    }
    fseek(fp, 0, SEEK_END);
    const long kernelLength = ftell(fp);
    fseek(fp, 0, SEEK_SET);
    
    kernelSource = malloc(kernelLength);
    
    fread(kernelSource, 1, kernelLength, fp);
    fclose(fp);
    
    program = clCreateProgramWithSource(context, 1, (const char**)&kernelSource, (const size_t*)&kernelLength, &ret);
    ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
    if (ret != CL_SUCCESS)
    {
        size_t len;
        char buffer[8 * 1024];
        
        printf("Error: Failed to build program executable!\n");
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        printf("%s\n", buffer);
        goto FINISH;
    }
    
    kernel = clCreateKernel(program, "test", &ret);
    if(kernel == NULL)
    {
        puts("Kernel failed to create!");
        goto FINISH;
    }
    
    const size_t contentLength = sizeof(*pHostBuffer) * 1024 * 1024;
    
    // 以下為在主機端分配輸入緩存
    pHostBuffer = malloc(contentLength);
    
    // 然后對此工作緩存進行初始化
    for(int i = 0; i < 1024 * 1024; i++)
        pHostBuffer[i] = i + 1;
    
    // 這里預分配的緩存大小為4MB，第一個參數是讀寫的
    memObj = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, contentLength, pHostBuffer, &ret);
    if(memObj == NULL)
    {
        puts("Memory object1 failed to create!");
        goto FINISH;
    }
    
    ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&memObj);
    
    if(ret != CL_SUCCESS)
    {
        puts("Set arguments error!");
        goto FINISH;
    }
    
    // 做存儲器映射
    int *pDeviceBuffer = clEnqueueMapBuffer(command_queue, memObj, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE, 0, contentLength, 0, NULL, NULL, &ret);
    if(pDeviceBuffer == NULL)
    {
        puts("Memory map failed!");
        goto FINISH;
    }
    if(pDeviceBuffer != pHostBuffer)
    {
        // 若從GPU端映射得到的存儲器地址與原先主機端的不同，則將數據從主機端傳遞到GPU端
        ret = clEnqueueWriteBuffer(command_queue, memObj, CL_TRUE, 0, contentLength, pHostBuffer, 0, NULL, NULL);
        if(ret != CL_SUCCESS)
        {
            puts("Data transfer failed");
            goto FINISH;
        }
        
        /** 如果主機端與設備端地址不同，我們不妨測試一下設備端存儲器的Cache情況 */
        
        // 先測試主機端的時間
        int sum = 0;
        
        // 先過一遍存儲器
        for(int j = 0; j < 1024; j++)
            sum += pHostBuffer[j];
        
        time_t t1 = time(NULL);
        for(int i = 0; i < 1000000; i++)
        {
            for(int j = 0; j < 1024; j++)
                sum += pHostBuffer[j];
        }
        time_t t2 = time(NULL);
        printf("The host delta time is: %f. The value is: %d\n", difftime(t2, t1), sum);
        
        // 測試設備端
        sum = 0;
        
        // 先過一遍存儲器
        for(int j = 0; j < 1024; j++)
            sum += pDeviceBuffer[j];
        
        t1 = time(NULL);
        for(int i = 0; i < 1000000; i++)
        {
            for(int j = 0; j < 1024; j++)
                sum += pDeviceBuffer[j];
        }
        t2 = time(NULL);
        printf("The device delta time is: %f. The value is: %d\n", difftime(t2, t1), sum);
    }
    else
    {
        // 若主機端與設備端存儲器地址相同，我們僅僅做CPU端測試
        int sum = 0;
        
        // 先過一遍存儲器
        for(int j = 0; j < 1024; j++)
            sum += pHostBuffer[j];
        
        time_t t1 = time(NULL);
        for(int i = 0; i < 1000000; i++)
        {
            for(int j = 0; j < 1024; j++)
                sum += pHostBuffer[j];
        }
        time_t t2 = time(NULL);
        printf("The host delta time is: %f. The value is: %d\n", difftime(t2, t1), sum);
    }

    // 這里指定將總共有1024 * 1024個work-item
    ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, (const size_t[]){1024 * 1024}, NULL, 0, NULL, NULL);
    
    // 做次同步，這里偷懶，不用wait event機制了～
    clFinish(command_queue);
        
    // 做校驗
    for(int i = 0; i < 1024 * 1024; i++)
    {
        if(pDeviceBuffer[i] != (i + 1) * 2)
        {
            puts("Result error!");
            break;
        }
    }
    
    puts("Compute finished!");
    
FINISH:
    
    /* Finalization */
    if(pHostBuffer != NULL)
        free(pHostBuffer);
    
    if(kernelSource != NULL)
        free(kernelSource);
    
    if(memObj != NULL)
        clReleaseMemObject(memObj);
    
    if(kernel != NULL)
        clReleaseKernel(kernel);
    
    if(program != NULL)
        clReleaseProgram(program);
    
    if(command_queue != NULL)
        clReleaseCommandQueue(command_queue);
    
    if(context != NULL)
        clReleaseContext(context);
    
    return 0;
}

以下是OpenCL內核源代碼：

__kernel void test(__global int *pInOut)
{
    int index = get_global_id(0);
    
    pInOut[index] += pInOut[index];
}

另外，主機端代碼部分中，OpenCL源文件路徑是寫死的。各位朋友可以根據自己環境來重新指定路徑。

當然，我們還可以修改主機端“clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, contentLength, pHostBuffer, &ret);”這段創建存儲器對象的屬性。比如，將CL_MEM_USE_HOST_PTR去掉。然后可以再試試效果。

倘若clCreateBuffer的flags參數用的是CL_MEM_ALLOC_HOST_PTR，那么其host_ptr參數必須為空。在調用clEnqueueMapBuffer之后，可以根據其返回的緩存地址，對存儲區域做數據初始化。

CL_MEM_ALLOC_HOST_PTR表示應用程序暗示OpenCL實現從主機端可訪問的存儲空間給設備端分配存儲緩存。這個與CL_MEM_USE_HOST_PTR還是有所區別的。CL_MEM_USE_HOST_PTR是完全從應用端當前的內存池分配存儲空間；而CL_MEM_ALLOC_HOST_PTR對於CPU與GPU共享主存的環境下，可以在CPU端留下一個訪問GPU端VRAM的入口點。我們通過以下程序來測試當前環境的OpenCL實現（以下代碼在調用調用了clEnqueueMapBuffer函數之后做了緩存數據初始化的時間比較）：

    long deltaTimes[10];
    
    for(int i = 0; i < 10; i++)
    {
        struct timeval tBegin, tEnd;
        gettimeofday(&tBegin, NULL);
        
        for(int i = 0; i < 1024 * 1024; i++)
            pDeviceBuffer[i] = i + 1;
        
        gettimeofday(&tEnd, NULL);
        
        deltaTimes[i] = 1000000 * (tEnd.tv_sec - tBegin.tv_sec ) + tEnd.tv_usec - tBegin.tv_usec;
    }
    
    long useTime = deltaTimes[0];
    for(int i = 1; i < 10; i++)
    {
        if(useTime > deltaTimes[i])
            useTime = deltaTimes[i];
    }
    
    printf("Device memory time spent: %ldus\n", useTime);
    
    int *pHostBuffer = malloc(contentLength);
    for(int i = 0; i < 10; i++)
    {
        struct timeval tBegin, tEnd;
        gettimeofday(&tBegin, NULL);
        
        for(int i = 0; i < 1024 * 1024; i++)
            pHostBuffer[i] = i + 1;
        
        gettimeofday(&tEnd, NULL);
        
        deltaTimes[i] = 1000000 * (tEnd.tv_sec - tBegin.tv_sec ) + tEnd.tv_usec - tBegin.tv_usec;
    }
    
    useTime = deltaTimes[0];
    for(int i = 1; i < 10; i++)
    {
        if(useTime > deltaTimes[i])
            useTime = deltaTimes[i];
    }
    
    printf("Host memory time spent: %ldus\n", useTime);

其中，對gettimeofday的調用需要包含頭文件<sys/time.h>。這個函數所返回的時間可以精確到μs（微秒）。

在Intel Core i7 4650U, Intel Graphics 5000環境下，花費時間差不多，都是2.6ms（毫秒）。因此，在內核真正執行的時候為了清空這部分存儲空間的Cache，驅動還是要做點工作的。當然，驅動也可為這塊內存區域分配Write-Combined類型的存儲器，這樣主機端對這部分數據的訪問不會被Cache，盡管速度會慢很多，但是通過non-temporal Stream方式讀寫還是會很不錯。況且大部分OpenCL應用對同一塊內存數據的讀寫都只有一次，這么做也不會造成Cache污染。