OpenCL 编写内核代码

对于刚接触OpenCL编程的读者而言,如何编写内核代码是他们很关心的一个问题。本文讲解如何编写内核代码。从OpenCL 创建Program对象中可以知道,OpenCL程序对象可由源码创建。

下面通过一个简单的例子来了解如何使用OpenCL C编写一个内核程序将两个浮点数组相加。串行版本代码求和时需要通过一个循环将两个数组中的各个元素相加:

void add_cpu(int n, const float *a, const float *b ,
              float *result)
{
    int i;
    for(i = 0; i 〈 n; i++)
    {
        result[i] = a[i] + b[i];
    }
}

使用OpenCL C并行代码如下:

kernel void add_gpu(global const float *a, global const float *b,
                        global float *result)
{
    int id = get_global_id(0);
    result[id] = a[id] + _b[id];
}

add_gpu函数声明使用kernel或(_kernel)修饰符来告诉编译器这是一个OpenCL C内核函数。add_gpu内核只包含计算单个元素求和的代码,也就是串行代码中的内循环,这是因为在数据并行模式下,同时会有多个工作项来参与计算,每个工作项只参与跟自己ID相关的计算。在代码中使用内建函数get_global_id()来获得当前工作项的ID。关于内建工作项函数,以及不同设备执行模型,这里就不展开。在此读者需要记住的是,内核函数使用kernel修饰符来限定,并且内核函数不能有返回值,只能是void类型。

通过上述的例子,我们知道了如何正确书写一个内核函数的规则,但内核源码到底应该用什么方式来保存呢?在OpenCL开发中,对于内核源码的处理一般有下列两种方式。

内核代码保存

我们可以把内核代码保存在一个文本文件中,该文本文件后缀习惯为“cl”,如kernel. cl。对于kernel.cl文件,与我们一般的文件读取处理一样。在kernel.cl中的内容书写格式与C语言风格类似。

例如,看一个规约的内核代码,文件Reduction_kernel.cl内容如下:

__kernel void reduce(__global uint4 *input,
                            __global uint4 *output,
                            __local uint4 *sdata)
{
    //load shared mem
    unsigned int tid = get_local_id(0);
    unsigned int bid = get_group_id(0);
    unsigned int gid = get_global_id(0);
    unsigned int localSize = get_local_size(0);
    unsigned int stride = gid * 2;
    sdata[tid] = input[stride] + input[stride + 1];
    barrier(CLK_LOCAL_MEM_FENCE);
    //do reduction in shared memory
    for(unsigned int s = localSize >> 1; s > 0; s >>= 1)
    {
        if(tid 〈 s)
        {
              sdata[tid] += sdata[tid + s];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }
    //write result for this block to global mem
    if(tid == 0) output[bid] = sdata[0];
}

代码书写规范与C语言一样。在程序中需要读取Reduction_kernel.cl文件内容。文件后缀名之所以为“cl”,只是习惯而已,不过很多OpenCL代码编辑器能自动识别.cl文件从而可以显示相应的语法高亮(如Xcode就是其中之一)。但是需要注意的是,如果采用这种方式,内核代码文件需要跟随程序一起发布。

如果开发者不想把OpenCL源代码直接暴露给用户,那么就需要在应用发布前使用某些工具来对.cl文件进行加密。笔者常用的是在开发阶段采用此种方式,最后应用程序发布时把源代码编译过后保存为二进制文件。

字符串保存

如果采用这种方式,需要把.cl文件内容读取到字符串中。所以内核代码也可以直接写在字符串中。例如:

const char *src[] =
{
    " __kernel void redution( \n"
    "    __global int *data,     \n"
    "    __global int *output,   \n"
    "    __local int *data_local   \n"
    "    ) \n"
    " {   \n"
    "    int gid=get_group_id(0);   \n"
    "    int tid=get_global_id(0);    \n"
    "    int size=get_local_size(0);   \n"
    "    int id=get_local_id(0);     \n"
    "    data_local[id]=data[tid];   \n"
    "    barrier(CLK_LOCAL_MEM_FENCE);   \n"
    "    for(int i=size/2;i>0;i>>=1){    \n"
    "         if(id〈i){   \n"
    "             data_local[id]+=data_local[id+i];   \n"
    "         }   \n"
    "         barrier(CLK_LOCAL_MEM_FENCE);   \n"
    "    }    \n"
    "    if(id==0){    \n"
    "         output[gid]=data_local[0];   \n"
    "    }    \n"
    " }   \n"
};

字符串中内容格式与写C代码格式一致。对于每一行末尾需要添加“\n”换行符。

两种编写内核代码方式作用一样。对于第一种方式,在写代码过程中比较方便,与我们常规的C代码风格一致,但源码在后期处理过程中需要增加文件读取;对于第二种方式,在写代码过程中有点烦琐,但源码可以直接交给函数clCreateProgramWithSource()处理。在这两种方式中,笔者更常用第一种方式。

赞(1)
未经允许不得转载:极客笔记 » OpenCL 编写内核代码
分享到: 更多 (0)

评论 抢沙发

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