OpenCL概述 续篇

摘要：Kernel就是在程序代码中的一个函数，这个函数能在OpenCL设备上执行。一个Kernel对象就是kernel函数以及其相关的输入参数。

Kernel对象：

Kernel就是在程序代码中的一个函数，这个函数能在OpenCL设备上执行。一个Kernel对象就是kernel函数以及其相关的输入参数。

Kernel对象通过程序对象以及指定的函数名字创建。注意：函数必须是程序源代码中存在的函数。

运行时编译：

在运行时，编译程序和创建kernel对象是有时间开销的，但这样比较灵活，能够适应不同的OpenCL硬件平台。程序动态编译一般只需一次，而Kernel对象在创建后，可以反复调用。

创建Kernel后，运行Kernel之前，我们还要为Kernel对象设置参数。我们可以在Kernel运行后，重新设置参数再次运行。

arg_index指定该参数为Kernel函数中的第几个参数(比如第一个参数为0，第二个为1,…)。内存对象和单个的值都可以作为Kernel参数。下面是2个设置Kernel参数的例子：

clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&d_iImage);

clSetKernelArg(kernel, 1, sizeof(int), (void*)&a);

在Kernel运行之前，我们先看看OpenCL中的线程结构：

大规模并行程序中，通常每个线程处理一个问题的一部分，比如向量加法，我们会把两个向量中对应的元素加起来，这样，每个线程可以处理一个加法。

下面我看一个16个元素的向量加法：两个输入缓冲A、B，一个输出缓冲C

在这种情况下，我们可以创建一维的线程结构去匹配这个问题。

每个线程把自己的线程id作为索引，把相应元素加起来。

OpenCL中的线程结构是可缩放的，Kernel的每个运行实例称作WorkItem(也就是线程），WorkItem组织在一起称作WorkGroup，OpenCL中，每个Workgroup之间都是相互独立的。

通过一个global id(在索引空间，它是唯一的）或者一个workgroup id和一个work group内的local id，我就能标定一个workitem。

在kernel函数中，我们能够通过API调用得到global id以及其他信息：

get_global_id(dim)

get_global_size(dim)

这两个函数能得到每个维度上的global id。

get_group_id(dim)

get_num_groups(dim)

get_local_id(dim)

get_local_size(dim）

这几个函数用来计算group id以及在group内的local id。

get_global_id(0) = column, get_global_id(1) = row

get_num_groups(0) * get_local_size(0) == get_global_size(0)

摘要：OpenCL的内存模型定义了各种各样内存类型，各种内存模型之间有层级关系。各种内存之间的数据传输必须是显式进行的，比如从host memory到device memory，从global memory到local memory等等。

OpenCL内存模型

OpenCL的内存模型定义了各种各样内存类型，各种内存模型之间有层级关系。各种内存之间的数据传输必须是显式进行的，比如从host memory到device memory，从global memory到local memory等等。

WorkGroup被映射到硬件的CU上执行（在AMD 5xxx系列显卡上，CU就是simd，一个simd中有16个pe，或者说是stream core），OpenCL并不提供各个workgroup之间的一致性，如果我们需要在各个workgroup之间共享数据或者通信之类的，要自己通过软件实现。

Kernel函数的写法

每个线程（workitem）都有一个kenerl函数的实例。下面我们看下kernel的写法：

 1: __kernel void vecadd(__global const float* A, __global const float* B, __global float* C)

 2: {

 3: int id = get_global_id(0);

 4: C[id] = A[id] + B[id];

 5: }

每个Kernel函数都必须以__kernel开始，而且必须返回void。每个输入参数都必须声明使用的内存类型。通过一些API，比如get_global_id之类的得到线程id。

内存对象地址空间标识符有以下几种：

__global – memory allocated from global address space

__constant – a special type of read-only memory

__local – memory shared by a work-group

__private – private per work-item memory

__read_only/__write_only – used for images

Kernel函数参数如果是内存对象，那么一定是__global,__local或者constant。

运行Kernel

首先要设置线程索引空间的维数以及workgroup大小等。

我们通过函数clEnqueueNDRangeKerne把Kernel放在一个队列里，但不保证它马上执行，OpenCL driver会管理队列，调度Kernel的执行。注意：每个线程执行的代码都是相同的，但是它们执行数据却是不同的。

该函数把要执行的Kernel函数放在指定的命令队列中，globald大小（线程索引空间）必须指定，local大小（work group）可以指定，也可以为空。如果为空，则系统会自动根据硬件选择合适的大小。event_wait_list用来选定一些events，只有这些events执行完后，该kernel才可能被执行，也就是通过事件机制来实现不同kernel函数之间的同步。

当Kernel函数执行完毕后，我们要把数据从device memory中拷贝到host memory中去。

释放资源：

大多数的OpenCL资源都是指针，不使用的时候需要释放掉。当然，程序关闭的时候这些对象也会被自动释放掉。

释放资源的函数是：clRelase{Resource} ，比如: clReleaseProgram(), clReleaseMemObject()等。

错误捕捉：

如果OpenCL函数执行失败，会返回一个错误码，一般是个负值，返回0则表示执行成功。我们可以根据该错误码知道什么地方出错了，需要修改。错误码在cl.h中定义，下面是几个错误码的例子.

CL_DEVICE_NOT_FOUND -1

CL_DEVICE_NOT_AVAILABLE -2

CL_COMPILER_NOT_AVAILABLE -3

CL_MEM_OBJECT_ALLOCATION_FAILURE -4

…

下面是一个OpenCL机制的示意图

程序模型

数据并行：work item和内存对象元素之间是一一映射关系；workgroup可以显示指定，也可以隐式指定。

任务并行：kernel的执行独立于线程索引空间；用其他方法表示并行，比如把不同的任务放入队列，用设备指定的特殊的向量类型等等。

同步：workgroup内work item之间的同步；命令队列中不同命令之间的同步。

完整代码如下：

 1: #include "stdafx.h"

 2: #include <CL/cl.h>

 3: #include <stdio.h>

 4: #include <stdlib.h>

 5: #include <time.h>

 6: #include <iostream>

 7: #include <fstream>

 8: 

 9: using namespace std;

 10: #define NWITEMS 262144

 11: 

 12: #pragma comment (lib,"OpenCL.lib")

 13: 

 14: //把文本文件读入一个string中

 15: int convertToString(const char *filename, std::string& s)

 16: {

 17: size_t size;

 18: char* str;

 19: 

 20: std::fstream f(filename, (std::fstream::in | std::fstream::binary));

 21: 

 22: if(f.is_open())

 23: {

 24: size_t fileSize;

 25: f.seekg(0, std::fstream::end);

 26: size = fileSize = (size_t)f.tellg();

 27: f.seekg(0, std::fstream::beg);

 28: 

 29: str = new char[size+1];

 30: if(!str)

 31: {

 32: f.close();

 33: return NULL;

 34: }

 35: 

 36: f.read(str, fileSize);

 37: f.close();

 38: str[size] = '\0';

 39: 

 40: s = str;

 41: delete[] str;

 42: return 0;

 43: }

 44: printf("Error: Failed to open file %s\n", filename);

 45: return 1;

 46: }

 47: 

 48: int main(int argc, char* argv[])

 49: {

 50: //在host内存中创建三个缓冲区

 51: float *buf1 = 0;

 52: float *buf2 = 0;

 53: float *buf = 0;

 54: 

 55: buf1 =(float *)malloc(NWITEMS * sizeof(float));

 56: buf2 =(float *)malloc(NWITEMS * sizeof(float));

 57: buf =(float *)malloc(NWITEMS * sizeof(float));

 58: 

 59: //初始化buf1和buf2的内容

 60: int i;

 61: srand( (unsigned)time( NULL ) );

 62: for(i = 0; i < NWITEMS; i++)

 63: buf1[i] = rand()%65535;

 64: 

 65: srand( (unsigned)time( NULL ) +1000);

 66: for(i = 0; i < NWITEMS; i++)

 67: buf2[i] = rand()%65535;

 68: 

 69: for(i = 0; i < NWITEMS; i++)

 70: buf[i] = buf1[i] + buf2[i];

 71: 

 72: cl_uint status;

 73: cl_platform_id platform;

 74: 

 75: //创建平台对象

 76: status = clGetPlatformIDs( 1, &platform, NULL );

 77: 

 78: cl_device_id device;

 79: 

 80: //创建GPU设备

 81: clGetDeviceIDs( platform, CL_DEVICE_TYPE_GPU,

 82: 1,

 83: &device,

 84: NULL);

 85: //创建context

 86: cl_context context = clCreateContext( NULL,

 87: 1,

 88: &device,

 89: NULL, NULL, NULL);

 90: //创建命令队列

 91: cl_command_queue queue = clCreateCommandQueue( context,

 92: device,

 93: CL_QUEUE_PROFILING_ENABLE, NULL );

 94: //创建三个OpenCL内存对象，并把buf1的内容通过隐式拷贝的方式

 95: //拷贝到clbuf1,buf2的内容通过显示拷贝的方式拷贝到clbuf2

 96: cl_mem clbuf1 = clCreateBuffer(context,

 97: CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,

 98: NWITEMS*sizeof(cl_float),buf1,

 99: NULL );

 100: 

 101: cl_mem clbuf2 = clCreateBuffer(context,

 102: CL_MEM_READ_ONLY ,

 103: NWITEMS*sizeof(cl_float),NULL,

 104: NULL );

 105: 

 106: status = clEnqueueWriteBuffer(queue, clbuf2, 1,

 107: 0, NWITEMS*sizeof(cl_float), buf2, 0, 0, 0);

 108: 

 109: cl_mem buffer = clCreateBuffer( context,

 110: CL_MEM_WRITE_ONLY,

 111: NWITEMS * sizeof(cl_float),

 112: NULL, NULL );

 113: 

 114: const char * filename = "add.cl";

 115: std::string sourceStr;

 116: status = convertToString(filename, sourceStr);

 117: const char * source = sourceStr.c_str();

 118: size_t sourceSize[] = { strlen(source) };

 119: 

 120: //创建程序对象

 121: cl_program program = clCreateProgramWithSource(

 122: context,

 123: 1,

 124: &source,

 125: sourceSize,

 126: NULL);

 127: //编译程序对象

 128: status = clBuildProgram( program, 1, &device, NULL, NULL, NULL );

 129: if(status != 0)

 130: {

 131: printf("clBuild failed:%d\n", status);

 132: char tbuf[0x10000];

 133: clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0x10000, tbuf, NULL);

 134: printf("\n%s\n", tbuf);

 135: return -1;

 136: }

 137: 

 138: //创建Kernel对象

 139: cl_kernel kernel = clCreateKernel( program, "vecadd", NULL );

 140: //设置Kernel参数

 141: cl_int clnum = NWITEMS;

 142: clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*) &clbuf1);

 143: clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*) &clbuf2);

 144: clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*) &buffer);

 145: 

 146: //执行kernel

 147: cl_event ev;

 148: size_t global_work_size = NWITEMS;

 149: clEnqueueNDRangeKernel( queue,

 150: kernel,

 151: 1,

 152: NULL,

 153: &global_work_size,

 154: NULL, 0, NULL, &ev);

 155: clFinish( queue );

 156: 

 157: //数据拷回host内存

 158: cl_float *ptr;

 159: ptr = (cl_float *) clEnqueueMapBuffer( queue,

 160: buffer,

 161: CL_TRUE,

 162: CL_MAP_READ,

 163: 0,

 164: NWITEMS * sizeof(cl_float),

 165: 0, NULL, NULL, NULL );

 166: //结果验证，和cpu计算的结果比较

 167: if(!memcmp(buf, ptr, NWITEMS))

 168: printf("Verify passed\n");

 169: else printf("verify failed");

 170: 

 171: if(buf)

 172: free(buf);

 173: if(buf1)

 174: free(buf1);

 175: if(buf2)

 176: free(buf2);

 177: 

 178: //删除OpenCL资源对象

 179: clReleaseMemObject(clbuf1);

 180: clReleaseMemObject(clbuf2);

 181: clReleaseMemObject(buffer);

 182: clReleaseProgram(program);

 183: clReleaseCommandQueue(queue);

 184: clReleaseContext(context);

 185: return 0;

 186: }

 187: 

也可以在http://code.google.com/p/imagefilter-opencl/downloads/detail?name=amdunicourseCode1.zip&can=2&q=#makechanges上下载完整版本。

原文作者：迈克老狼