OpenCL 共享虚拟存储器示例|极客笔记

链表之类的数据结构中，结构体中有一个指向其他位置的指针。在例子中，数据结构体如下：

typedef struct _Element
{
    float ＊internal;             //指向结构体数组内的其他位置
    float ＊external;             //指向其他数组的位置
    float value;
} Element;

主机代码中会分配两个SVM：inputElements和inputFloats，SVM inputElements中保存的是Element结构体数据，SVM inputFloats中保存的是浮点型数据。SVM inputElements的Element结构体元素external指针指向SVM inputFloats中的位置。
笔者实验GPU为AMD R9 285，支持粗粒地缓冲SVM和细粒度缓冲SVM。在本节中，笔者就针对这两种情况分别列举两个例子。对于这两个例子，内核代码都是一样的，内核代码如下：

//SVMbasic.cl文件内容
typedef struct _Element
{
    global float ＊internal;
    global float ＊external;
    float value;
} Element;
kernel void svmbasic (global Element ＊elements,
global float ＊dst)
{
    int id = (int)get_global_id(0);
    float internalElement =
                            ＊ (elements[id].internal);     //访问本SVM其他位置
    float externalElement = ＊(elements[id].external); //访问其他SVM值
    dst[id] = internalElement + externalElement;
}

粗粒度缓冲SVM

创建粗粒度缓冲SVM时，主机代码如下：

int SVMCoarseGrained(cl_context context,
                        cl_command_queue cmdqueue,
                        cl_kernel kernel)
{
    cl_int err;
    int size = 16;
//创建三个粗粒度缓冲SVM
Element ＊inputElements = (Element ＊)clSVMAlloc(context,
                                  CL_MEM_READ_ONLY,
                                  sizeof(Element) ＊ size, 0);
    float ＊inputFloats = (float ＊)clSVMAlloc(context,
                              CL_MEM_READ_ONLY,
                              sizeof(float) ＊ size, 0);
    float ＊output = (float ＊)clSVMAlloc(context,
                                                CL_MEM_READ_ONLY,
                                                sizeof(float) ＊ size, 0);
    //粗粒度缓冲SVM映射
    err = clEnqueueSVMMap(cmdqueue, CL_TRUE,
                              CL_MAP_WRITE,
                              inputElements,
                              sizeof(Element) ＊ size,
                              0, 0, 0);
err = clEnqueueSVMMap(cmdqueue, CL_TRUE,CL_MAP_WRITE, inputFloats,
                              sizeof(float) ＊ size,
                              0, 0, 0);
//初始化输入的SVM值
for(int i = 0; i 〈 size; i++)
{
        inputElements[i].value = i;
        inputFloats[i] = i + size;
}
for(int i = 0; i 〈 size; i++)
{
        int index = rand_index(i);    //产生0～size-1间的随机数
        inputElements[i].internal = &(inputElements[index].value);
        inputElements[i].external = &(inputFloats[index]);
}
//粗粒度缓冲SVM解映射
err = clEnqueueSVMUnmap(cmdqueue, inputFloats, 0, 0, 0);
err = clEnqueueSVMUnmap(cmdqueue, inputElements, 0, 0, 0);
//粗粒度缓冲SVM inputElements和output作为参数传入
err = clSetKernelArgSVMPointer(kernel, 0, inputElements);
err = clSetKernelArgSVMPointer(kernel, 1, output);
//内核会使用inputFloats中的值,但不作为参数传入
err = clSetKernelExecInfo(kernel,
    CL_KERNEL_EXEC_INFO_SVM_PTRS,
    sizeof(inputFloats), &inputFloats);
size_t globalsize = size;
err = clEnqueueNDRangeKernel(cmdqueue, kernel, 1, 0,
        &globalsize, NULL,
        NULL, NULL, NULL);
err = clFinish(cmdqueue);
err = clEnqueueSVMMap(cmdqueue, CL_TRUE,
CL_MAP_READ, output,
sizeof(float) ＊ size, 0, 0, 0);
//验证计算结果
for(int i = 0; i 〈 size; i++)
{
float expectValue = ＊(inputElements[i].external) + ＊
  (inputElements[i].internal);
if(expectValue != output[i])
{
printf("Mistch at %d\n", i);
break;
}
}
err = clEnqueueSVMUnmap(cmdqueue, output, 0, 0, 0);
clSVMFree(context, output);
clSVMFree(context, inputFloats);
clSVMFree(context, inputElements);
}

在创建SVM时，没有指定为细粒度缓冲SVM，则创建的SVM为默认的粗粒度缓冲SVM。对于粗粒度缓冲SVM，主机与设备间同步数据需要使用映射/解映射操作。对于SVM，可以使用clSetKernelArgSVMPointer函数作为内核参数传入。对于内核中会使用的SVM数据，但又不是作为参数传入内核，需要使用clSetKernelExecInfo函数。

细粒度缓冲SVM

创建细粒度缓冲SVM时，主机代码如下：

int SVMCoarseGrained(cl_context context,
                        cl_command_queue cmdqueue,
                        cl_kernel kernel){
    cl_int err;
    int size = 16;
    //创建细粒度缓冲SVM
    Element ＊inputElements = (Element ＊)clSVMAlloc(context,
                                  CL_MEM_READ_ONLY |
                                  CL_MEM_SVM_FINE_GRAIN_BUFFER,
                                  sizeof(Element) ＊ size, 0);

float ＊inputFloats = (float ＊)clSVMAlloc(context,
                              CL_MEM_READ_ONLY |
                              CL_MEM_SVM_FINE_GRAIN_BUFFER,
                              sizeof(float) ＊ size, 0);
    float ＊output = (float ＊)clSVMAlloc(context,
                                                CL_MEM_READ_ONLY |
                                                CL_MEM_SVM_FINE_GRAIN_BUFFER,
                                                sizeof(float) ＊ size, 0);
    //初始化输入的SVM值
    for(int i = 0; i 〈 size; i++)
    {
        inputElements[i].value = i;
        inputFloats[i] = i + size;
    }
    for(int i = 0; i 〈 size; i++)
    {
        int index = rand_index(i);
        inputElements[i].internal = &(inputElements[index].value);
        inputElements[i].external = &(inputFloats[index]);
    }
    err = clSetKernelArgSVMPointer(kernel, 0, inputElements);
    err = clSetKernelArgSVMPointer(kernel, 1, output);
    err = clSetKernelExecInfo(kernel,
                                    CL_KERNEL_EXEC_INFO_SVM_PTRS,
                                    sizeof(inputFloats), &inputFloats);
    size_t globalsize = size;
    err = clEnqueueNDRangeKernel(cmdqueue, kernel, 1, 0,
                                        &globalsize, NULL,
                                        NULL, NULL, NULL);
    err = clFinish(cmdqueue);
    //验证计算结果
    for(int i = 0; i 〈 size; i++)
    {
        float expectValue = ＊(inputElements[i].external) +
                            ＊(inputElements[i].internal);
        if(expectValue != output[i])
       {
        printf("Mismatch at %d\n", i);
        break;
    }
 }
 clSVMFree(context, output);
 clSVMFree(context, inputFloats);
 clSVMFree(context, inputElements);
return 0;
}

在创建SVM时指定为CL_MEM_SVM_FINE_GRAIN_BUFFER，即为细粒度缓冲SVM。对于细粒度缓冲SVM，主机与设备间可以无映射直接访问。对比粗粒度缓冲SVM操作代码，细粒度缓冲SVM操作代码更简洁。如果想要硬件支持的SVM特性越多、越高级，相应的硬件也越复杂，成本也越高。