CUDA(七) 周斌 CUDA 编程

内置类型和函数Built-ins and functions

线程同步问题Synchronizing threads

线程调度问题Scheduling threads

存储模型Memory model

重访Matrix multiply

原子函数Atomic functions

CUDA函数声明

	执行位置 Executed on the	调用位置 Only caliable from the
_device_ float DeviceFunc()设备端函数	device	device
_global_ void KernelFunc()入口函数	device  GPU	host
_host_ float HostFunc()主机端函数	host主机	host

 

_global_返回类型必须是void

_device_以前是默认内联，现在有了变化

Global和device函数

•   尽量少用递归
•   不要用静态变量
•   少用malloc（现在允许但是不鼓励）
•   小心通过指针实现的函数调用

向量数据类型

char[1-4], uchar[1-4]符号

short[1-4], ushort[1-4]短

int[1-4], unit[1-4]

long[1-4], ulong[1-4]长

longlong[1-4], ulonglong[1-4]长长

float[1-4]单浮点型

double1 ,double2双精度浮点型

同时适用于host 和 device代码

通过函数 make_<type name> 构造

int2 i2 = make_int2(1, 2);
float4 f4 = make_float4(1.0f, 2.0f, 3.0f, 4.0f);

t通过.x, .y, .z, and .w访问

int2 i2 = make_int2(1, 2);
int x = i2.x;
int y = i2.y;

数学函数

   部分函数列表

   sqrt, rsqrt均方根

   exp, log指数

  sin, cos, tan, sincos三角函数

  asin, acos, atan2

  trunc, ceil, floor

Intrinsic function内建函数

   仅面向Device设备端

    更快，但精度低

    以_为前缀，例如：_exp, _log, _sin, _pow,_

线程层次

A thread block is a batch of threads that can cooperate with each other by:

                                            Sychronizing their exectution

                                                             For hazard-free shared memory accesses

                                                 Efficiently sharing data through a low latency shared memory

Two threads from two different blocks cannot cooperate

一个Kernel启动在设备端启动一个完整的线程grid，一个线程的grid包含了若干个线程块block，线程块的数目和每个线程块里面线程的数目都是开发者指定的，在global函数调用前面有三个尖括号里面的两个数字。

int threadID = blockIdx.x * blockDim.x + threadIdx.x;//通过grid和block坐标计算线程ID，为了方便索引二维三维数据
float x = input[threadID];//通过线程id把输入位置元素读入到线程的局部变量
float y = func(x);//在输入数据上执行函数；数据可并行
output[threadID] = y;//用线程id存储输出结果

线程索引标志着线程在整个程序中的位置，上面的设备端代码实际上是一个线程代码。假设它是一维的，就可以计算在整个线程块上线程的位置，确定线程索引，通过input写入局部变量，通过变量引用相关的函数func处理，把结果在写出到output对应的位置。

线程同步

  块内线程可以同步

         调用__syncthreads创建一个barrier栅栏,GPU端的线程代码里调用

         每个线程在调用点等待块内所有线程执行到这个地方，然后所有线程继续执行后续指令

 

Mds [i] =Md[j];
_syncthreads();//要求所有线程都执行到这个位置之后，在继续往下执行
func(Mds[i], Mds[i+1]);

同步对线程代码准确度很高。

线程同步要求线程的执行时间尽量接近，负载均衡很重要，某一个线程运行很长，其他的线程都会等待。只在一个块内进行，全局同步开销很大，块内同步可以自主调度，不用等待其他块，对可扩展性和适应性有好处。

同步破坏了并行性，独立性，__syncthreads()，导致线程暂停，也可能导致线程死锁。对程序执行造成致命的错误。

 

线程调度：从软件启动的线程数，远远大于整个硬件可用的执行部件。

一个硬件处理核心SM（Streaming Multi-Processor），一个SM有若干个处理核心SP/ALU  (Streaming Processing)  ，每一个SP可承载一个实际的线程。

G80 

   16个SMs

    每个含有8个SPs

        总共128个SPs

    每个SM驻扎多达768个线程，上下文空间

   总共同时执行12288（16*128）个线程

调度不意味着就在执行，同一个时钟周期上同时执行。

GT200 

   30个SMs

    每个含有8个SPs

        总共240个SPs

    每个SM驻扎多达1024个线程（8个block），上下文空间

   总共同时执行30720（240block）个线程

 

warp - 块内的一组线程

    -G80/GT200-32个线程

    -运行于同一个SM

    -线程调度的基本单位

    -threadIdx值连续,下标值连续

   -一个实现细节--理论上

                  warpSize

  因为一个SM的SP数目是固定的，调度过程，block线程数很多，调度基本单位不能用block，应该用更小的单元。

一个warp内部的线程天然就是同步的。

一个warp线程执行到相同位置。

同一个SM上，调度3块block时的warp。

 

线程调度主要的目的，利用线程独立的相同代码充分占据停滞的空隙，达到延迟掩藏效果。

SM implements zero-overhead warp scheduling在一个硬件上warp开销是0开销，所有的warp上下文实际上存在于物理空间内，需要执行的时候直接切换过来

•         At any time , only one of the warps is executed by SM 在一个SM上，在任何时刻都只有一个warp在执行
•        Warps whose next instruction has its operands ready for consumption are eligible for execution
•        Eligible Warps are selected for execution on a prioritized scheduling policy
•        All threads in a warp execute the same instruction when selected。

当某个warp停下来的时候SM硬件资源被其他warp占用。

如果warp内部线程沿着不同分支执行，divergent warp，warp必须步调一致，调度器没法为每一个ALU设计一个调度机构，如果为每一个ALU设计一个调度机构，芯片成本就会发生很大的变化，出现1/8性能，其他线程等到某一线程。

 

如果一个SM分配了3个block，其中每个block含有256个线程，总共有多少个warp？一个block有多少个warp？warp大小是32，一个block=256/32=8，一个block有8个warp，一个sm上有3个block，则有3*8=24个warp

GT200的一个SM最多可以烛照1024个线程相当于多少个warp  ，1024/32=32warp

每个warp有32个线程，但是每个SM只有8个SPs（ALU），32/8=4次，分批次处理。

当一个SM调度一个warp时

    指令已经预备

    在第一个周期8个线程进入SPs

    在第二、三、四个周期个进入8个线程

    因此，分发一个warp需要4个周期

一个kernel包含

    1次对 global memory的读写操作（200cycles）

    4次独立的multiples/adds操作

需要多少个warp才可以隐藏内存

    每个warp包含4个multiples/adds操作，每一个假定4个周期

         16个周期

     需要覆盖200个周期

         200/16=12.5

       ceil(12.5)=13

需要13个warp

   

 Device code can :GPU设备，不同部分存储器

          -R/W per-thread registers  读写每一个线程的私有寄存器

          -R/W per-thread local memory  

          -R/W per-block shared memory每一个线程块有一个公共的共享存储

         -R/W per-grid global memory  读写所有线程共享的显存上的global memory

         -Read only per-grid constant memory独立的存储空间，固定值的存储器，能够在多个线程在使用一个不太变化的内存，只能读

Host code can： CPU主机端代码

        -R/W per grid global and constant memories 读写global memory核constant memory

寄存器Registers

        每个线程专用，私有寄存器

        快速，片上，可读写，寄存器在芯片上面的

         增加了kernel的寄存器，寄存器增加了，每一个线程的速度会增加，计算单元减少，SM线程数减少

 

寄存器Registers-G80

    每个SM   

       多达768threads

       8k个寄存器

 每个线程可以分到多少资源8k/768=10个寄存器/线程。

超出限制后台线程数将因为block的减少而减少。

例如每个线程用到了11个寄存器，并且每个block含有256个线程

    一个SM可以驻扎768/256=3，但是只有2个block（寄存器增大），因此只有2*256=512个线程

    一个SM可以驻扎512/32=16个warp

warp数量变少意味着有资源的浪费，效率下降了。

本开是可以装三个768个线程，3个block，但只要2个block，一共有512个线程可以同时驻扎，剩下的线程只能在寄存器里面，不够分，只能在寄存器里闲着。

局部存储器localmemory

    存储于globalmemory

       作用域是每个thread私有

   用于存储自动变量数组

          通过常量索引访问。A[5]编译器通过localmemory放到globalmemory位置中去。

 

 

共享存储器shared memory

       每一个块block

       快速、片上、可读写在cache在同一个层次

       全速随机访问

 

 

内存是竞争资源，约束block数目和线程数目

共享存储器shared memory -G80

  每一个SM包含

          多达8个blocks

          16KB共享存储器

       每个block分配了16/8=2KB

若每个block用5KB，则只能用16/5=3个block

全局存储器global memory

    长时延（100个周期）

    片外。可读写

    随机访问影响性能

Host主机端可读写

   GT200

      带宽：150GB/s带宽很大

      容量：4GB

G80-86.4GB/S带宽很大

 

 

常量存储器constant memory

   延时短，高带宽，当所有线程访问同一位置时只读

  存储与global memory 但是有缓存

  host主机端可读写

容量64KB

存储常量，滤波器系数等

 

 

变量声明	寄存器	作用域	生命周期
必须是单独的自动变量而不能是数组	register	thread	kernel
自动变量数组	local	thread	kernel
__shared__ int sharedVar;	shared	block	kernel
__device__ int deviceVar;	global	grid	application
__constant__ int constantVar;	constant	grid	application

Global和Constant变量

Host 可以通过以下函数进行访问

cudaGetSymbolAddress()；主机在设备上找到特定变量的地址，

cudaGetSymbolSize()；大小

cudaMemcpyToSymbol()；拷贝变量位置

cudaMemcpyFromSymbol()；变量位置拷贝回来

Constants必须在函数外声明：