要点初见：OpenCL 2.0 异构计算 [第三版] 知识点整理
时间：2022-11-22 12:30:00
Word 版Github项目地址:https://github.com/BingLiHanShuang/Chinese-Knowledge-Collation-of-Heterogeneous-Computing-with-OpenCL-2.0
OpenCL 2.0 异构计算 [第三版] （Heterogeneous Computing with OpenCL 2.0）
https://www.bookstack.cn/read/Heterogeneous-Computing-with-OpenCL-2.0/README.md
Intel opencl sdk下载安装
https://software.intel.com/content/www/us/en/develop/tools/opencl-sdk/choose-download.html
目录
OpenCL 2.0 异构计算 [第三版] （Heterogeneous Computing with OpenCL 2.0）
Intel opencl sdk下载安装
OpenCL较大的改变
第1章 异构计算简介
异构计算：串行处理、并行处理
加速的方法
CMOS
并发与并行
线程
共享内存模型
瓶颈
信息通信机制
并行计算的粒度
细粒度并行
粗粒度并行
选择粗细粒度
将数据视为向量
用于数据共享
OpenCL2.共享虚拟内存的三
OpenCL2.0新特性
OpenCL多层次并行化
OpenCL 设备端-主机端语言
OpenCL并行基于任务和数据
OpenCL2.共享虚拟内存
OpenCL2.0内存一致模型
第2章 设备架构
执行超标量的方法
向量操作
三种常见的并行形式
实现硬件多线程有两种方法
片上系统的优点
局部内存访问方式
第3章 介绍OpenCL
OpenCL标准
clGetPlatformIDs()
clGetDeviceIDs()
clCreateContext()
clCreateContextFromType()
clGetContextInfo()
命令指定的行为
clCreateCommandQueueWithProperties()
clEnqueue
clFlush()
clFinish()
事件
命令状态
CL_COMPLETE
clGetEventInfo()
APIclWaitForEvents()
设备端入队
向量加法
条带处理
工作项
get_global_id(0)
每个维度上工作项的数量(NDRange)
工作组
使用源码创建内核的步骤
创建内核对象
clBuildProgram()创建NVIDIA GPU的中间码
clGetProgramInfo()
clCreateProgramWithBinary()
clSetKernelArg()
clEnqueueNDRangeKernel()
OpenCL三种内存类型
clCreateBuffer()
clCreateImage()
通道序
OpenCL读写图像数据的内置函数
clCreatePipe()
clGetPipeInfo()
内核对管道的读写
OpenCL访问管道
clEnqueueWriteBuffer()和clEnqueueReadBuffer()
OpenCL内存
OpenCL设备内存分类
通用地址空间
创建并执行OpenCL应用的步骤
使用C API实现的OpenCL向量相加 完整代码
使用C++ Wapper实现的OpenCL向量相加
第4章 OpenCL案例
OpenCL计算直方图
思路
计算直方图的OpenCL内核代码
内存栅栏
clEnqueueFillBuffer()
直方图统计的主机端代码
OpenCL旋转图像
图像旋转内核
图像采样器(sampler_t sampler)
CLK_NORMALIZED_COORDS_FALSE
CL_ADDRESS_CLAMP
CLK_FILTER_NEAREST和CLK_FILTER_LINEAR
为旋转例程创建图像对象
图像旋转主机端的完整代码
主机端创建采样器
图像卷积主机端完整代码
生产者-消费者
管道
卷积内核(生产者)
卷积内核(消费者)
创建管道对象
生产者-消费者主机端完整代码
打印编译错误信息
clGetProgramBuildInfo()
clProgramBuildInfo()
查询程序对象编译日志的函数封装
创建一个程序字符串
第5章 OpenCL运行时和并发模型
OpenCL命令队列的主要同步点
clEnqueueReadBuffer()中的blocking_read
事件命令状态
clGetEventInfo()
clWaitForEvents()
clEnqueueBarrierWithWaitList()
clEnqueueMarkerWithWaitList()
clSetEventCallback()
开启命令队列的计时功能
clGetEventProfilingInof()
clCreateUserEvent()
clSetUeserEventStatus()
设置乱序的标志位
OpenCL同步事件对象
OpenCL下的多设备编程(使用异构设备进行并行编程)
clEnqueueNDRangeKernel()
分歧
clGetKernelWorkGroupInfo()
获取工作项在执行区域内的具体位置
get_local_size()
get_enqueue_local_size()
size_t get_global_linear_id()
size_t get_local_linear_id()
work_group_barrier()
flags可以设置成的实参
三类工作组评估函数
谓词评估函数
广播函数
并行原语函数
闭扫描和开扫描
原生内核
内置内核
嵌套并行机制(nested parallelism)
设备端入队的好处
enqueue_kernel()
Clang块
设备端设置内核参数需要动态分配局部内存时
第6章 OpenCL主机端内存模型
图像内存对象的优势
clCreateBuffer()
子数组对象(subbuffer)
OpenCL中图像与数组对象的不同
clEnqueueReadImage()
clGetSupportedImageFormats()
clCreateImage()
clCreatePipe()
在数组对象创建和初始化、传递内核参数，从内存对象中读回数据中，数据转移的过程
显式的从主机端或设备端将数据拷贝到设备端或主机端
clEnqueueMigrateMemObjects()
CL_MIGRATE_MEM_OBJECT_HOST
零拷贝(zero-copy)数据
clEnqueueMapBuffer()
AMD处理内存标志
OpenCL 2.0三种SVM类型
OpenCL共享数据的行为对比
粗粒度数组SVM
粗粒度SVM使用
clSVMAlloc()
clSVMFree()
细粒度数组SVM
第7章 OpenCL设备端内存模型
OpenCL内存空间四种类型
OpenCL C保证内存的一致性的同步操作
层级的一致性描述
work_group_barrier()
原子操作的方式
OpenCL 2.0原子操作
atomic_fetch()
全局数组
OpenCL图像类型
读取图像的内置函数
采样器对象
写入图像函数
管道读写函数
reserve_read_pipe()和reserve_write_pipe()
commit_read_pipe()和commit_write_pipe()
work_group_reserve_write_pipe()和work_group_reserve_read_pipe()
work_group_commit_read_pipe()和work_group_commit_write_pipe()
使用常量地址空间的方式
OpenCL查询常量参数个数的限制、常量数组的最大尺寸
clSetKernelArg()
统一地址空间
OpenCL的一致性顺序
可以作为内存域指定的选项
atomic_compare_exchange_strong_explicit()
OpenCL C对原子操作的操作域进行初始化
atomic_work_item_fence()
第8章 异构系统下解析OpenCL
GPU相比于CPU的特性
使用GPU处理图像获取较高性能的需求
合并访问
计算内核的带宽大小公式
查询对应设备上局部内存的大小
第9章 案例分析：图像聚类
图像聚类步骤
CPU串行实现SURF
OpenCL原子操作实现SURF
OpenCL合并访问内存（转置矩阵）实现SURF
OpenCL向量化实现SURF
局部内存的优势
OpenCL将SURF特征放入局部内存实现SURF
OpenCL将聚类中点坐标放入常量内存实现SURF
第10章 OpenCL的分析和调试
性能优化要点
clGetEventProfilingInfo()
使用OpenCL事件获取内核的时间信息
CodeXL的主要功能
CodeXL的使用方式
OpenCL应用的跟踪数据可了解的信息
时间线对于调试OpenCL的好处
总结页面的主要汇总信息
性能计数器的作用
内核分析器的收益
ISA码的益处
不让编译器对OpenCL内核进行任何的优化
内核分析器推断方式
异构应用的代码构成
API级别调试的特性
CodeXL有对OpenCL应用进行调试的方式
使用printf调试
第11章 高级语言映射到OpenCL2.0 —— 从编译器作者的角度
C++ AMP为C++编程语言添加的规则
C++ AMP向量相加
C++ AMP特性
C++ AMP array_view
C++ AMP parallel_for_each
C++ AMP restrict
CLamp开源的组件
clCreateProgramWithBinary()
OpenCL的相关步骤在C++ AMP的对应情况
CLamp编译器编译和链接C++ AMP程序
线程划分的显式隐式
支持线程划分的显式编程模型的特点
C++ AMP和OpenCL地址空间的区别
OpenCL创建只读只写内存对象
discard_data()
array_view
第12章 WebCL：使用OpenCL加速Web应用
WebCL编程的构成
WebCL定义的对象
WebCL处理对象
WebCL注意点
本设备的OpenCL是否支持乱序
WebCL中的enqueueNDRange
WebCL中webCLCommandQueue类的参数
WebGL、WebCL能够共享的内存对象
运行时调用WebCL和WebGL
两个能保护内存的准则
保证程序安全性的机制
第13章 其他高级语言中OpenCL的使用
OpenCL较大的改变
1、共享虚拟内存
主机和设备端可以共享复杂数据结构指针，比如：树和链表；以减少花在数据结构转换上的精力。
2、动态并行
可以在不用主机代码参与的情况下，进行内核的加载，为的就是减小加载内核所产生的瓶颈。
3、统一地址空间
这样同样的函数就可以处理GPU和CPU上的数据，以减少编程的难度。
4、C++原子操作
工作项可以跨越工作组共享数据和设备，可以让更多的算法使用OpenCL实现。
第1章 异构计算简介
异构计算：串行处理、并行处理
加速的方法
1、分而治之
将一个问题递归的划分为数个子问题，直到可用的计算资源能够解决划分后的子问题。
2、分散-聚合
先发送一部分输入数据到每个并行资源中，然后将这些输入数据处理后的结果进行收集，再将这些结果合并到一个结果数据集中。
CMOS
互补金属氧化物半导体，Complementary Metal Oxide Semiconductor
并发与并行
并行编程必须是并发的，不过并发编程不一定并行。虽然很多并发程序可以并行执行，但是互相有依赖的并发任务就不能并行了。比如：交错执行就符合并发的定义，而不能并行的执行。所以，并行是并发的一个子集，并发程序是所有程序集的一个子集。
线程
进程中的所有线程都会共享一些资源(比如：内存、打开的文件、全局变量)，不过他们也有属于自己的资源(比如：堆栈、自动变量)。
线程使用全局共享地址空间分配出的变量进行通讯。
通讯时需要有同步机制来保证同一个内存区域的内容不会被多个线程更新。
共享内存模型
关键的特性：不需要编程者去管理数据的移动
在这样的系统中，线程如何去更新全局变量，底层硬件和编程者要达成共识，并遵守“内存一致性模型”
瓶颈
内存一致性模型下，能够使用的处理器相对较少，这是因为共享总线和相干性协议将会成为性能瓶颈。
系统中多一些松散的共享内存，相对来说会好一些；当有成规模的内核在共享内存系统中时，其会让系统变得复杂，并且内核间的互相通讯也要付出很高的代价。
大多数多核CPU平台都支持某一种共享内存，OpenCL也可以在支持共享内存的设备上运行。
消息通讯机制
消息传递接口(MPI)库在当今环境下依旧是很受欢迎的消息传递中间件
并行计算的粒度
线程中计算量与通讯量(比如，需要在线程间进行同步的变量)
细粒度并行
使用细粒度并行时，需要考虑如下几点：
1、计算量不要过大，这样每个线程都有足够的工作可做。
2、尽可能减少在数据同步上的开销，这样每个线程能够独立的完成自己的任务。
3、负载的划分也很重要，因为有大量的独立任务需要并行执行，设计良好的任务调度器都能灵活控制负载，并保证在多任务运行的同时，让线程上的达到均衡。
粗粒度并行
使用粗粒度并行时，需要考虑如下几点：
1、计算量肯定要高于细粒度并行时的计算量，因为不会像细粒度并行那样，有很多线程同时执行。
2、编程者使用粗粒度编程时，需要了解应用的整体结构，让粗粒度中的每个线程作为任务，服务于应用。
粗细粒度选择
当同步和通讯的开销大于计算，那么将并行粒度加大，更有利于控制同步和通讯所需的过高开销。
细粒度并行能减少负载不均和访存延迟在性能上的开销(对于GPU来说更是如此，其为粒度之细可以切换线程时达到零开销，同时也能隐藏访存延迟)。
将数据视为向量
当一个算法需要对大量数据进行同一组操作时，就可将数据视为向量，执行同一操作时，多个数据作为输入，经过向量操作后输出多个数据。
这种执行方式就是利用单指令多数据(SIMD)的方式对数据进行处理，可并行硬件可以利用这种执行方式，并行的对不同的数据进行处理。这种粒度的并行与向量的大小有关，通过SIMD执行单元的多数据处理来获得应用性能的提升。
数据共享的用途
1、一个任务中的输入，依赖于另外一个任务的输出——例如：生产者-消费者模型、流水线执行模型。
2、需要将中间结果进行综合(比如：归约，还有图1.5里的例子)。
理想状态下，可以尝试将应用能并行的部分剥离出来，确保并行的部分没有数据依赖，不过这只是理想状态而已。栅栏和锁有时会用来做数据的同步。
OpenCL2.0的三种共享虚拟内存
1、粗粒度共享缓存
2、细粒度共享缓存
3、系统级细粒度共享缓存
使用共享虚拟内存需要OpenCL实现将系统中主机和设备端的地址链接起来。这就允许在OpenCL内核中使用数据结构的指针(比如：链表)，之前版本的OpenCL标准是不支持内核中使用自定义数据结构指针。
粗粒度缓存支持通过API的调用，来更新整个缓存的内容。
细粒度缓存和系统级细粒度的粒度为字节级，这就意味着无论是主机端，还是设备端，对数据的更新将会立即同步。
细粒度共享内存是按内存一致模型中定义好的内存序，在同步点对数据进行同步。
OpenCL2.0新特性
1、嵌套并行化
2、共享虚拟内存
3、管道内存对象
4、C11原子操作
5、增强图像支持
OpenCL多级别的并行化
能够将应用高效的映射到同构或异构、单独或多个CPU或GPU上，以及其他硬件供应商提供的硬件系统当中。
OpenCL “设备端-主机端”语言
主机作为对其他设备的管理者。
设备端语言被设计用来映射内存系统和执行模型。
主机端语言用较低的开销，为复杂的并发程序搭建任务管道。
OpenCL基于任务和基于数据的并行
OpenCL的内核类似于SPMD模型，内核就是并行单元(OpenCL称为工作项)中执行的实例，内核实例创建并入队后，会被映射到支持标量和矢量的硬件上运行。
OpenCL2.0共享虚拟内存
共享虚拟内存是一项很重要的特性，其对能减轻编程者的负担，特别是在类似APU这样使用统一物理内存的系统上。
OpenCL2.0内存一致模型
提供“获取/释放”语义以缓解编程者在不明确的锁上耗费不必要的精力
第2章 设备架构
超标量执行方式
VLIW和硬件管理：通过查询地址，并行执行同指令流中不相关的指令
SIMD与向量并行：让指令在数据上并行执行
向量操作
将向量化操作通用化，并且向量操作通常会用来处理较长连续的数据序列，通常会使用流水线的形式进行向量操作，而非同时对多个数据进行操作，并且向量操作对连续、密集的内存读写给予了更多的支持。
并行的三种常见的形式
指令并行，数据并行、线程并行。
并行的方式就是执行多个独立的指令流
两种方式实现硬件多线程
1、并发多线程(SMT)
多线程的指令在执行资源上交替执行(通过超标量扩展的调度逻辑和线程资源)
缺点：需要对更多的状态信息进行保存，并且会让指令间得依赖关系和调度逻辑变得更加复杂
2、时域多线程
每个线程都能通过轮询的方式连续执行
目的：两个线程共享一个ALU。
优势：
1、调度逻辑简单。
2、流水线的延迟可以隐藏对多个线程的切换(调度)，减少转发逻辑。
3、当有线程缓存未命中，或等待另一个分支计算的结果等之类事件，都能通过改变线程指令顺序进行掩盖，并且执行更多的线程能更好的掩盖流水线上的延迟。
原理：以延迟换取最大的吞吐量，通过时域多线程的方式，来你做吞吐量的计算：多个线程交替执行，以保证设备处于忙碌状态，不过每个独立线程的执行时间要多于其理论最小执行时间。
片上系统优点
1、将很多元素融合在一个设备上，这样在生产的时候一次就能搞定，避免在制造阶段耗费太多成本。
2、更少的功能将会降低设备上的面积占用率，可以节省功耗和降低设备体积，这对于移动通讯领域很重要。
3、更短的距离意味着数据交互和信息传递完全可以在一个内存系统中进行，并且降低了交互的功耗。
4、低通讯延迟可以增加处理器对负载分发的次数。
内存访问的局部方式
1、空域：两个以上(包括两个)的读或写操作，对内存上(一定程度上)附近的地址。
2、时域：在相对小的时间窗内对同一地址进行两个以上(包括两个)的读或写操作。
第3章 介绍OpenCL
OpenCL标准
1、平台模型
指定一个host处理器，用于任务的调度。以及一个或多个device处理器，用于执行OpenCL任务(OpenCL C Kernel)。这里将硬件抽象成了对应的设备(host或device)。
2、执行模型
定义了OpenCL在host上运行的环境应该如何配置，以及host如何指定设备执行某项工作。这里就包括host运行的环境，host-device交互的机制，以及配置内核时使用到的并发模型。并发模型定义了如何将算法分解成OpenCL工作项和工作组。
3、内核编程模型
定义了并发模型如何映射到实际物理硬件。
4、内存模型
定义了内存对象的类型，并且抽象了内存层次，这样内核就不用了解其使用内存的实际架构。其也包括内存排序的要求，并且选择性支持host和device的共享虚拟内存。
clGetPlatformIDs()
查找制定系统上的可用OpenCL平台的集合。
在具体的OpenCL程序中，这个API一般会调用两次，用来查询和获取到对应的平台信息：
第一次调用这个API需要传入num_platforms作为数量参数，传入NULL作为平台参数。这样就能获取在该系统上有多少个平台可供使用。编程者可以开辟对应大小的空间(指针命名为platforms)，来存放对应的平台对象(类型为 cl_platform_id)。
第二次调用该API时，就可将platforms传入来获取对应数量的平台对象。平台查找完成后，使用clGetPlatformInfo()API可以查询对应供应商所提供的平台，然后决定使用哪个平台进行运行OpenCL程序。
clGetPlatformIDs()这个API需要在其他API之前调用。
格式：
cl_int
clGetPlatformIDs(
cl_uint num_entries,
cl_platform_id *platforms,
cl_uint *num_platforms)
clGetDeviceIDs()
查询平台上可用的设备，多了平台对象和设备类型作为入参。
需要三步创建device：
第一，查询设备的数量；
第二，分配对应数量的空间来存放设备对象；
第三，选择期望使用的设备(确定设备对象)。
device_type参数可以将设备限定为GPU(CL_DEVICE_TYPE_GPU)，限定为CPU(CL_DEVICE_TYPE_CPU)，或所有设备(CL_DEVICE_TYPE_ALL)，当然还有其他选项。这些参数都必须传递给clGetDeviceIDs()。相较于平台的查询API，clGetDeviceInfo()API可用来查询每个设备的名称、类型和供应商。
格式：
cl_int
clGetDeviceIDs(
cl_platform_id platform,
cl_device_type device_type,
cl_uint num_entries,
cl_device_id *devices,
cl_uint *num_devices)
clCreateContext()
创建上下文对象。
properties参数用于限制上下文作用的范围。这个参数可由特定的平台提供，其能够使能与图像的互用性，或使能其他能力。
通过限制给定平台的上下文，允许编程者使用多个平台创建的不同的上下文，并且能在一个平台中混用多个供应商提供的设备。
另外，创建上下文时必须要使用设备对象，并且编程者可以设置一个用户回调函数，还可以额外传递一个错误码(需要在错误码对象的生命周期内)用于获取API运行的状态。
格式：
cl_context
clCreateContext(
const cl_context_properties *properties,
cl_uint num_devices,
const cl_device_id *devices,
void (CL_CALL_BACK *pfn_notify)(
const char *errinfo,
const void *private_info,
size_t cb,
void *user_data),
void *user_data,
cl_int *errcode_ret)
clCreateContextFromType()
可以使用所有的设备类型(CPU、GPU和ALL)创建上下文。
clGetContextInfo()
创建上下文之后，可以用来查询上下文中设备的数量，以及具体的设备对象。
命令指定的行为
执行内核、进行数据传递、执行同步
clCreateCommandQueueWithProperties()
创建命令队列，且将命令队列与一个device进行关联。
peoperties参数是由一个位域值组成，其可使能命令性能分析功能(CL_QUEUE_PROFILING_ENABLE)，以及/或允许命令乱序执行(CL_QUEUE_OUT_OF_DRDER_EXEC_MODE_ENABLE)。
格式：
cl_command_queue
clCreateCommandQueueWithProperties(
cl_context context,
cl_device_id device,
cl_command_queue_properties peoperties,
cl_int *errcode_ret)
clEnqueue
任何以clEnqueue开头的OpenCL API都能向命令队列提交一个命令，并且这些API都需要一个命令队列对象作为输入参数。例如，clEnqueueReadBuffer()将device上的数据传递到host，clEnqueueNDRangeKernel()申请一个内核在对应device执行。
所有的clEnqueue开头的API，均有三个共同的参数：
1、事件链表的指针，其指定了当前命令依赖的事件列表
2、等待列表的长度
3、表示当前命令执行的事件指针，这个指针用于依赖该命令的其他命令。
clFlush()
对命令队列进行栅栏操作，需要一个命令队列作为参数。
clFlush()将会阻塞host上的执行线程，直到命令队列上的命令都从队列上移出。移出命令队列后的命令，就已经提交到device端，不过不一定完全执行完成。
格式：
cl_int clFlush(cl_command_queue command_queue);
clFinish()
对命令队列进行栅栏操作，需要一个命令队列作为参数。
clFinish()的调用将会阻塞host上的执行线程，直到命令队列上的所有命令执行完毕，其功能就是和同步栅栏操作一样。
格式：
cl_int clFinish(cl_command_queue command_queue);
事件
OpenCL API中，用来指定命令之间依赖关系的对象称为事件(event)
命令状态
1、Queued：命令处于命令队列中。
2、Submitted：命令从命令队列中移除，已经提交到设备端执行。
3、Ready：命令已经准备好在设备上执行。
4、Running：命令正在设备上执行。
5、Ended：命令已经在设备上执行完成。
6、Complete：所有命令以及其子命令都执行完成。
CL_COMPLETE
当命令成功的执行完成，事件的状态将会被设置为CL_COMPLETE。
如果命令非正常终止，事件的状态将会为一个负数值。这种情况下，有非正常终止的命令队列，以及其他在同一上下文上创建的命令队列，都将不能正常使用或运行。
clGetEventInfo()
查询事件
APIclWaitForEvents()
用于和host同步，该API阻塞host的执行线程，等待指定事件队列上的所有命令执行完毕。
格式：
cl_int
clWaitForEvents(
cl_uint num_events,
const cl_event *event_list)
设备端入队
执行中的内核现在可以让另外一个内核进入命令队列中。这种情况下，正在执行的内核可以称为“父内核”，刚入队的内核称为“子内核”。虽然，父子内核是以异步的方式执行，但是父内核需要在子内核全部结束后才能结束。我们可通过与父内核关联的事件对象来对执行状态进行查询，当事件对象的状态为CL_COMPLETE时，就代表父内核结束执行。
设备端的命令队列是无序命令队列，其具有无序命令
