OpenCL并行编程基础 第五章

GPU架构体系请参见相关文档。此处不赘述

 

5.1 并行线程组织

OpenCL并行程序的基本并行执行单元式线程（即 计算机上能够独立调度的一段顺序执行的指令流，GPU上使用硬件调度）

GPU的一个内核总是装备SIMD硬件，包括多个计算和内存处理流水线，支持多个线程同时执行，这些线程即波前（wavefront 又为warp（线程束——NVIDIA））,宽度（线程个数）为SIMD的宽度。NVIDIA多为32，AMD多为64。

准备多个波前线程，以便隐藏内存延时。这一批线程为工作组（work-group），工作组内的线程被分配到同一个内核中，它们可通过局部存储器交换数据。

工作组大小有限，全部线程个数由总工作量 和 每个线程的工作量 共同决定，所以一个工作组不能喂饱GPU。引入网格（grid），又名 N维范围（N-Dimensional Range i.e. ND-Range）。一个OpenCL内核函数 对应一个网格。

（图片来自《异构处理器OpenCL编程导论》）

clEnqueueNDRangeKernel(任务队列，内核，并行组织，offset，网格尺寸，维度尺寸，)

const size_t global[] = {4, 4};     // 网格中有4*4个工作组
const size_t local[] = {16, 16};    // 每个工作组中有16*16个线程
const size_t work_dim = 2;          // 工作维度
clEnqueueNDRangeKernel(queue,kernel,work_dim,offset,&global,&local,0,NULL,NULL);

1，2，3维组织方式再编译后都键索引值转为一维，再由GPU硬件调度。

5.2 OpenCL存储器模型

3个物理空间和4个逻辑空间。物理空间：GPU有 寄存器 ，片上高速存储器 ， 片外存储器 （有常数存储器 & 全局存储器）。

1. 私有存储器：每个GPU内核的寄存器堆，对程序员透明。OpenCL中定义的变量，默认为私有存储器空间的变量，只能被相应线程使用。

   int i;
   float x[4];

    

2. 局部存储器（共享存储器）：每个内核配备一定的片上存储器，用__local 标识符。为一个工作组（work-group）共享。不同工作组间不可见，访问速度与寄存器相当。常48KB or 64KB /内核。可组织多个内存块（bank）来提高访问速度（提高带宽）。

   __local int i;
   __local float x[4];

    

3. 常数存储器：GPU片外存储器的一部分，初始化后不可更改，用__constant修饰，能自动获得GPU内部专用的常数缓存支持，速度较快，所有线程可访问。

   constant int i;
   constant float x[4];

    

4. 全局存储器：GPU片外存储器，由CPU定义，GPU用。有三种常见类型Buffer（数组或自定义结构体），Image（图像格式存储的一维二维三维对象，专门的函数处理），Pipe（一个数据队列，访问太慢）

   __global float4 *color;    //4浮点数矢量
   struct foo_t{
       float a[3];
       int b[2];
   };
   __global foo_t *my_info;    //foo_t 类型数组

    

5.3 数据类型

标量数据，矢量数据，抽象数据，保留数据

内容太多，全是查表用的，直接搬书

 

 

5.4 运算符

OpenCL提供了矢量运算操作，矢量变量可以直接使用算术和逻辑运算符

float4 v,u;
float4 y;

y = u + v;
//四元矢量相加

//对应的标量现价代码如下：
y.x = v.x + u.x;
y.y = v.y + u.y;
y.z = v.z + u.z;
y.w = y.w + u.w;

 

float4 v,u;
float f;

v = u + v;
//四元矢量相加

//对应的标量现价代码如下：
v.x = u.x + f;
v.y = u.y + f;
v.z = u.z + f;
v.w = u.w + f;

 

5.5 函数

__attribute__表示可选属性说明

__kernel
__attribute__ ((vec_type_hint(<typen>)))        //说明编译器对SIMD宽度矢量化参数，默认GPU整数位宽，NVDIA为32，INTEL为8
__attribute__ ((work_group_size_hint(X,Y,Z)))   //内核函数一般采用的线程组织方式
__attribute__ ((reqd_work_group_size(X,Y,Z)))   //内核函数必须使用的线程组织方式
void foo(...){
    ...
}

OpenCL有一组获取并行维度核线程索引的函数，get_domain_size(uint dim)和get_domain_id(uint dim)，domain表示访问域，可以是global（对应全部线程）或local（对应整个工作组），dim表示维度。e.g. 网格以三位方式组织，dim可以等于0、1、2，分别对应3个维度的值。

GPU硬件上提供了特殊函数计算单元，但精度会降低，可牺牲速度运算软件实现来保证精度。

附图5-9