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特性越多、越高级,相应的硬件也越复杂,成本也越高。

赞(0)
未经允许不得转载:极客笔记 » OpenCL 共享虚拟存储器示例
分享到: 更多 (0)

评论 抢沙发

  • 昵称 (必填)
  • 邮箱 (必填)
  • 网址