7 内核执行

时间：2024-03-25 16:25:04

文章目录

  • 内核程序如何在GPU上执行的。
    • 它们如何启动,有哪些执行特性,
    • 如何由thread组成block
    • 再由线程块组成grid
    • 资源管理涉及哪些因素。
  • 还包括对动态并行的描述
    • 它是CUDA5.0的新特性
    • 可使CUDA内核程序为GPU启动工作任务。

7.1 概况

  • CUDA内核程序在GPU上执行,
  • 且从早期的CUDA版本开始就一直与CPU并发执行。
  • 即：
    • 内核启动是异步的:
    • 控制权会在GPU完成请求的操作之前返给CPU。
  • CUDA最初引进时,开发者不需关心
    • 内核启动的异步执行(或缺少此机制)。
  • 数据显式地复制到GPU并从GPU复制回,内存复制命令会在
    • 启动内核程序的命令请求后进行排队。
  • 想通过编写CUDA代码来暴露内核启动的异步特性是不可能的,
    • 异步执行的主要附带作用是
    • 在连续执行多个内核启动时
      • 隐藏驱动程序的开销。

 

  • 从前面对映射锁页内存(主机内存可直接被GPU访)可看出
  • 对于写入主机内存的(相反的,从其中读取的)内核程序而言,内核启动的异步执行很重要。
  • 若一个内核在没有显式同步(例如采用CUDA事件)的情况下启动或写入主机内存,代码将面临CPU与GPU之间的竞争并会出现运行错误。
  • 对于通过映射锁页内存进行读操作的内核,一般不要显式同步,
    • 因为CPU的任何待定写操作都会在内核启动之前被发送出来。
  • 要是内核程序通过写入映射锁页内存来给CPU返回结果,就必须使用同步避免读后写了。

 

  • 一旦内核启动,
  • 它以一个网格的形式运行,网格含多个线程块,线程块由多个线程
    • 并非线程块的运行都并发
    • 每个线程块都被分配给一个SM
    • 每个SM可为多线程块维持上下文。
  • 为掩盖内存和指令的延时,SM常会需要多于单线程块可以包含的线程束
    • (SM2.0及以上,一个线程块可含1024线程)
  • 每个SM中线程块的最大数无法通过API查询,但英伟达文档写明,SM3.x版本之前的硬件中为8,SM3.x及其之后的为16

 

  • 编程模型不保证执行次序,
  • 或是某些线程块或线程是否可并发
  • 开发者不要假定内核启动中的所有线程都是并发执行的。
  • 通常很容易就启动了比机器所能维持的要多的线程,
    • 且其中的一些在其他线程结束之前不会执行。
  • 线程次序无法确定,
    • 甚至连内核启动开始处进行全局内存的初始化都是困难任务

 

  • 动态并行————特斯拉K20(GK110)(第一个支持SM3.5的GPU)中新增特性
  • 使内核可以启动其他的内核程序,
    • 并完成它们之间的同步。
  • 这些特性解决了一些以前的硬件上CUDA所呈现的局限性。
    • 如,一个动态并行内核可通过启动和等待一个子网格来完成初始化。

7.2 语法

  • 用CUDA运行时的时候,内核这样启动
    • Kernel<<<gridSize, blockSize, sharedMem,Stream>>>(Parameters…)
  • 待启动的内核名
  • 指定一个dim3结构形式的网格的大小;
  • 指定了一个dim3结构形式的线程块的维度;
  • 指定为每个线程块预留的附加的共享内存
    • 内核程序可用的共享内存的数目为
    • 在此参数值上加上内核中静态申请的共享内存数
  • 指定内核启动所属于的流

 

  • 指定网格和块大小的dim3结构含有3个成员变量(x、y和z)。
  • 用C++编译时,一个带有系统默认参数的构造函数会将x、y和z均默认初始化为1。
  • 详见清单7-1,摘自 NVIDIA SDK中的 vector_types.h
    7 内核执行

 

  • 内核程序可通过驱动程序API使用cuLaunchKernel来启动
  • cuLaunchKernel将网格和线程块的维度作为独立参数而不是dim3结构。

7 内核执行

 

  • 对于三对尖角括号语法,
  • cuLaunchKenel参数包括
    • 调用的内核、网格和线程块大小、共享内存数目及流。
  • 区别在于内核本身的这些参数如何给出:
    • ptxas发射出的内核微码包含了描述各个内核参数的元数据（注！！）,
    • so kernelParams是个void*类型的数组,
    • 每个元素对应内核的一个参数。
  • 参数的类型可被驱动程序识别,所以参数所占的内存数据将会作为用来调用内核的、具有硬件特性的命令的一部分复制到命令缓冲区中。

注

  • cuLaunchKernel()不适用于那些没在CUDA3.2或更高版本上编译的二进制映像文件,
    • CUDA3.2是第一个可包含内核参数元数据的版本

7.2.1 局限性

  • 所有参与内核启动的C+类都必须是带有以下特性的“简单旧数据”( plain old data POD)。

  • 无用户声明的构造函数;

  • 无用户定义的复制分配操作符;

  • 无用户定义的析构函数;

  • 无非POD的非静态数据成员;

  • 无私有或保护型的非静态数据

  • 无基类;

  • 无虚函数。

 

  • 注意,违背了这些规则的类也可用于CUDA,甚至是CUDA内核程序中。
  • 只是,它们绝对不能使用于内核启动中。
  • 如此,一个CUDA内核使用的类可以通过使用来自内核启动的POD输入数据来构造。

 

  • CUDA内核程序也没有返回值。
  • 它们必须通过设备内存(必须显式地复制回CPU)或映射主机内存来传回结果。

7.2.2 高速存和一致性

 

  • GPU含多个高速缓存以便于在发生重用时加速计算。
  • 常量缓存被充分利用起来以便于
    • 广播式传输到同一个SM的执行单元;
  • 纹理缓存减少了外部带宽的使用。
  • 这两种缓存都不能很好地同GPU的写内存操作保持一致性。
  • 没有促使这些缓存和一、二级缓存之间保持一致性的协议,
    • 所以无法减少全局内存的延迟和支持聚合带宽。

 

  • 这意味着两件事:
  • 当一个内核运行时,
    • 不要对那些同时(或者是被一个并发运行的内核)正在通过常量内存和纹理内存进行访问的内存执行写内存的操作。
  • CUDA驱动程序必须在每个内核启动之前使常量缓存和纹理缓存无效。

 

  • 对不含TEX指令的内核程序,
  • CUDA驱动程序不需要使常量缓存和纹理缓存无效。
  • 因此,未使用纹理的内核程序引发更少的驱动程序开销。

7.2.3异步与错误处理

  • 内核启动是异步的,因为一旦内核被提交到硬件就会立即与CPU并行执行（注！！）。
  • 这一异步会使错误处理变得很复杂。
  • 如果一个内核遇到一个错误(例如,它读入ー个无效内存位置),该错误有时会在内核启动后的一段时间才能传输到驱动程序(和应用程序)中。
  • 检査这种错误最可靠的方法是用 cudaDeviceSynchronize()或 cuctxsynchronize)函数以使其与GPU同步。
  • 如果在内核执行中出现一个错误,会返“ unspecified launch failure”的错误代码。

 

  • 除了 cudaDeviceSynchronize()或 cuCtxSynchronize()等显式CPU/GPU同步函数外,
  • 这个错误代码还可能是来自与CPU隐式同步的函数,如同步的内存复制调用。

注

  • 大多数平台上,内核会在CPU处理完启动命令之后的数个微秒开始在GPU执行。
  • 但在WDDM( Windows Display Driver Model)平台上,它可能需要更长的时间。
  • 因为驱动程序必须执行一个内核转换以便将启动提交到硬件中,并且在用户模式下GPU的工作任务会进入队列,平摊用户态到内核态过渡的开销。

无效内核启动

  • 有可能所请求的内核启动无法为硬件执行。
  • 例如,指定比硬件能够支持的块内线程数还要多的线程就会出现这种情况。
  • 驱动程序会尽可能检测到这些情况并报告错误,而不是试图将该启动提交给硬件。

 

  • CUDA运行时和驱动程序API以不同的方式处理这种情况。
  • 当一个无效的参数被指定时,驱动程序API的显式API调用(例如 culaunchgrid()和 culaunchkernel(等函数)返回
    错误代码。
  • 但是,当使用CUDA运行时的时候,由于内核是按CC+一行代码启动的,故而没有API调用来返回错误代码。
  • 对应地,那个错误会被“记录”到本地线程槽中,而应用程序则可以通过 cudagetlasterror()函数来查询该错误值。
  • 与此相同的错误处理机制还被用在因其他原因(例如,内存访问冲突)导致的无效内核启动上。

7.2.4 超时

这儿没写！！！！

7.2.5 本地内存

  • 本地内存是线程私有的,且CUDA中的网格可以包含数千个线程,因此CUDA网格需要的本地内存数目是相当多的。
  • CUDA开发者会用心预先分配资源以降低内核启动等操作因为缺乏资源而失败的可能性。
  • 但就本地内存而言,仅一个保守的分配就已经消耗了太多的内存。
  • 缘于此,使用大量本地内存的内核需要更长的时间且有可能是同步执行的,因为CUDA驱动程序必须在执行内核启动之前分配好内存。
  • 此外,如果内存分配失败,内核也将会因缺乏资源而启动失败。

 

  • 默认情况下,当CUDA驱动程序必须分配本地内存以运行内核时,它会在内核完成之后释放内存。
  • 另外,这种行为还会使得内核启动同步。
  • 但这种行为可以通过给cuCtxCreate()指定 CU_CTX_LMEM_RESIZE_TO_MAX标志
    • 或在创建主上下文之前调用cudaSetDeviceFlags()函数置cudaDeviceLmemResizeToMax来禁止。
  • 它带来的结果是,在一个需要比默认更多本地内存的内核启动之后,增加的可用本地内存不会被释放。

7.2.6 共享内存

  • 共享内存在内核启动时分配,在内核执行期间一直保持。
  • 除了可在内核中声明的静态分配方法之外,共享内存也可被声明为一个未指定大小的extern变量。
  • 用于分配大小不定的数组的共享内存量
    • 由内核启动的第三个参数
    • 或cuLaunchKernel的 sharedMemBytes指定。

7.3 线程块、线程、线程束、束内线程

  • 内核以线程块构成的网格进行启动。
  • 线程可分为32个线程组成的线程束(warp),
    • 线程束中的单个线程被称为一个束内线程。

7.3.1线程块网格

  • 线程块独立被调度到SM中,来自于同一线程块的线程在同一SM中执行。
  • 图7-1是二维线程块(8Wx8H)组成的二维网格(8Wx6H)。
  • 图7-2是三维线程块(8Wx8Hx4D)组成的三维网格(8Wx6H×6D)。

7 内核执行

7 内核执行

  • 网格由65535×65535个线程块(SM1.0的硬件)或65535×65535×65535的线程块(SM2.0的硬件)组成。
  • 每个线程块由512或1024个线程组成,线程块中的线程之间可通过SM的共享内存通信。
  • 一个网格中的线程块有可能会被分配给不同的SM。
  • 为使硬件吞吐量最大化,一个给定的SM可以在同一时间内运行来自不同线程块的线程与线程束。
  • 当所需要的资源变得可用,线程束的调度器会分派指令。

 

  • 网格的最大尺寸 CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X等查询,
    • 亦可通过cudaGetDeviceGetproperties并检查cudaDeviceProp::maxGridSize查询
  • 线程块的最大尺寸通过CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK
    • 或 deviceProp.maxThreadsPerBlock

1 线程

here!!!

相关文章