找回密码
 立即注册

QQ登录

只需一步,快速开始

查看: 42|回复: 2

DAY47:阅读Atomic Functions

[复制链接]
发表于 2018-7-9 11:24:56 | 显示全部楼层 |阅读模式

B.12. Atomic Functions
An atomic function performs a read-modify-write atomic operation on one 32-bit or 64-bit word residing in global or shared memory. For example, atomicAdd() reads a word at some address in global or shared memory, adds a number to it, and writes the result back to the same address. The operation is atomic in the sense that it is guaranteed to be performed without interference from other threads. In other words, no other thread can access this address until the operation is complete. Atomic functions do not act as memory fences and do not imply synchronization or ordering constraints for memory operations (see Memory Fence Functions for more details on memory fences). Atomic functions can only be used in device functions.
On GPU architectures with compute capability lower than 6.x, atomics operations done from the GPU are atomic only with respect to that GPU. If the GPU attempts an atomic operation to a peer GPU’s memory, the operation appears as a regular read followed by a write to the peer GPU, and the two operations are not done as one single atomic operation. Similarly, atomic operations from the GPU to CPU memory will not be atomic with respect to CPU initiated atomic operations.
Compute capability 6.x introduces new type of atomics which allows developers to widen or narrow the scope of an atomic operation. For example, atomicAdd_system guarantees that the instruction is atomic with respect to other CPUs and GPUs in the system. atomicAdd_block implies that the instruction is atomic only with respect atomics from other threads in the same thread block. In the following example both CPU and GPU can atomically update integer value at address addr:
  1. __global__ void mykernel(int *addr) {
  2.   atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
  3. }

  4. void foo() {
  5.   int *addr;
  6.   cudaMallocManaged(&addr, 4);
  7.   *addr = 0;

  8.    mykernel<<<...>>>(addr);
  9.    __sync_fetch_and_add(addr, 10);  // CPU atomic operation
  10. }
复制代码


The new scoped versions of atomics are available for all atomics listed below only for compute capabilities 6.x and later.
Note that any atomic operation can be implemented based on atomicCAS() (Compare And Swap). For example, atomicAdd() for double-precision floating-point numbers is not available on devices with compute capability lower than 6.0 but it can be implemented as follows:
  1. #if __CUDA_ARCH__ < 600
  2. __device__ double atomicAdd(double* address, double val)
  3. {
  4.     unsigned long long int* address_as_ull =
  5.                               (unsigned long long int*)address;
  6.     unsigned long long int old = *address_as_ull, assumed;

  7.     do {
  8.         assumed = old;
  9.         old = atomicCAS(address_as_ull, assumed,
  10.                         __double_as_longlong(val +
  11.                                __longlong_as_double(assumed)));

  12.     // Note: uses integer comparison to avoid hang in case of NaN (since NaN != NaN)
  13.     } while (assumed != old);

  14.     return __longlong_as_double(old);
  15. }
  16. #endif
复制代码




回复

使用道具 举报

 楼主| 发表于 2018-7-9 14:22:59 | 显示全部楼层
这章节开始将要介绍原子操作函数.原子操作系列函数是非常重要的内容, 同样对于CPU和GPU方面来说.
自从多核的普及(还能想象一下正好10年前, AMD推出了巴塞罗那系列4核CPU时候, 给大家带来的震撼吗),以及, GPU上还量并行的应用, 以及, 到多卡系统, CPU-GPU交互的普及, 不使用原子操作将会越来越变得寸步难行.
大约在Sisiy我刚刚尝试开始使用CUDA的时候, 我曾经面临过8800GTX(初代的1.0计算能力的卡, 不支持任何原子操作), 和当年的小珍珠GT240的艰难选择.
因为当时要解决从一个特定的图像中搜索匹配特定的小图像内容, 没有原子操作的8800GTX当时用起来非常艰难(因为要将可能的数量不定的结果compact起来), 最后只能选择后者.这就是曾经的一个例子.
回到今天这里. 我们已经恍惚间遭遇了6.0+了.从6.0开始, 原子操作性能如同本章所说, 得到了很大的提升,主要体现在应用范围的扩大, 以前只能在一张卡内使用,现在扩展到了系统内: 其他的伙伴卡(例如4卡系统), 以及, 和CPU之间交互.这可能是因为现在需要解决的问题规模越来越复杂, 特别是深度学习的逐渐应用,
能将卡的显存联合起来(PCI-E/NVLink P2P Access), 等于有效的扩大了容量,而通过跨卡的原子操作, 又能解决缓冲区边界的拼接问题(大型图像处理问题很常见)。
但是本章说的比较简单, 只告诉你现在可以有两种后缀可以选择了:
(1)原本的函数+_block后缀
(2)正常的原本函数(无后缀)
(3)原本的函数+_system后缀
变成了这三个级别. 是不是有点眼熟?(忘记了可以查看一下之前的memory fence章节以提供你熟悉的记忆场景来源)
根据手册里没有提供的资料, 但被NV官泄的GTC历年大会上的视频和幻灯片,类似atomicAdd_system()这种级别的(例如CPU+GPU), 是依靠于新的Pascal的细粒度的Unified Memory能力.NV暗示了不需要PCI-E 3.0 Atomics的支持, 即可正常使用.这对一些场合, 例如CPU和GPU需要同时处理一个任务多个部分的场合十分有用. 曾经我们的一个客户使用TX1(不是Pascal. Maxwell), 曾经抱怨过一个问题.她说, 我的当前GPU上的kernel实现比我的CPU快10倍. 但是还是不能满足我的要求,但是如果能同时算上CPU的计算性能, 则差不多正好,(CPU的核心们提供了好几十个百分点提升, 刚才的性能比较是单核的),然后客户继续说, 我尝试使用unified memory来做这个, 但是我有两个问题无法解决1)没法同时CPU上的原子操作和GPU上的原子操作协同起来(Pascal的系统级的原子操作支持). (2)Unified Memory不能同时使用, GPU在用unified memory的时候, CPU一用就segmentaion fault了.然后我们当时的建议是, 使用Zero-copy, 这个能同时用. 然后再想法规避掉CPU上的原子操作,可惜最后这个用户说, 能用是能用了, 但是很慢很慢.
如今到了Pascal+的时候, 原子操作得到了提升为系统级别.此时如果该用户继续购买了TX2, 应当这种问题很容易就解决了吧.因为6.0+允许GPU-GPU, CPU-GPU间的原子操作协作了.当然TX2没有多卡. 不过考虑到它是较弱的CPU(256 SP), 加上6个CPU核心很多时候也是很好的.
回复 支持 反对

使用道具 举报

 楼主| 发表于 2018-7-9 14:46:16 | 显示全部楼层
这也是AMD当年在还热心推广OpenCL的时候, 推出的OpenCL 2.0里面的一个重要的演示例子.它们的APU(类似TX2的, 也是CPU+GPU, 不过CPU是x86的, 不是arm, GPU也不是N卡), 单独CPU需要X秒, 单独GPU部分需要Y秒, 两者一起上全局原子操作, 只需要Z秒(Z比前两者都小不少).
NV目前此atom*_system(), 支持普通PCI-E(不需要PCI-E 3.0 Atomics硬件特性支持), 以及更好的NVLink.
然后本章节还说了一点:原子操作(atomic*()系列函数)本身无任何memory fence作用.(还记得memory fence吗? 之前的章节说过它的2大作用. 一个是软件(编译器)控制上的. 控制生成的语句的顺序,另外一个是硬件执行上的)。
atomic本身不会暗示编译器进行任何特殊处理.也不会在GPU上硬件执行的时候导致任何特殊fence效果.换句话说, 普通读写和原子操作完全不维持一致性的.不建议任何时候混用它们, 除非你知道你在做什么.特别是注意很多卡的shared memory(例如Maxwell)的普通写入和原子操作不维持一致性.同一个地址先普通写入值A, 然后再做一次, 例如atomicAdd + 1的操作, 将会导致未知的结果(Maxwell之前的shared memory上的原子操作是用普通读写模拟的).
此问题在Pascal+上得到了解决。
Maxwell可以暂时使用原子交换(写入+读取旧值)来patch一下.不要旧值部分即可.不过Maxwell/Pascal+的shared memory上的原子操作性能的确得到了海量提升.在之前的版本中(Maxwell之前的Kepler和Fermi), 这个操作是用的SP的模拟循环进行的: 尝试锁定, SP们计算, 尝试竞争性的写入, 如果竞争失败重新来这个过程.以前的这个过程在密集的冲突的时候(例如对邻近的多个甚至1个地址上进行原子操作)性能会很惨的.SP将大量无用的空转.Maxwell开始引入了GCN(没错. A家的)的"远程原子操作", 所有原子操作指令直接提交给shared memory自己执行.感兴趣的可以看一下当年的GTC的介绍的细节, 当时这个叫"remote atomics", 直译过来是"远程原子操作"
这也是为何推荐GTC的原因, NV不像AMD那样的公开一切, 很多时候直接手册里一句, 怎么怎么用就完了,甚至有的时候直接还会丢下一句, 需要****特性的可以用介绍的PTX嵌入来实现.而GTC是很好的一个信息来源, 避免对细节了解不清楚写出错误的实现,或者在实现的时候疑神疑鬼(例如很多客户不敢下手写东西, 总是怀疑: ****这样会*****吗——太常见了。
此外, 从计算能力6.0(Pascal)开始, 引入了double版本的原子操作.以前的这个操作是不能被直接支持的.但是可以通过"万能"的CAS(比较为真则交换)这个操作来模拟出来.所谓CAS原子操作是万能的原因是, 其他的原子操作都可以通过它来模拟出来(不考虑效率),就如同数字电路中,NOR和NAND门, 这两个门是万能的一样.你可以有AND, OR, NOT, 做为3个基础.但也可以只需要NOR或者NAND之一, 即可实现其他所有逻辑.CAS也同理.本章种给了一段传统的实现atomicAdd(double版本)的方式. 欢迎大家阅读. 代码还是很好理解的.
需要注意的是, 正常情况下的原子操作精度较低(和计算能力有关),例如在对subnormal number的处理上(subnormal是指当浮点数的绝对值小到很小一个程度的时候, 浮点数(half, float, double)自动切换成定点数的形式, 此时指数固定为一个特殊值),很多时候, 原子操作不能处理subnormal数据的, 将被强制当成0处理.如果用户需要较高的数据精度的话, 可以切换到下一级表示(例如half -> float, 或者float -> double).也可以选择通过手工计算来规约(无atomic), SP的精度还是较高的.

需要补充的是:本章节的例子写的不好, 里面使用了非标准的CPU端的原子操作扩展函数,也就是GCC的__sync_fetch_and_add(), (非标准扩展__sync*()系列),你应当考虑C++11起的标准原子操作,或者C11起的标准原子操作, 也就是atomic*()系列.以取得跨平台和跨系统时候的可移植性.
此外, 还需要补充的是:一些书或者网上有一些特别的技巧,例如可以通过atomicCAS, 来实现一些互斥锁之类的东西,请在使用前确保你已经充分了解了这些所谓的技巧, 我们不建议使用它们的.


回复 支持 反对

使用道具 举报

您需要登录后才可以回帖 登录 | 立即注册

本版积分规则

关闭

站长推荐上一条 /1 下一条

快速回复 返回顶部 返回列表