投稿
  1. 首页
  2. 随笔

cuda并行程序设计复习（基础概念、矩阵相乘）

csi • 2022-12-17 • 随笔 • 阅读 26

cuda并行程序设计复习（基础概念、矩阵相乘）,第1张

cuda并行程序设计复习（基础概念、矩阵相乘） 第一章

  1. CPU和GPU的设计非常不同

    CPU：面向延时的内核设计，有较大的控制单元与缓存空间

     	强大的ALU可以较少 *** 作延时，

 	大型的缓存，减少长延迟的内存访问转换为断延时的高速缓存访问

 	复杂的控制单元：用于分支延迟和预测，减少数据转发延迟

    GPU：面向吞吐量的设计核心，具有较多的SIMD单元

     		小型的缓存为了提高内存的访问量；简单的控制单元，没有分支预测与数据				   

 		转发；高能效的ALU,大量延时长但是大量的流水线型运行吞吐量巨大；大量

 		的线程运行以减少线程逻辑与状态的时间花销

  2. 系统成本的增加

    每个芯片的SW线每10个月翻倍

    HW门数每18个月翻倍

  3. 可扩展性

    同样的应用可以在新一代的核心上高效运行，同样的应用在更多的相同核心上高效运行。

    随之硬件的不断改进

    • 越来越多的计算单元（核心）数量
    • 越来越多的线程数量
    • 矢量长度增加
    • pipeline长度增长
    • DRAM大小增加
    • DRAM通道数增加
    • 减少数据移动的延迟

    便携性：同一个应用可以在不同类型的核心上高效运行；在不同的组织和界面上高效运行
    ―跨越ISA(指合集架构)–X86与-ARM，等等。

    -面向延迟的CPU与面向吞吐量的GPU的比较

    –跨平行度模型–VLIW vs.SIMD vs.线程

    ―跨越内存模型–共享内存与分布式内存

第二章：

  1. CUDA加速应用：CUDA库，编译指导；编程语言

  2. CUDA库：

    更加简单使用：可以直接使用CUDA库而不用了解GPU编程

    Drop-in：许多的GPU加速库遵循标准的API，可以使用少量的代码改动而实现加速

    质量优：CUDA库对广泛的应用场景提供了高质量的实现，例如线性代数、数值计算，数据分析与人工智能、计算机视觉处理

  3. 编译指导： Easy, Portable Acceleration

    易用性：编译器负责并行性的管理与数据移动的细节

    便携性：代码通用，不特定指定硬件类型，可以部署到多种语言之中

    不确定性：不同的编译器版本的代码性能会有所差异

    例如 openACC　C、C++和ＦＯＲＴＲＡＮ的编译器指令

    #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];
}

  4. 编程语言：更加灵活与与高效

    性能优：程序员可以对数据并行与移动有最好的控制

    灵活性：并行计算并不限制于有限的库以及指令集，可以自定义实现

    冗余性：程序员需要书写更多的细节内容

    例如　Numerical analytics > MATLAB Mathematica,LabVIEWCUDA Fortran
          Fortran > CUDA Fortran
          C > CUDAC
          C++ > CUDA C++
          Python > PyCUDA,Copperhead,Numba
          F# > Alea.cubase

  1. 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并行模型

  1. 函数前缀

    执行调用__device__devicedevice__global__devicehost__host__hosthostA kernel function must return void
    __device__  and __host__ 可以同时使用

  2. 灰度颜色计算公式P=0.21*r+0.71*g+0.07*b

  3. 图像模糊代码

__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-1 && colc-1 &&rowc  

  1.  
    可扩展性
     
  
    • 每个区块可以以相对于其他区块的任何顺序执行。
    • 硬件可以在任何时候自由地将块分配给任何处理器-
    • 一个内核可以扩展到任何数量的并行处理器
     
  2.  
    线程执行块
     
    一般情况一个SM最多8个块
     
    费米架构一个SM由1536个线程，可以分为256*6或者512*3
     
    SM负责维护线程块的idx，管理线程执行~~~~
     
  3.  
    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的线程容量。
 
第四章 内存与局部数据 

  1. 矩阵乘法 不适用共享内存
 

__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 

  1. CUDA 变量
 
val存储范围生命周期int localregisterthreadthread__device__ __shared__ int sharevalsharedblockblock__device__ int Globalvalsharedgridapplicaton__device__ __constant__ int contconstantgrid–applicaton

  •  
    多个修饰时__device__可选，自动变量为寄存器变量，数组变量为全局变量
     
  •  
    global constant变量在函数外声明 shared local在函数内声明
     
 

  1.  
    共享内存
     
    每个SM中有一个，访问速度远高于全局内存；访问范围为blcok；生命周期也为block；由内存加载/存储；计算机体系结构中的暂存存储器。
     
  2.  
    TIle思想
     
  
    • 通过多线程来表示全局内存的块
    • 通过Tile将全局内存加载至on-chip 内存
    • 通过栅栏同步实现通信
    • 让多个线程共同访问相同数据
     
  3.  
    共享内存矩阵乘法
     
    __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(rowN/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 
 

  •  
    对于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()可以减少活动的线程数量’
     
					
										


					

						
欢迎分享，转载请注明来源：内存溢出
原文地址: http://outofmemory.cn/zaji/5688634.html
赞 (0)
打赏 微信扫一扫 微信扫一扫 支付宝扫一扫 支付宝扫一扫
csi csi 一级用户组
0 0
上一篇 2022-12-17
下一篇 2022-12-17

发表评论

登录后才能评论

评论列表（0条）

保存