GPU编程和流式多处理器

GPU编程和流式多处理器

流式多处理器(SM)是运行CUDA内核的GPU的一部分。本章重点介绍SM的指令集功能。

流式多处理器(SM)是运行我们的CUDA内核的GPU的一部分。每个SM包含以下内容。

  • 可以在执行线程之间划分的数千个寄存器
  • 几个缓存:
    • 共享内存,用于线程之间的快速数据交换
    • 恒定高速缓存,用于快速广播恒定内存中的读取
    • 纹理缓存,以聚合纹理内存的带宽
    • – L1缓存,可减少对本地或全局内存的延迟
  • Warp调度程序可以快速在线程之间切换上下文,并向准备执行的Warp发出指令
  • 用于整数和浮点运算的执行核心:
    • –整数和单精度浮点运算
    • –双精度浮点
    • –用于单精度浮点先验功能的特殊功能单元(SFU)

存在许多寄存器,以及硬件可以如此高效地在线程之间进行上下文切换的原因,可以最大程度地提高硬件的吞吐量。GPU被设计为具有足够的状态,以覆盖执行等待时间和数百个时钟周期的存储等待时间,执行读取指令后,设备存储器中的数据可能需要数百个时钟周期才能到达。

SM是通用处理器,设计与CPU中的执行内核有很大不同:目标时钟频率低得多;支持指令级并行性,但不支持分支预测或推测性执行;如果根本没有缓存,则缓存较少。对于合适的工作负载,GPU中强大的计算能力足以弥补这些缺点。

自2006年推出首款支持CUDA的硬件以来,SM的设计一直在迅速发展,其代号为Tesla,Fermi和Kepler的三个主要修订版本。开发人员可以通过调用cudaGetDeviceProperties()并检查cudaDeviceProp.majorcudaDeviceProp.minor或通过调用驱动程序API函数cuDeviceComputeCapability()来查询计算能力。计算能力1.x,2.x和3.x分别对应于Tesla类,Fermi类和Kepler类硬件。表8.1总结了每一代SM硬件中添加的功能。

表1 SM功能

 GPU编程和流式多处理器

 

 GPU编程和流式多处理器

 

 以后会显示不同SM的框图。CUDA内核可以执行整数和单精度浮点指令。一个双精度单元实现双精度支持(如果可用);和特殊功能单元实现倒数,倒数平方根,正弦/余弦和对数/指数函数。当执行指令所需的资源可用时,warp调度程序会将指令调度到这些执行单元。

本文重点介绍SM的指令集功能。因此,有时会引用“ SASS”指令,即ptxas或CUDA驱动程序将中间PTX代码转换成的本机指令。开发人员无法直接编写SASS代码;相反,NVIDIA通过cuobjdump实用程序使这些说明对开发人员可见,可以通过检查编译的微代码来指导其源代码的优化。

1. 存储Memory

1.1. 寄存器

每个SM包含数千个32位寄存器,这些内核在启动内核时指定给线程分配。寄存器是SM中最快,最丰富的内存。例如,开普勒类(SM 3.0)SMX包含65,536个寄存器或256K,而纹理缓存仅为48K。

CUDA寄存器可以包含整数或浮点数据。对于能够执行双精度算术(SM 1.3和更高版本)的硬件,操作数包含在偶数值寄存器对中。在SM 2.0和更高版本的硬件上,寄存器对也可以保存64位地址。

CUDA硬件还支持更广泛的内存事务:内置int2 / float2int4 / float4数据类型分别位于对齐的寄存器对或四元组中,可使用单个64位或128位宽的加载或存储来读取或写入。一旦进入寄存器,就可以将各个数据元素称为.x / .y(对于int2 / float2)或.x / .y / .z / .w(对于int4 / float4)。

通过指定命令行选项--ptxas-options --verbose开发人员可以使nvcc报告内核使用的寄存器数。内核使用的寄存器数量会影响SM中可容纳的线程数量,通常必须仔细调整,以获得最佳性能。可以使用--ptxas-options --maxregcount N指定用于编译的最大寄存器数。

注册采样Register Aliasing

寄存器可以保存浮点或整数数据,某些内在函数仅用于迫使编译器更改其变量视图。__int_as_float()__float_as_int()内部函数,导致32位整数和单精度浮点之间“改变性能”的变量。

float__int_as_float(int i);
int __float_as_int(float f);

__double2loint() __double2hiint() ,和__hiloint2double()内部函数类似原因寄存器来改变个性(通常就地)。__double_as_longlong()__longlong_as_double()就地强制寄存器对;__double2loint()__double2hiint()分别返回输入操作数的最低和最高32位;__hiloint2double()从高半部分和低半部分中构造一个精度型

int double2loint(double d);
int double2hiint(double d);
int hiloint2double(int hi,int lo);
double long_as_double(long long int i);
long long int __double_as_longlong(double d);

1.2. 本地Local Memory

本地存储器用于溢出寄存器,还用于存储已索引且无法在编译时计算其索引的局部变量。本地内存由与全局内存相同的设备内存池支持,因此它具有与Fermi和更高版本的硬件上的L1和L2缓存层次结构相同的延迟特性和优势。本地内存的寻址方式是自动合并内存事务。硬件包括加载和存储本地内存的特殊说明:SASS变体是Tesla的LLD / LST和Fermi和Kepler的LDL / STL

1.3. 全局Global Memory

SM可以使用GLD / GST指令(在Tesla上)和LD / ST指令(在Fermi和Kepler上)读取或写入全局内存。开发人员可以使用标准的C运算符来计算和取消引用地址,包括指针算法和取消引用运算符*[]->。对64位或128位内置数据类型(int2 / float2 / int4 / float4)进行操作,自动使编译器发出64位或128位加载和存储指令。通过合并内存事务可实现最大的内存性能。

特斯拉级硬件(SM 1.x)使用​​特殊的地址寄存器来保存指针。后来的硬件实现了一种加载/存储架构,该架构使用相同的寄存器文件来存储指针。整数和浮点值;以及用于恒定内存,共享内存和全局内存的相同地址空间。1个

费米级硬件包括旧硬件不具备的一些功能。

  • 通过“宽”加载/存储指令支持64位寻址,其中地址保存在偶数寄存器对中。在32位主机平台上不支持64位寻址。在64位主机平台上,将自动启用64位寻址。结果,针对为32位和64位主机平台编译的,相同内核生成的代码,可能具有不同的寄存器计数和性能。
  • L1缓存的大小可以配置为16K或48K。2(Kepler添加了将缓存拆分为32K L1 / 32K共享功能。)加载指令可以包括可缓存性提示(告诉硬件将读取的内容拖入L1或绕过L1,并将数据仅保留在L2中)。可通过嵌入式PTX或通过命令行选项–X ptxas –dlcm = ca(默认设置在L1和L2中缓存)或–X ptxas –dlcm = cg(仅在L2中缓存)访问这些文件。

即使多个GPU线程在同一内存位置上运行,原子操作(或仅仅是“原子”)也可以正常工作地更新内存位置。在操作期间,硬件会在内存位置强制执行互斥。由于不能保证操作顺序,因此通常支持的运算符是关联的。3

Atomics首先可用于SM 1.1和更高版本的全局内存以及SM 1.2和更高版本的共享内存。但是,在开普勒一代硬件出现之前,全局内存原子太慢而无法使用。

当通过--gpu-architecturenvcc指定了适当的体系结构时,表2中汇总的全局原子内在函数将自动变为可用。所有这些内在函数都可以对32位整数进行操作。SM 1.2中添加了对atomicAdd()atomicExch()atomicCAS()的64位支持。在SM 2.0中添加了32位浮点值(float)的atomicAdd()。在SM 3.5中添加了对atomicMin()atomicMax()atomicAnd()atomicOr()atomicXor()的64位支持。

表2原子操作

 GPU编程和流式多处理器

 

 注意

由于原子操作是使用GPU的集成内存控制器中的硬件实现的,无法在PCI Express总线上运行,无法在与主机内存或对等内存相对应的设备内存指针上正确运行。

在硬件级别,原子有两种形式:原子操作返回执行算子之前在指定内存位置的值,以及归约操作,开发人员可以在内存位置“触发并忘记”而忽略返回值值。由于如果不需要返回旧值,则硬件可以更有效地执行操作,编译器将检测是否使用了返回值,如果未使用,则发出不同的指令。例如,在SM 2.0中,指令分别称为ATOMRED

1.4. 恒定内存Memory

常量内存驻留在设备内存中,但由另一个只读缓存支持,该缓存经过优化,可以将读取请求的结果广播,到均引用同一内存位置的线程。每个SM均包含一个经过延迟优化的小型缓存,用于处理这些读取请求。将内存(和缓存)设置为只读可简化缓存管理,因为硬件无需实施回写策略来处理已更新的内存。

SM 2.x和后续硬件包括针对内存的特殊优化,该优化未表示为常量,但编译器已将其标识为(1)只读和(2)其地址不依赖于块或线程ID。“均匀加载”(LDU)指令使用恒定的缓存层次结构读取内存,并将数据广播到线程。

1.5. 共享内存

共享内存非常快,是SM中的片上内存,线程可以将其用于线程块内的数据交换。由于它是每个SM资源,因此共享内存的使用会影响占用率,即SM可以保留的warp数量。SM使用特殊指令加载和存储共享内存:SM 1.x上为G2R / R2G,SM 2.x及更高版本上为LDS / STS

共享内存安排为交错的存储体,通常针对32位访问进行了优化。如果经线中有多个线程引用同一存储体,则会发生存储体冲突,并且硬件必须连续处理内存请求,直到为所有请求提供服务为止。为了避免存储区冲突,应用程序将基于线程ID以交错模式访问共享内存,如下所示。

extern __shared__ float shared[];

float data = shared[BaseIndex + threadIdx.x];

从相同的32位共享内存位置读取扭曲中的所有线程也很快。硬件包括广播机制以针对这种情况进行优化。写入同一存储区的操作会由硬件进行序列化,降低性能。写入同一地址会导致争用情况,应避免。

对于2D访问模式(例如图像处理内核中的像素图块),最好填充共享内存分配,以便内核可以引用相邻行,而不会引起存储体冲突。SM 2.x和后续硬件具有32个存储区,其中每个存储区可供2个图块使用,其中同一warp中的线程可以按行访问数据,这是将图块大小填充为33个32位字的倍数的好策略。

在SM 1.x硬件上,共享内存的大小约为16K。在更高版本的硬件上,总共有64K的L1高速缓存,可以配置为16K或48K共享内存,其余部分用作L1高速缓存。

在过去的几代硬件中,NVIDIA改进了硬件对除32位以外的操作数大小的处理。在SM 1.x硬件上,来自同一存储区的8位和16位读取导致存储区冲突,而SM 2.x和更高版本的硬件可以从同一存储区广播中,任何大小的读取。同样,共享内存中的64位操作数(例如double)比SM 1.x上的32位操作数慢得多,以至于开发人员有时不得不求助于将数据分别存储为上下半部分。SM 3.x硬件为主要在共享内存中使用64位操作数的内核添加了一项新功能:将存储体大小增加到64位的模式。

共享内存中的原子算子

SM 1.2添加了在共享内存中执行原子操作的功能。全局内存与使用单个指令(取决于GATOMGRED,取决于是否使用返回值)来实现原子的全局存储器不同,共享内存原子是通过显式的锁定/解锁语义实现的,编译器发出的代码会导致每个线程循环这些锁定操作,直到线程执行其原子操作为止。

清单1将源代码提供给atomic32Shared.cu,该程序专门用于编译以突出显示共享内存原子的代码生成。清单2显示了为SM 2.0生成的最终微代码。注意,LDSLK(与锁共享负载)指令如何返回谓词,该谓词说明是否获取了锁,确定了要执行更新的代码,并且代码循环运行,直到获取了锁并执行了更新。

锁是按32位字执行的,锁的索引由共享内存地址的位2–9确定。注意避免争用,否则清单2中的循环最多可以迭代32次。

list1. atomic32Shared.cu

__global__ void

Return32( int *sum, int *out, const int *pIn )

{

    extern __shared__ int s[];

    s[threadIdx.x] = pIn[threadIdx.x];

    __syncthreads();

    (void) atomicAdd( &s[threadIdx.x], *pIn );

    __syncthreads();

    out[threadIdx.x] = s[threadIdx.x];

}

Listing 2. atomic32Shared.cubin (microcode compiled for SM 2.0)

code for sm_20

    Function : _Z8Return32PiS_PKi

/*0000*/     MOV R1, c [0x1] [0x100];

/*0008*/     S2R R0, SR_Tid_X;

/*0010*/     SHL R3, R0, 0x2;

/*0018*/     MOV R0, c [0x0] [0x28];

/*0020*/     IADD R2, R3, c [0x0] [0x28];

/*0028*/     IMAD.U32.U32 RZ, R0, R1, RZ;

/*0030*/     LD R2, [R2];

/*0038*/     STS [R3], R2;

/*0040*/     SSY 0x80;

/*0048*/     BAR.RED.POPC RZ, RZ;

/*0050*/     LD R0, [R0];

/*0058*/     LDSLK P0, R2, [R3];

/*0060*/     @P0 IADD R2, R2, R0;

/*0068*/     @P0 STSUL [R3], R2;

/*0070*/     @!P0 BRA 0x58;

/*0078*/     NOP.S CC.T;

/*0080*/     BAR.RED.POPC RZ, RZ;

/*0088*/     LDS R0, [R3];

/*0090*/     IADD R2, R3, c [0x0] [0x24];

/*0098*/     ST [R2], R0;

/*00a0*/     EXIT;

    ...................................

1.6. 障碍和连贯性

熟悉的__syncthreads()内部函数会等待,直到继续执行线程块中的所有线程为止。需要保持线程块内共享内存的一致性。其他类似的内存屏障指令,也可以用于在更大范围的内存上执行某些排序,如表3所述。

表3内存屏障本质

 GPU编程和流式多处理器

 

上一篇:一个Markdown文件处理程序


下一篇:Java中的System类总结