内置类型和函数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必须在函数外声明: