引言
OpenCL简介
- kernel和OpenCL执行模型
- 平台和设备
- - 主机-设备交互
- 执行环境
- - 上下文
  - 命令队列
  - 事件
  - 内存对象
  - - buffer
    - image
  - flush命令和finish命令
  - 新建一个OpenCL程序对象
  - OpenCL的kernel
- 内存模型
- 编写kernel
- 本章向量相加程序源码
OpenCL的并发与执行模型
- OpenCL同步：kernel，fence和barrier
- 队列与全局同步
- - OpenCL内存一致性
  - 事件
  - - 多个设备上的队列
    - 事件回调
- 主机端内存模型
- - buffer对象
  - image对象

引言

以下仅为书中一些基本概念，具体应参考官方specification文档、C++ binding文档和api参考文档。以下OpenCL基于OpenCL 1.1

OpenCL简介

kernel和OpenCL执行模型

kernel是OpenCL程序在设备上实际执行的代码，每个kernel称为work-item，多个work-item组成一个work-group
kernel执行期间，需指定一个n维索引空间（NDRange）来声明需要新建多少个work-item
NDRange是work-item的索引空间，被定义为数据类型为size_t，长度为N的数组，N代表所描述的work-item的维度

//work-item维度为2维，有1024个work-item
size_t SpaceSize[2]={1024,1};

一个work-group包含多个work-item，work-group维度与NDRange维度一致，NDRange元素长度要能被work-group长度整除

//一个work-group有64个work-item 总共分为16个work-group(1024/64)
size_t WorkGroupSize[2]={64,1};

每个work-item都是独立运行的，即使在同一个work-group内

平台和设备

主机-设备交互

平台模型将一个设备定义一系列计算单元（功能独立），每个计算单元中进一步划分为处理部件
clGetPlatformIDs()用来获取指定系统上的可用计算平台，在应用程序中应当调用两次，从而获取平台相关信息
clGetDeviceIDs()获取平台上支持的设备，如GPU、CPU和其他等要参考OpenCL 1.1 参考文档，与clGetPlatformIDs()类似调用方法

执行环境

上下文

上下文（context）存在于主机端，协调主机-设备的交互机制，管理设备上的可用内存对象，跟踪针对每个设备新建的kernel和程序
clCreateContext()用于新建上下文。其参数properties用于限定上下文范围，将上下文局限于某个平台可令开发者为多个平台提供上下文。开发者必须提供与context关联的设备数量和设备ID，也可设置用户回调函数

命令队列

一旦主机端指定运行kernel的设备且context已经新建，则每个设备必须新建命令队列才可运行，每个命令队列仅关联一个设备
主机端需要对设备进行操作，就需要把命令提交到设备对应的命令队列上
clCreateCommandQueue()用于新建命令队列并且关联到某个设备，支持乱序执行命令
所有指定主机与设备交互的api，会以’clEnqueue’开头，且需要一个命令队列作为参数

事件

任何操作作为命令入队到命令队列中，都会产生事件
作用：表示依赖；提供程序剖析机制（profiling）

内存对象

新建的内存对象，只在一个上下文中有效，运行时根据数据依赖关系的需要管理特定设备的数据传输

buffer

类似数组，由malloc()新建，buffer中数据在内存中连续存储
clCreateBuffer()分配缓冲区，返回内存对象；clEnqueueWriteBuffer()将主机端内存的数据传到OpenCL buffer中；clEnqueueReadBuffer()将OpenCL buffer数据传回主机端内存
如果在独立的加速器设备（如GPU）上执行的kernel依赖某个buffer，该buffer必须传输到设备上
buffer链接的是context，而不是设备，应当在OpenCL运行时决定数据移动

image

不透明对象，利于进行数据填充和其他可能的优化
未必所有OpenCL设备支持image，应当先使用clGetDeviceInfo()查看设备是否支持image对象
image不能像数组直接引用，相邻元素不保证放在连续内存中
image元素使用cl_image_format表示
通过clCreateImage2D()或clCreateImage3D()新建OpenCL image
读取image数据时，要用sampler对象；数据写入image时需强制转换为对应格式

flush命令和finish命令

clFinish()函数，阻塞直到命令队列中的所有命令完成
clFlush()函数，阻塞直到命令队列中的所有命令被移出队列，命令准备就绪但不保证命令已经执行完毕

新建一个OpenCL程序对象

OpenCL C代码被称为Program，是kernel函数的集合，kernel是被调度到设备上的执行单位
通过clCreateProgramWithSource()（较方便，直接读取源代码文件）或clCreateProgramWithBinary()使OpenCL 源代码转换为cl_program对象
通过clBuildProgram()编译对应Program，如果编译错误，会报告错误信息

OpenCL的kernel

从cl_program中抽取kernel，kernel的名称和program对象传递给clCreateKernel()，如果对象有效且kernel被发现，会返回kernel对象
kernel执行前，需先用clSetKernelArg()指定kernel的参数
通过clEnqueueNDRangeKernel()执行kernel，该函数的参数需根据参考文档认真研究

内存模型

全局内存，对设备上所有计算单位可见，主机->设备和设备->主机的数据驻留在全局内存上。__global关键字定义的数据，驻留在全局内存上
常量内存，方便所有work-item同时访问常量数据，全局内存的一部分。__constant关键字定义的数据，驻留在常量内存上
本地内存，供work-group共享。比全局内存访问延迟短，带宽更高。__local关键字定义的数据，驻留在本地内存上
私有内存，对单个work-item可见，局部变量和非指针类型的kernel参数默认为私有

编写kernel

函数以__kernel开始，返回值必须为void，必须指定指针所指向的地址空间（全局、常量和本地）
同一个work-group中的work-item多次使用的数据缓存到本地内存（__local）上可以提高性能。一旦work-item执行完成，其本地内存上数据是临时的，要保存结果则需把结果传回全局内存

本章向量相加程序源码

//kernel file code
__kernel void vecadd(__global int* A, __global int* B, __global int* C)
{int idx = get_global_id(0); //get the unique ID of the work-itemC[idx] = A[idx] + B[idx];
}

//main file code
#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>int main()
{int *A = NULL;  //inputint *B = NULL;  //inputint *C = NULL;  //outputconst int elements = 2048; //elements in each arraysize_t datasize = sizeof(int)*elements; //size of data//allocate space for arraysA = (int*)malloc(datasize);B = (int*)malloc(datasize);C = (int*)malloc(datasize);for(int i=0;i<elements;i++){ //initialize the inputA[i] = i;B[i] = i;}cl_int status; //for error checking but not checked in this example//step 1//discover and initialize platformscl_uint numPlatforms=0;cl_platform_id *platforms=NULL;status = clGetPlatformIDs(0,NULL,&numPlatforms); //get the number of platformsplarforms = (cl_platform_id*)malloc(numPlatforms*sizeof(cl_platform_id)); //allocate space for each platformstatus = clGetPlatformIDs(numPlatforms,platforms,NULL); //fill in platforms//step 2//discover and initialize devicescl_uint numDevices = 0;cl_device_id *devices = NULL;status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices); //get the number of all devicesdevices = (cl_device_id*)malloc(numDevices*sizeof(cl_device_id)); //allocate space for each devicestatus = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, numDevices, devices, NULL); //fill in devices//step 3//create contextcl_context context=NULL;context = clCreateContext(NULL, numDevices, devices, NULL, NULL, &status); //create a context and associate it with the devices//step 4//create a command queuecl_command_queue cmdQueue;cmdQueue = clCreateCommandQueue(context, devices[0], 0, &status); //create a command queue and associate it with the device you want to execute on
}//step 5//create device bufferscl_mem bufferA;cl_mem bufferB;cl_mem bufferC;bufferA = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status); //create a buffer object that will contain the data from the host array AbufferB = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, &status); bufferC = clCreateBuffer(context, CL_MEM_WRITE_ONLY, datasize, NULL, &status); //step 6//write host data to device buffersstatus = clEnqueueWriteBuffer(cmdQueue, bufferA, CL_FALSE, 0, datasize, A, 0, NULL, NULL);status = clEnqueueWriteBuffer(cmdQueue, bufferB, CL_FALSE, 0, datasize, B, 0, NULL, NULL);//step 7//create and compile the programcl_program program = clCreateProgramWithSource(context, 1, 'kernel file path', NULL, &status); //create the program from the kernel filestatus = clBuildProgram(program, numDevices, devices, NULL, NULL, NULL);//step 8//create the kernelcl_kernel kernel = NULL;kernel = clCreateKernel(program, "vecadd", &status); //create the kernel from the function named "vecadd"//step 9//set the kernel argumentsstatus = clSetKernelArg(kernel, 0, sizeof(cl_mem), &bufferA); //associate the buffers with the kernelstatus = clSetKernelArg(kernel, 1, sizeof(cl_mem), &bufferB);status = clSetKernelArg(kernel, 2, sizeof(cl_mem), &bufferC);//step 10//configure the work-item structuresize_t globalWorkSize[1];globalWorkSize[0] = elements; //number of work-item//step 11//enqueue the kernel executionstatus = clEnqueueNDRangeKernel(cmnQueue, kernel, 1, NULL, globalWorkSize, NULL, 0, NULL, NULL); //globalWorkSize is the 1D dimension of the work-item//step 12//read the output buffer back to the hostclEnqueueReadBuffer(cmdQueue, bufferC, CL_TRUE, 0, datasize, C, 0, NULL, NULL);bool result = true;for(int i=0;i<elements;i++){if(C[i] != i+i){result = false;break;}}if(result) printf("correct\n");else printf("incorrect\n");//step 13//release the OpenCL ResourcesclReleaseKernel(kernel);clReleaseProgram(program);clReleaseCommandQueue(cmdQueue);clReleaseMemObject(bufferA);clReleaseMemObject(bufferB);clReleaseMemObject(bufferC);clReleaseContext(context);//free host resourcesfree(A);free(B);free(C);free(platforms);free(devices);
}

OpenCL的并发与执行模型

OpenCL同步：kernel，fence和barrier

全局同步只定义在kernel执行边界处，即不同work-group的work-item无法进行同步
若同一个work-group内有一个work-item调用了barrier函数，则同一个work-group内的其他work-item到达barrier函数前，该kernel不会继续进行

队列与全局同步

同步点：
（1）调用clFinish函数，阻塞直至命令队列执行完毕
（2）等待特定事件完成
（3）阻塞访存操作，在入队函数中定义参数CL_TRUE

OpenCL内存一致性

如果kernel在第二个设备上运行，那么在第一个设备上产生的任何结果数据在第二个设备上都是随需可用的
第一个数据结构上时间的完成标示着数据可以移动，不需要进行单独的缓冲区拷贝操作

事件

入队函数的最后一个参数可以指定一个封装了入队命令状态的事件，所有事件执行完成之前，命令不会继续执行

多个设备上的队列

事件只能实现在同一个上下文的命令间同步，不同上下文（设备）见共享数据应调用clFinish函数，并对数据进行显式拷贝

多设备变成，多为两种执行模型
两个或多个设备以流水线的形式工作，一个设备等待另一个设备的运行结果
多个设备相互独立运行的任务模型

事件回调

事件回调可用于将新的命令入队，还可用于调用主机端函数
设置回调函数的clSetEventCallback函数必须在clEnqueueNDRangeKernel函数后调用

主机端内存模型

buffer对象

可理解为C语言中的数组，数据是连续分配，可随机存取
若没有定义采用阻塞方式存取，对buffer的存取是异步的

如图代码所示，不保证A和B的值是不一样的，但C是可以保证的。

image对象

对设备代码不可见，不可通过指针访问
多维结构
仅限于图像数据相关数据类型

《OpenCL异构计算》学习笔记相关推荐

第二行代码学习笔记——第六章:数据储存全方案——详解持久化技术
本章要点任何一个应用程序,总是不停的和数据打交道. 瞬时数据:指储存在内存当中,有可能因为程序关闭或其他原因导致内存被回收而丢失的数据. 数据持久化技术,为了解决关键性数据的丢失. 6.1 持久化技 ...
第一行代码学习笔记第二章——探究活动
知识点目录 2.1 活动是什么 2.2 活动的基本用法 2.2.1 手动创建活动 2.2.2 创建和加载布局 2.2.3 在AndroidManifest文件中注册 2.2.4 在活动中使用Toast ...
第一行代码学习笔记第八章——运用手机多媒体
知识点目录 8.1 将程序运行到手机上 8.2 使用通知 * 8.2.1 通知的基本使用 * 8.2.2 通知的进阶技巧 * 8.2.3 通知的高级功能 8.3 调用摄像头和相册 * 8.3.1 调用 ...
第一行代码学习笔记第六章——详解持久化技术
知识点目录 6.1 持久化技术简介 6.2 文件存储 * 6.2.1 将数据存储到文件中 * 6.2.2 从文件中读取数据 6.3 SharedPreferences存储 * 6.3.1 将数据存储到 ...
第一行代码学习笔记第三章——UI开发的点点滴滴
知识点目录 3.1 如何编写程序界面 3.2 常用控件的使用方法 * 3.2.1 TextView * 3.2.2 Button * 3.2.3 EditText * 3.2.4 ImageView ...
第一行代码学习笔记第十章——探究服务
知识点目录 10.1 服务是什么 10.2 Android多线程编程 * 10.2.1 线程的基本用法 * 10.2.2 在子线程中更新UI * 10.2.3 解析异步消息处理机制 * 10.2.4 ...
第一行代码学习笔记第七章——探究内容提供器
知识点目录 7.1 内容提供器简介 7.2 运行权限 * 7.2.1 Android权限机制详解 * 7.2.2 在程序运行时申请权限 7.3 访问其他程序中的数据 * 7.3.1 ContentRe ...
第一行代码学习笔记第五章——详解广播机制
知识点目录 5.1 广播机制 5.2 接收系统广播 * 5.2.1 动态注册监听网络变化 * 5.2.2 静态注册实现开机广播 5.3 发送自定义广播 * 5.3.1 发送标准广播 * 5.3.2 发 ...
第一行代码学习笔记第九章——使用网络技术
知识点目录 9.1 WebView的用法 9.2 使用HTTP协议访问网络 * 9.2.1 使用HttpURLConnection * 9.2.2 使用OkHttp 9.3 解析XML格式数据 * 9 ...
安卓教程----第一行代码学习笔记
安卓概述系统架构 Linux内核层,还包括各种底层驱动,如相机驱动.电源驱动等系统运行库层,包含一些c/c++的库,如浏览器内核webkit.SQLlite.3D绘图openGL.用于java运行 ...

《OpenCL异构计算》学习笔记

目录

引言

OpenCL简介

kernel和OpenCL执行模型

平台和设备

主机-设备交互

执行环境

上下文

命令队列

事件

内存对象

buffer

image

flush命令和finish命令

新建一个OpenCL程序对象

OpenCL的kernel

内存模型

编写kernel

本章向量相加程序源码

OpenCL的并发与执行模型

OpenCL同步：kernel，fence和barrier

队列与全局同步

OpenCL内存一致性

事件

多个设备上的队列

事件回调

主机端内存模型

buffer对象

image对象

《OpenCL异构计算》学习笔记相关推荐

最新文章

热门文章