cuda并行程序设计复习(基础概念、矩阵相乘)
第一章
-
CPU和GPU的设计非常不同
CPU:面向延时的内核设计,有较大的控制单元与缓存空间
强大的ALU可以较少操作延时, 大型的缓存,减少长延迟的内存访问转换为断延时的高速缓存访问 复杂的控制单元:用于分支延迟和预测,减少数据转发延迟
GPU:面向吞吐量的设计核心,具有较多的SIMD单元
小型的缓存为了提高内存的访问量;简单的控制单元,没有分支预测与数据 转发;高能效的ALU,大量延时长但是大量的流水线型运行吞吐量巨大;大量 的线程运行以减少线程逻辑与状态的时间花销
-
系统成本的增加
每个芯片的SW线每10个月翻倍
HW门数每18个月翻倍
-
可扩展性
同样的应用可以在新一代的核心上高效运行,同样的应用在更多的相同核心上高效运行。
随之硬件的不断改进
- 越来越多的计算单元(核心)数量
- 越来越多的线程数量
- 矢量长度增加
- pipeline长度增长
- DRAM大小增加
- DRAM通道数增加
- 减少数据移动的延迟
便携性:同一个应用可以在不同类型的核心上高效运行;在不同的组织和界面上高效运行
―跨越ISA(指合集架构)–X86与-ARM,等等。-面向延迟的CPU与面向吞吐量的GPU的比较
–跨平行度模型–VLIW vs.SIMD vs.线程
―跨越内存模型–共享内存与分布式内存
第二章:
-
CUDA加速应用:CUDA库,编译指导;编程语言
-
CUDA库:
更加简单使用:可以直接使用CUDA库而不用了解GPU编程
Drop-in:许多的GPU加速库遵循标准的API,可以使用少量的代码改动而实现加速
质量优:CUDA库对广泛的应用场景提供了高质量的实现,例如线性代数、数值计算,数据分析与人工智能、计算机视觉处理
-
编译指导: Easy, Portable Acceleration
易用性:编译器负责并行性的管理与数据移动的细节
便携性:代码通用,不特定指定硬件类型,可以部署到多种语言之中
不确定性:不同的编译器版本的代码性能会有所差异
例如 openACC C、C++和FORTRAN的编译器指令
#pragma acc parallel loop copyin(input1[0:inputLength],input2[0:inputLength]), copyout(output[0:inputLength]) for(i = 0; i < inputLength; ++i) { output[i] = input1[i] + input2[i]; }
-
编程语言:更加灵活与与高效
性能优:程序员可以对数据并行与移动有最好的控制
灵活性:并行计算并不限制于有限的库以及指令集,可以自定义实现
冗余性:程序员需要书写更多的细节内容
例如 Numerical analytics > MATLAB Mathematica,LabVIEWCUDA Fortran
Fortran > CUDA Fortran
C > CUDAC
C++ > CUDA C++
Python > PyCUDA,Copperhead,Numba
F# > Alea.cuBase
-
CUDA C主机代码基本API
设备代码可以 R/W每个线程的寄存器,所有共享的全局存储
主机代码可以向每个网格传输全局存储器数据
cudaFree(d)
cudaMalloc(To,From,size,type)
cudaMemcpy(To,From,size,Type)
(异步传输)
cudaError_t
错误类型
-
程序都是存储在内存中的一组指令,可以由硬件读取、解释和执行
CPU和GPU都是基于不同的指令集设计
-
程序指令可以对存储在内存中或者寄存器中的数据进行操作
-
冯诺依曼结构:存储器,输入、输出、控制器、运算器
-
一个网格中的所有线程都运行相同的内核代码(单程序多数据)
-
一个块内的线程通过共享内存、原子操作和栅栏同步与不同的区块内的线程互不影响
-
CUDA 4.0以后支持三维id
-
CUDA核心包含三个重点抽象:线程组层次、共享
存储器和栅栏同步。 -
一个kernel函数中只有一个grid
第三章 CUDA并行模型
-
函数前缀
执行 调用 __device__
device device __global__
device host __host__
host host A kernel function must return void __device__ and __host__ 可以同时使用
-
灰度颜色计算公式
P=0.21*r+0.71*g+0.07*b
-
图像模糊代码
__global__ void blurKernerl(uchar in,uchar out,int w,int h){
int col = blockIdx.x*blockDim.x+threadIdx.x;
int row = blockIdx.y*blockDim.y+threadIdx.y;
if(col<w && row<h){
int pixVal = 0;
int pix = 0;
for(int i = -blur_size;i<blur_size+1;i++)
for(int i = -blur_size;i<blur_size+1;i++){
int colc = col+i;
int rowc = row+j;
if( (colc >-1 && colc<w) && (rowc>-1 &&rowc <h)){
pixVal+ = in[rowc*w+colc];
pix++;
}
}
out[row*w+col] = (uchar)pixval/pix;
}
}
-
可扩展性
- 每个区块可以以相对于其他区块的任何顺序执行。
- 硬件可以在任何时候*地将块分配给任何处理器-
- 一个内核可以扩展到任何数量的并行处理器
-
线程执行块
一般情况一个SM最多8个块
费米架构一个SM由1536个线程,可以分为
256*6或者512*3
SM负责维护线程块的idx,管理线程执行~~~~
-
Warp
Warp一般为32个线程,warp为SM的一个调度单位,由两个16线程并排,未来可能会变化。
例如:在三个块的SM中,每块由256个线程,一共多少warp
每块由256/32=8 个warp 一共就是 3*8=24 个warp
例子 费米架构下 8x8,16x16,32x32谁能最好充分利用线程数
- 对于8×8,我们每块有64个线程。由于每个SM最多可容纳1536个线程,这相当于24个Block。然而,每个SM只能占用8个区块,每个SM只能有512个线程。
- 对于16X16,我们每块有256个线程。由于每个SM最多可以占用1536个线程,它最多可以占用6个Block并实现满负荷运转,除非有其他资源考虑。
- 对于32×32,我们将有每块1024个线程。对于费米来说,只有一个区块可以装入一个SM。只使用一个SM的2/3的线程容量。
第四章 内存与局部数据
- 矩阵乘法 不适用共享内存
__global__ void Mat(float *M,float* N,float* P,int wid){
int col = blockIdx.x*blockDim.x + threadIdx.x;
int row = blockIdx.y*blockDim.y + threadIdx.y;
if(row < wid && col < wid){
float pvalue=0;
for (int k = 0;k<wid;k++){
pvalue += M[row*wid +k] * N[k*wid + col]
}
P[row * wid +col] = pvalue;
}
}
- CUDA 变量
val | 存储 | 范围 | 生命周期 |
---|---|---|---|
int local | register | thread | thread |
__device__ __shared__ int shareval |
shared | block | block |
__device__ int Globalval |
shared | grid | applicaton |
__device__ __constant__ int cont |
constant | grid | –applicaton |
-
多个修饰时
__device__
可选,自动变量为寄存器变量,数组变量为全局变量 -
global constant变量在函数外声明 shared local在函数内声明
-
共享内存
每个SM中有一个,访问速度远高于全局内存;访问范围为blcok;生命周期也为block;由内存加载/存储;计算机体系结构中的暂存存储器。
-
TIle思想
- 通过多线程来表示全局内存的块
- 通过Tile将全局内存加载至on-chip 内存
- 通过栅栏同步实现通信
- 让多个线程共同访问相同数据
-
共享内存矩阵乘法
__global__ Mat_gl(float* M,float* N,float* P,int wid){ __shared__ int Ms[tile_wid][tile_wid] __shared__ int Ns{tile_wid][tile_wid] int tx = threadIdx.x,ty = threadIdx.y; int bx = blockIdx.x, by = blockIdx.y; int col = bx*blockDim.x+tx; int row = by*blockDim.y+ty; if(row<wid &&col<wid){ float pvalue = 0; for(int p = 0;p>N/tile_wid;p++){ Ms[ty][tx] = M[row * wid + p * tile_wid + tx]; Ns[ty][tx] = N[(p*tile_wid + ty) * wid + col]; __syncthreads(); for(int i=0;i<tile_wid;i++) pvalue += Ms[ty][i] * Ns[i][tx]; __syncthreads(); } P[row*wid + col] = pvalue; }}
-
对于16X16的tile,一共256个线程,在每个阶段加载全局内存的数量为2X256=512 次,执行乘法/加法运算次数为 256 X (2X16) = 8192 每次加载内存有16次浮点运算
-
对于32X32的tile,一共1024个线程,在每个阶段加载全局内存的数量为 2X1024=2048次,执行乘法/加法运算次数为 1024 X (2X16) = 65536,每次加载内存有32次浮点运算
-
一个共享内存的大小为16KB的SM,对于TILE_WIDTH = 16 ,每个线程块使用2x256x4B = 2K的共享内存数据,所以最大支持8个线程块执行,一共加载了8x512=4096次
-
对于TILE_WIDTH = 32,每个线程块使用 2x512x4B最多运行两个块,但是总的线程数为1536,这就只能让一个块运行
-
每个
__syncthreads()
可以减少活动的线程数量’