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

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/461670.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

AndroidManifest.xml文件解析(转帖)

AndroidManifest.xml文件解析(转帖)

原帖地址&#xff1a;http://www.cnblogs.com/pilang/archive/2011/04/20/2022932.html 一、关于AndroidManifest.xml AndroidManifest.xml 是每个android程序中必须的文件。它位于整个项目的根目录&#xff0c;描述了package中暴露的组件&#xff08;activities, services, 等…
阅读更多...
windows 7系统搭建本地SVN服务器的过程

windows 7系统搭建本地SVN服务器的过程

1、安装svn:TortoiseSVN-1.7.9.23248-x64-svn-1.7.6; 1、在PC机的D盘创建如下目录&#xff1a;D:\svn\project\workspace; 2、右键点击目录workspace文件&#xff0c;选择Tortoise->Create repository here,即在此创建版本库&#xff08;Y&#xff09;; 然后你就会看到D:\sv…
阅读更多...
x210开发板的BSP（其中使用buildroot文件夹建立rootfs）

x210开发板的BSP（其中使用buildroot文件夹建立rootfs）

以下内容源于朱有鹏嵌入式课程的学习与整理&#xff0c;如有侵权请告知删除。 参考博客&#xff1a;buildroot详解和分析_Alex-wu的博客-CSDN博客_buildroot 板级支持包&#xff08;BSP&#xff0c;Board Support Package&#xff09;&#xff0c;是由引导程序&#xff08;Boo…
阅读更多...
Android中Activity启动模式详解

Android中Activity启动模式详解

在Android中每个界面都是一个Activity&#xff0c;切换界面操作其实是多个不同Activity之间的实例化操作。在Android中Activity的启动模式决定了Activity的启动运行方式。 Android总Activity的启动模式分为四种&#xff1a; Activity启动模式设置&#xff1a; <activity and…
阅读更多...
EDM的九大用途盘点

EDM的九大用途盘点

对于什么是EDM&#xff0c;前面博主的博文已经有所介绍。那么&#xff0c;EDM的九大用途是什么呢&#xff1f;本文博主就为大家介绍一下。 交叉营销交叉营销是通过把时间&#xff0c;金钱、构想、活动或演示空间等资源整合&#xff0c;为任何企业&#xff0c;包括家庭式小企业、…
阅读更多...
分析根文件系统中的/etc/inittab文件

分析根文件系统中的/etc/inittab文件

以下内容源于朱有鹏嵌入式课程的学习与整理&#xff0c;如有侵权请告知删除。 1、文件简介 &#xff08;1&#xff09;/etc/inittab文件属于运行时配置文件。 &#xff08;2&#xff09;这个文件是文本格式的&#xff0c;即内容由一系列遵照某格式的字符组成。 &#xff08;3&…
阅读更多...
MySQL导入.sql文件及常用命令

MySQL导入.sql文件及常用命令

MySQL导入.sql文件及常用命令 在MySQL Qurey Brower中直接导入*.sql脚本&#xff0c;是不能一次执行多条sql命令的&#xff0c;在mysql中执行sql文件的命令&#xff1a; mysql> source d:/myprogram/database/db.sql; 另附mysql常用命令&#xff1a; 一) 连接MYSQL&…
阅读更多...
第一次软工作业展示——潘学

第一次软工作业展示——潘学

第一次软工作业完成啦&#xff01; 回首这个作业的完成过程&#xff0c;我是很有收获。这个作业有几个难点&#xff1a;1、在给定目录下读取TXT文件的内容&#xff1b;2、从读到的内容中分析出单词&#xff1b;3、统计单词的出现频率并输出。 我之前只学习过C和java&#xff0c…
阅读更多...
【0802 | Day 7】Python进阶（一）

【0802 | Day 7】Python进阶（一）

目 录 数字类型的内置方法 一、整型内置方法&#xff08;int&#xff09; 二、浮点型内置方法&#xff08;float&#xff09; 字符串类型内置方法 一、字符串类型内置方法&#xff08;str&#xff09; 二、常用操作和内置方法 优先掌握&#xff1a; 1.索引取值 2.切片 3.长度le…
阅读更多...
bzoj 1084 DP

bzoj 1084 DP

首先对于m1的情况非常容易处理&#xff08;其实这儿因为边界我错了好久。。。&#xff09;&#xff0c;直接DP就好了&#xff0c;设f[i][k]为这个矩阵前i个选k个矩阵的最大和&#xff0c;那么f[i][k]max(f[j][k-1]sum[j1][i])&#xff0c;那么对于m2的时候类似与m1的时候&#…
阅读更多...
uboot源码——命令体系

uboot源码——命令体系

以下内容源于朱有鹏嵌入式课程的学习&#xff0c;如有侵权&#xff0c;请告知删除。 参考资料&#xff1a;http://www.cnblogs.com/biaohc/p/6394710.html 一、uboot命令体系基础 1、使用uboot命令 uboot启动后进入命令行环境&#xff0c;在此输入命令按回车结束&#xff0…
阅读更多...
RCP：如何移除Toolbar中的Quick Access

RCP：如何移除Toolbar中的Quick Access

问题 自4.x开始&#xff0c;Quick Access搜索框成为Toolbar的“标准装备”&#xff0c;一般删除Actionset的方式似乎不起作用&#xff0c;通过Quick Access&#xff0c;用户很容易访问到RCP程序本来想隐藏的功能。 解决方法 在WorkbenchWindowAdvisor的openIntro中加入以下代码…
阅读更多...
XMPP文件传输（XEP-0096协议说明）

XMPP文件传输（XEP-0096协议说明）

XMPP XEP-0096协议是XMPP中的文件传输协议。 关于文件传输&#xff0c;在xmpp协议中有不少协议可以实现&#xff0c;而XEP-0096协议是其中非常简单的一个协议。由于邮件被删&#xff0c;我的代码demo丢失&#xff0c;因此只能在这里给大家讲一下其中的逻辑实现&#xff0…
阅读更多...
[笔记]VI编辑器的学习

[笔记]VI编辑器的学习

来源&#xff1a;http://team.youthol.cn/?p453 2013-03-22 09:20:00 在Vim中利用替换功能就可以将“^M”都删掉&#xff0c;键入如下替换命令行&#xff1a;:% s/\r//g 就可以类似的“&#xff5c;&#xff5c;”都删掉&#xff0c;键入如下替换命令行&#xff1a;%s/Tab键…
阅读更多...
uboot源码——环境变量

uboot源码——环境变量

以下内容源于朱有鹏嵌入式课程的学习&#xff0c;如有侵权&#xff0c;请告知删除。 参考资料&#xff1a;http://www.cnblogs.com/biaohc/p/6398515.html。 一、uboot的环境变量基础 1、环境变量的作用 在不改变源码、不用重新编译的情况下&#xff0c;可以通过设置环境变量…
阅读更多...
Linq To Sql进阶系列 -目录导航

Linq To Sql进阶系列 -目录导航

博客园CLR基础研究团队|CLR团队精品系列|C# 3.0专题 [Linq To Sql进阶系列] 目录导航 1 Linq To Sql进阶系列&#xff08;一&#xff09;-从映射讲起 本系列&#xff0c;或多或少&#xff0c;直接或间接依赖入门系列知识。但&#xff0c;依然追求独立成章。因本文作者水平有限&…
阅读更多...
uboot源码——mmc驱动分析

uboot源码——mmc驱动分析

以下内容源于朱有鹏《物联网大讲坛》课程的学习&#xff0c;以及博客http://www.cnblogs.com/biaohc/p/6409197.html的学习整理&#xff0c;如有侵权&#xff0c;请告知删除。 一、uboot与linux驱动 1、uboot是裸机程序 狭义的驱动的概念&#xff1a;操作系统中用来具体操控硬…
阅读更多...
VB与Java颜色值的转换

VB与Java颜色值的转换

正常的 RGB 颜色的有效范围&#xff0c;是从 0 到 16,777,215 (&HFFFFFF&)。每种颜色的设置值&#xff08;属性或参数&#xff09;都是一个四字节的整数。对于这个范围内的数&#xff0c;其高字节都是 0&#xff0c;而低三个字节&#xff0c;从最低字节到第三个字节&am…
阅读更多...
Dreamweaver MX显示汉字为乱码的解决方法

Dreamweaver MX显示汉字为乱码的解决方法

推荐几种解决方法:a.在“编辑”&#xff0d;“首选参数”中设置“新建文档”->默认编码&#xff1a;utf-8或者gb2312&#xff08;取决于你的网页编码&#xff09;&#xff0c;并勾选“当打开未指定编码的现有文件时使用”;此时每次打开文件时都没有乱码了&#xff0c;也不额…
阅读更多...
RDIFramework.NET（.NET快速开发框架） 答客户问（2014-02-23）

RDIFramework.NET（.NET快速开发框架） 答客户问（2014-02-23）

1、框架的部署安装&#xff0c;服务器端和客户端 答&#xff1a;开发版以上版本支持SOA模式&#xff0c;也即真正的面向服务端的模式&#xff0c;在实际使用过程中&#xff0c;可根据项目的实际需要&#xff0c;来选择性的进行部署&#xff08;直连模式或SOA模式&#xff09;&a…
阅读更多...
最新文章

Copyright @ 2022~2023