OpenCL向量相加

原文http://www.olcf.ornl.gov/training_articles/opencl-vector-addition/

本文仅仅是为了学习OpenCL而做的的相关翻译。

由于原文中的例子不能在我的环境中运行，因此做了一些改动。

通过这个例子能很好地了解OpenCL的编程模型。

1. 简介

这个例子是表示了两个向量相加，可以认为是OpenCL中的"hello world"。为了使程序更容易理解，没有加入错误处理机制。

//vecAdd.c#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <CL/opencl.h>// OpenCL kernel. Each work item takes care of one element of c
const char *kernelSource =                                      "\n" \
"__kernel void vecAdd(  __global float *a,                       \n" \
"                       __global float *b,                       \n" \
"                       __global float *c,                       \n" \
"                       const unsigned int n)                    \n" \
"{                                                               \n" \
"    //Get our global thread ID                                  \n" \
"    int id = get_global_id(0);                                  \n" \
"                                                                \n" \
"    //Make sure we do not go out of bounds                      \n" \
"    if (id < n)                                                 \n" \
"        c[id] = a[id] + b[id];                                  \n" \
"}                                                               \n" \"\n" ;int main( int argc, char* argv[] )
{// 向量长度int n = 8;// 输入向量int *h_a;int *h_b;// 输出向量int *h_c;// 设备输入缓冲区cl_mem d_a;cl_mem d_b;// 设备输出缓冲区cl_mem d_c;cl_platform_id cpPlatform;        // OpenCL 平台cl_device_id device_id;           // device IDcl_context context;               // contextcl_command_queue queue;           // command queuecl_program program;               // programcl_kernel kernel;                 // kernel//（每个向量的字节数）size_t bytes = n*sizeof(int);//（为每个向量分配内存）h_a = (int*)malloc(bytes);h_b = (int*)malloc(bytes);h_c = (int*)malloc(bytes);//（初始化向量）int i;for( i = 0; i < n; i++ ){h_a[i] = i;h_b[i] = i;}size_t globalSize, localSize;cl_int err;//（每个工作组的工作节点数目）localSize = 2;//（所有的工作节点）globalSize = (size_t)ceil(n/(float)localSize)*localSize;printf("%d\n",globalSize);//（获得平台ID）err = clGetPlatformIDs(1, &cpPlatform, NULL);//（获得设备ID，与平台有关）err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);//（根据设备ID，得到上下文）context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);//（根据上下文，在设备上创建命令队列）queue = clCreateCommandQueue(context, device_id, 0, &err);//（根据OpenCL源程序创建计算程序）program = clCreateProgramWithSource(context, 1,(const char **) & kernelSource, NULL, &err);//（创建可执行程序）clBuildProgram(program, 0, NULL, NULL, NULL, NULL);//（在上面创建的程序中创建内核程序）kernel = clCreateKernel(program, "vecAdd", &err);//（分配设备缓冲）d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);// （将向量信息写入设备缓冲）err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,bytes, h_a, 0, NULL, NULL);err |= clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,bytes, h_b, 0, NULL, NULL);// （设置计算内核的参数）err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);err |= clSetKernelArg(kernel, 3, sizeof(int), &n);// （在数据集的范围内执行内核）Execute the kernel over the entire range of the data seterr = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,0, NULL, NULL);// （在读出结果之前，等待命令队列执行完毕）Wait for the command queue to get serviced before reading back resultsclFinish(queue);// （从设备缓冲区读出结果）Read the results from the deviceclEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,bytes, h_c, 0, NULL, NULL );//（输出读出的结果）float sum = 0;for(i=0; i<n; i++)printf("%d ",h_c[i]);// （释放资源）clReleaseMemObject(d_a);clReleaseMemObject(d_b);clReleaseMemObject(d_c);clReleaseProgram(program);clReleaseKernel(kernel);clReleaseCommandQueue(queue);clReleaseContext(context);//（释放内存）free(h_a);free(h_b);free(h_c);system("pause");return 0;
}

2. 基本解释

2.1 内核：

Kernel是OpenCL代码的核心。全部kernel必须要作为一个C字符串读入，最容易的方式就是将整个kernel用引号包起来，行尾回车。在真正的程序中，应该将kernel放在一个独立的文件中。

// OpenCL kernel. Each work item takes care of one element of c
const char *kernelSource =                                      "\n" \
"__kernel void vecAdd(  __global float *a,                       \n" \
"                       __global float *b,                       \n" \
"                       __global float *c,                       \n" \
"                       const unsigned int n)                    \n" \
"{                                                               \n" \
"    //Get our global thread ID                                  \n" \
"    int id = get_global_id(0);                                  \n" \
"                                                                \n" \
"    //Make sure we do not go out of bounds                      \n" \
"    if (id < n)                                                 \n" \
"        c[id] = a[id] + b[id];                                  \n" \
"}                                                               \n" \"\n" ;

查看一下这个简单的内核由什么内容组成：

__kernel void vecAdd(  __global float *a, __global float *b,                       __global float *c, const unsigned int n)

__kernel 指明这是一个OpenCL内核，__global 说明指针指向的是全局的设备内存空间，其它的就是C语言的函数的语法。kernel必须返回空类型。

int id = get_global_id(0);

得到第0维全局工作节点的ID。

if (id < n)    c[id] = a[id] + b[id];

工作组的数目必须是一个整数，或者每个工作组的工作节点数目必须能被全部工作节点数目整除。由于共组的的大小被用来协调性能，没有必要一定能被所有线程数目整除，所以通常启用的线程比所需要的线程多一些，并忽略掉多余的。在考察了问题域之后，就能访问、操作设备内存了。
2.2 内存：

    // 输入向量    int *h_a;int *h_b;// 输出向量    int *h_c;

// 设备输入缓冲区    cl_mem d_a;    cl_mem d_b;// 设备输出缓冲区    cl_mem d_c;

CPU和GPU有不同的内存空间，所以必须支持对内存分别引用，一个集市主机数组指针，另外一个集是设备内存的操作句柄。这儿我们用 h_和d_前缀来区分。

2.3 线程映射：

    //（每个工作组的工作节点数目）    localSize = 2;

//（所有的工作节点）    globalSize = (size_t)ceil(n/(float)localSize)*localSize;

为了将问题映射到底层硬件，必须指明局部的大小，和全局的大小。局部大小定义了工作组中节点的数目，子NVIDIA GPU上这相当于线程块内线程的数目。全局大小定义了所有启动的工作节点数目。localSize大小必须能被globalSize整除，所以我们计算了一个最小的整数能覆盖问题域，并且能被localSize整除。

2.4 环境配置：

    //（绑定平台）    err = clGetPlatformIDs(1, &cpPlatform, NULL);

每个硬件提供商都有不同的平台，在用之前就应该给定，这儿clGetPlatformIDs()会将cpPlatform赋予系统可用的平台。例如，如果系统包含了AMD CPU和NVIDIA GPU，且这两个平台都安装了合适的OpenCL驱动，那这里平台都是可用的。（注：要使用不同的平台驱动，必须安装相关的驱动，在本例中我安装了AMD(ATI)的app SDK v2.5和Intel的intel_ocl_sdk_1.5_runtime_setup，所以会有两个平台，但是由于我的ATI的显卡GPU不能被app SDK v2.5支持，所以的获得设备ID时没有用GPU设备，而是用了CPU设备。如果这里配置不正确，下面的可能就无法进行）

//（获得设备ID，与平台有关）err = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_CPU, 1, &device_id, NULL);

可以查询平台来得到它包含什么样的设备。在这个例子中，用枚举值CL_DEVICE_TYPE_CPU来查询平台上的CPU设备。

    //（根据上下文，在设备上创建命令队列）    queue = clCreateCommandQueue(context, device_id, 0, &err);

在使用OpenCL设备之前，必须要配置context，context被用来管理命令队列，内存和内核的活动。一个context可以包含不止一个设备。

    //（根据上下文，在设备上创建命令队列）    queue = clCreateCommandQueue(context, device_id, 0, &err);

命令队列被用来将命令从主机放入指定的设备。内存的转移和内核的活动都能被放入命令队列在合适的时候在指定的设备上执行。

2.5 编译内核：

    //（根据OpenCL源程序创建计算程序）    program = clCreateProgramWithSource(context, 1,                            (const char **) & kernelSource, NULL, &err);

//（创建可执行程序）    clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

//（在上面创建的程序中创建内核程序）    kernel = clCreateKernel(program, "vecAdd", &err);

为了保证对于大多数设备的可移植性，默认运行内核的方式就是用即时（Just-in-time）编译我们必须为给定上下文的设备准备源码。首先，创建程序，这是一个内核程序的的集合，然后根据程序来创建各自的内核程序。

2.6 准备数据：

    //（分配设备缓冲）    d_a = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);    d_b = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, NULL);    d_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bytes, NULL, NULL);

// （将向量信息写入设备缓冲）    err = clEnqueueWriteBuffer(queue, d_a, CL_TRUE, 0,                                   bytes, h_a, 0, NULL, NULL);    err |= clEnqueueWriteBuffer(queue, d_b, CL_TRUE, 0,                                   bytes, h_b, 0, NULL, NULL);

// （设置计算内核的参数）    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &d_a);    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &d_b);    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &d_c);    err |= clSetKernelArg(kernel, 3, sizeof(int), &n);

在启动内核之前，必须在设备和主机之间创建缓冲区，将主机数据绑定到新创建的设备缓冲区上，最后，设置内核参数。

2.7 启动内核：

// （在数据集的范围内执行内核）
err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, &localSize,
0, NULL, NULL);

一旦内存驻留在设备上之后，内核就能排队启动了。

2.8 取回结果：

    // （在读出结果之前，等待命令队列执行完毕）    clFinish(queue);

// （从设备缓冲区读出结果）    clEnqueueReadBuffer(queue, d_c, CL_TRUE, 0,                                bytes, h_c, 0, NULL, NULL );

可以进行阻断，直到所有的命令队列执行完毕，然后将设备上的结果取回到主机。

3. 运行环境

3.1 OpenCL:

　　AMD app sdk v2.5

　　intel_ocl_sdk_1.5_runtime

3.2 Visual Studio 2010 express

转载于:https://www.cnblogs.com/wangshide/archive/2011/11/04/2235204.html

OpenCL向量相加相关推荐

OpenCL 第5课:向量相加
OpenCL程序分为两个部份,一部份是内核代码,负责具体算法.另一部份是主程序负责初始化OpenCL和准备数据.主程序加载内核代码,并按照即定方法进行运算. 内核代码可以写在主程序里面,也可以写在另一 ...
OpenCL学习笔记——整体流程（向量相加）
OpenCL学习笔记--整体流程 OpenCL可以实现混合设备的并行计算,这些设备包括CPU,GPU,以及其他处理器,比如Cell处理器,DSP等.使用OpenCL编程,可以实现可以值得并行加速代码. ...
python 向量_关于Python中的向量相加和numpy中的向量相加效率对比
直接使用Python来实现向量的相加 # -*-coding:utf-8-*- #向量相加 def pythonsum(n): a = range(n) b = range(n) c = [] for ...
python坐标系向量分量_关于Python中的向量相加和numpy中的向量相加效率对比
直接使用Python来实现向量的相加 # -*-coding:utf-8-*- #向量相加 def pythonsum(n): a = range(n) b = range(n) c = [] for ...
像素颜色和颜色向量相加相乘的理解
首先计算机图形学的一条原则是,如果它看上去是对的,那它就是对的.所以,对于一些计算模型来说,最重要的是最后拟合的效果.它不一定符合物理,但通常计算都以物理上的考虑为基础. 像素颜色和相关运算的理解我 ...
向量相加后是否与目标向量平行
一.这是一个验证向量组中是否含有相加能与目标向量平行的代码问题描述: 给定n个向量的起点和终点x1,y1,x2,y2,再给出一个目标向量,求能否由n个向量中的两个相加构造出一个与目标向量平行的向量( ...
r语言逻辑向量相加_R语言基础教程——第3章：数据结构——向量
如果学过像JAVA或者C这样的高级语言,都知道,数据类型的概念,包括,整数型.浮点型.字符串.布尔类型.这些语言中,定义变量需要定义数据类型,而在R中不需要.只需要直接赋值即可.在给变量赋值时,R中可 ...
要点初见：OpenCL 2.0 异构计算 [第三版] 知识点整理
Word 版Github项目地址: https://github.com/BingLiHanShuang/Chinese-Knowledge-Collation-of-Heterogeneous-Co ...
OpenCL 第6课:矩阵转置
上一节我们写了个一维向量相加的程序.这节我们来看一个4×4矩阵转置程序. 4X4矩阵我们采用二维数组进行存储,在程序设计上,我们让转置过程分4次转置完成,就是一次转一行.注意这里的OpenCL的工作维 ...

OpenCL向量相加

OpenCL向量相加相关推荐

最新文章

热门文章