DAY51:阅读Warp Shuffle Functions

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第51天,我们正在讲解CUDA C语法,希望在接下来的49天里,您可以学习到原汁原味的CUDA,同时能养成英文阅读的习惯。

B.15. Warp Shuffle Functions

__shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync exchange a variable between threads within a warp.

Supported by devices of compute capability 3.x or higher.

Deprecation Notice: __shfl, __shfl_up, __shfl_down, and __shfl_xor have been deprecated as of CUDA 9.0.

B.15.1. Synopsis

T can be int, unsigned int, long, unsigned long, long long, unsigned long long, float or double. With the cuda_fp16.h header included, T can also be __half or __half2.

B.15.2. Description

The __shfl_sync() intrinsics permit exchanging of a variable between threads within a warp without use of shared memory. The exchange occurs simultaneously for all active threads within the warp (and named in mask), moving 4 or 8 bytes of data per thread depending on the type.

Threads within a warp are referred to as lanes, and may have an index between 0 and warpSize-1 (inclusive). Four source-lane addressing modes are supported:

  • __shfl_sync()
  • Direct copy from indexed lane
  • __shfl_up_sync()
  • Copy from a lane with lower ID relative to caller
  • __shfl_down_sync()
  • Copy from a lane with higher ID relative to caller
  • __shfl_xor_sync()
  • Copy from a lane based on bitwise XOR of own lane ID

Threads may only read data from another thread which is actively participating in the __shfl_sync() command. If the target thread is inactive, the retrieved value is undefined.

All of the __shfl_sync() intrinsics take an optional width parameter which alters the behavior of the intrinsic. width must have a value which is a power of 2; results are undefined if width is not a power of 2, or is a number greater than warpSize.

__shfl_sync() returns the value of var held by the thread whose ID is given by srcLane. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. If srcLane is outside the range [0:width-1], the value returned corresponds to the value of var held by the srcLane modulo width (i.e. within the same subsection).

__shfl_up_sync() calculates a source lane ID by subtracting delta from the caller's lane ID. The value of var held by the resulting lane ID is returned: in effect, var is shifted up the warp bydelta lanes. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. The source lane index will not wrap around the value of width, so effectively the lower delta lanes will be unchanged.

__shfl_down_sync() calculates a source lane ID by adding delta to the caller's lane ID. The value of var held by the resulting lane ID is returned: this has the effect of shifting var down the warp by delta lanes. If width is less than warpSize then each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. As for __shfl_up_sync(), the ID number of the source lane will not wrap around the value of width and so the upper delta lanes will remain unchanged.

__shfl_xor_sync() calculates a source line ID by performing a bitwise XOR of the caller's lane ID with laneMask: the value of var held by the resulting lane ID is returned. If width is less than warpSize then each group of width consecutive threads are able to access elements from earlier groups of threads, however if they attempt to access elements from later groups of threads their own value of var will be returned. This mode implements a butterfly addressing pattern such as is used in tree reduction and broadcast.

The new *_sync shfl intrinsics take in a mask indicating the threads participating in the call. A bit, representing the thread's lane id, must be set for each participating thread to ensure they are properly converged before the intrinsic is executed by the hardware. All non-exited threads named in mask must execute the same intrinsic with the same mask, or the result is undefined.

B.15.3. Return Value

All __shfl_sync() intrinsics return the 4-byte word referenced by var from the source lane ID as an unsigned integer. If the source lane ID is out of range or the source thread has exited, the calling thread's own var is returned.

B.15.4. Notes

Threads may only read data from another thread which is actively participating in the __shfl_sync() command. If the target thread is inactive, the retrieved value is undefined.

width must be a power-of-2 (i.e., 2, 4, 8, 16 or 32). Results are unspecified for other values.

B.15.5. Examples

B.15.5.1. Broadcast of a single value across a warp

B.15.5.2. Inclusive plus-scan across sub-partitions of 8 threads

B.15.5.3. Reduction across a warp

本文备注/经验分享:

今天的主要内容是warp shuffle, 这是一个从计算能力3.0(Kepler)开始提供的特性.如同之前的warp vote一样, 同样可以进行warp内部的数据交换.但有几个区别: (1)没有数据的规约处理功能. (2)交换的数据是32-bit的(4B), 比warp vote的1-bit要大很多. 较新版本的CUDA(例如CUDA 9+), 具有增强版本的warp shuffle功能,例如可以交换64-bit的数据, 此时编译器将自动为你拆分成2个32-bit的shfl指令,但并不排除将来的硬件, 会直接实现64-bit的shfl版本.用户也可以无视本章节, 因为和warp vote一样, 这不是一个必选的特性.用户可以完全无视warp shuffle, 而写出功能完备的代码.不使用warp shuffle的时候, 需要通过shared memory进行数据交换. 后者这种交换不如warp shuffle高效.换句话说, warp shuffle有如下特色(相比shared memory上的交换): (1)不需要为参与数据交换的warp(和内部的线程们--lanes)分配shared memory.这样可以减少shared memory的使用. 甚至有些代码, 使用满了48KB的每个block最大的shared memory, 此时无法继续分配空间进行数据交换使用.但依然可以进行warp shuffle.(请注意, 你也可以临时将shared memory中的内容交换到寄存器, 然后将空出来的shared memory用作数据交换. 然后交换完成后, 再将被破坏掉的shared memory中的内容保存回去. 但这样比较低效)。 所以完全不占用空间, 相比shared memory上的方式, 是一个很大的特色. (2)不需要计算地址, 用户可以直接按值目标线程的位置之类的信息(例如, lane id, lane id是warp中的分量线程的编号, 0-31一般)进行交换. 而传统的shared memory上的方式则需要先计算写入地址(位置, 例如线程编号 *4 + 基地址, 如果要每个线程交换4B大小的话), 然后写入,写入完成后需要继续计算读取地址(同上),然后读取,这样的话需要较高的SP来进行辅助. 而shfl可以直接计算常见类型的地址, 等于免费了(例如向上shuffle 1个4B位置的时候, 该计算可以免费)。很多时候, 访问shared memory密集的代码并没有卡在shared memory上(可以通过profiler看), 而是卡在地址计算上. 这有的时候非常尴尬.而warp shuffle不仅仅节省了空间, 还节省了计算量. (3)点则是, shuffle的本质依然是使用shared memory, 它被编译后, 生成的指令被GPU执行的时候, 依然是提交一条指令给shared memory,只是shared memory不进行任何操作, 就地将数据按照一定的方式打乱(shuffle么. 类似洗牌), 然后直接返回.相比普通的shared memory上的1次写入+1次读取, shuffle可以一条指令内直接完成,这样就算SP上的地址计算不是一个瓶颈, 当shared memory成为瓶颈的时候, 通过shuffle可以等效的提高一倍效率. (4)点则是, 你现在不需要同步了. 以前通过shared memory上的交换, 除非使用locked-step的warp之类的 + volatile关键字之类的技巧, 你需要block同步的, 这样的代价就很大.而warp shuffle可以直接交换, 不需要block级别的同步, 也没有之前的技巧容易导致的BUG(例如用户忘记了volatile), 非常易用.正因为至少有这4方面的好处, 你能用warp shuffle的地址就应当使用. 除非你用不了(例如, 数据的交换需要在block大范围内, 而不是warp范围内, 则此时你只能走shared memory,这也是我们之前说过的shared memory的三大用途之一)。 回到具体的CUDA 9+, warp shuffle从这个版本起,引入了不兼容的改变(多了_sync后缀和需要warp内部同步的线程的掩码, 因为从计算能力7.0起, warp不一定必须完全步伐一致的执行). 对于新的CUDA 9.2 + 老卡(例如GTX1080), 用户可以指定0xffff之类的mask,(为了配合新硬件, 老硬件总会慢慢老去并从市面上消失的, 将来只会有新卡)此时将可以模拟原本的无后缀, 无第一个参数的老shuffle系列函数行为.或者用户依然可以维持老用法, 但会得到一个警告, 建议用户尽快迁移. 本章节所提供的shfl系列函数, 具有4个主要版本. 分别是up/down, 即目标线程的编号可以通过简单的加减法来确定的时候.还有一种是叫xor, 这个是将线程编号进行异或操作。另外一个则是直接索引(这个是没有后缀的那个), 用户可以直接要求指定一个线程编号.这4种. 其中的xor版本需要注意一下, 这也叫buffterfly型操作(蝴蝶),因为蝴蝶的翅膀是类似X形状交错的。网上有一些教程, 对xor版本里面的线程编号变化情况有图,你会看到真的很像蝴蝶. 像是很多操作, 例如前序和,像是常见的一些规约操作,都可以全部或者部分通过shuffle来进行.手册上的一些例子, 包括本章节最后的:那个规约求和的例子, 可以适合用在一个长的block内部的规约的最后(当缩减到warp规模的时候).实际上手册之前这个例子有个对比的.是一个通过shared memory上的规约,后者版本需要多次的shared memory的读写.而warp shuffle版本没有这个要求.

继续补充一点:以前很多代码需要 #if __CUDA_ARCH < 300 进行shared memory上的空间分配, 以及, 后续的shared memory上的交换 #else 直接warp shuffle #endif

现在这种代码已经基本上不需要了. 这是为了Kepler之前的卡进行的. 现在从CUDA 9开始, 已经不再支持对这些卡进行开发了(但可以运行), 用户看一些书的时候, 应当直接看#else的部分, 而不需要再考虑进行前者了. 此外, 还需要注意的是, 从计算能力3.0(最初支持shuffle的版本)到计算能力7.2, warp shuffle的速率均是32条/SM/clock, 但是因为实际上SM里面的SP数量是变化的. 例如从Kepler的192个, 到Maxwell/Pascal(6.1)的128个, 到Pascal(6.0)和GV100(7.0)的64个SP, 实际上warp shuffle的等效速率是从1/6到1/4到1/2发生变化的,换句话说, 越来越快了。用户应当考虑在不同的硬件上, 进行warp shuffle, 和就地直接重新计算数据(如果能通过直接计算算出来的话), 两者之间的代价权衡. 此外, 关于32这点, 也可以看出的确是走shared memory的,(shared memory这些年一直是32个banks,之前的32是按线程单位提供的---请注意profiler的指令单位是warp单位, 也就是1条/SM/周期),该指令具有典型的shared memory的周期,当用户的代码卡shared memory操作,或者对延迟非常敏感而无法掩盖的时候, 应当注意shuffle本身的延迟.(比计算指令高不少, 而且可变)但是一般来说, 除非特殊情况, 使用shuffle总是有正面效果的. 大致这样.几个例子上的应用, 用户可以看一下. 都比较简单。

这里需要说一下. 所有的本章节的shfl函数都接受一个width参数, 可以在warp进步进一步的切分大小. 很多时候, 例如在邻近的16个线程或者4个线程之类的场合, 直接指定大小可以等效的减轻用户计算线程ID的计算量.此时则应当考虑使用. 此外, 为了未来着想, 用户当应当只需要warp内部的较小的范围的时候, 应当直接指定较小的范围, 例如4, 8, 16这种. 未来的一些硬件可能对这些有特殊的优化. 例如A家的一些硬件, 当在特定的sub-wave的范围内进行交换的时候, 可以通过某些特殊的设计, 直接将额外的指令执行消除, 同时减少了指令量, 也消除了延迟. 例如连续的16个wave内的某些交换操作, 可以直接不走LDS(等于N卡的shared memory), 直接0成本.虽然目前N卡没有对特定的范围情况进行优化处理。 但用户能这样写就应当保持这样写, 以取得未来的更好效果.此外, 几乎如同一个惯例,虽然A家的硬件如此优秀, 不仅仅提供了wave内的shuffle功能, 还提供了特殊情况下的硬件优化,但很遗憾的是, 无论是这种优化, 还是普通的shuffle功能, 至今依然不能使用,AMD至今没有将它们导出到OpenCL C种.因此虽然N卡的warp shuffle比AMD的wave shuffle更往后推出,但上来就可以用的。 选择N卡, 的确是你明智的选择. 小声说一句: CUDA C版本的warp shuffle虽然已经很强了, 但PTX版本的功能更强.PTX版本是双返回值的.除了CUDA C这里能返回交换后的数据外,还能返回是否真的参与了交换(例如因为越界),需要额外功能的shuffle支持, 应当考虑PTX嵌入.

有不明白的地方,请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-07-18

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏我是攻城师

如何使用neo4j存储树形无限级菜单

50060
来自专栏王亚昌的专栏

Go语言入门之路

Go is designed from first principles to advance the practice of software enginee...

38450
来自专栏玩转全栈

flutter使用platform-channels制作插件

一、flutter使用platform-channels制作插件是否是一种完美的体验?

51040
来自专栏GreenLeaves

EF基础知识小记一

1、EF等ORM解决方案出现的原因 因为软件开发中分析和解决问题的方法已经接近成熟,然后关系型数据库却没有,很多年来,数据依然是保存在表行列这样的模式里,所以,...

18290
来自专栏云飞学编程

Python爬虫的简单实现!用python爬虫自己做天气预报查询

最近小编在学习爬虫,就想找个东西练练手,小说、图片、音乐什么的都烂大街了,正好最近天气是越来越冷,小编窝家里自己敲了个天气简单查询的代码,请大家指正下!

13310
来自专栏AI派

早上起床后不想动,让 Python 来帮你朗读网页吧

之所以用 Python,就是因为 Python 有着丰富的库,网页正文识别也不在话下。这里我尝试了 readability、goose3。

20620
来自专栏木子昭的博客

Python3好用的原生api

对列表进行反序是一个很常见的操作, 但python反向切片的玩法实在是非常简洁, 让人无法拒绝, 其实对某一数据结构进行"反向"是一个很有意...

9310
来自专栏java达人

Java 理论与实践: JDK 5.0 中更灵活、更具可伸缩性的锁定机制

多线程和并发性并不是什么新内容,但是Java 语言设计中的创新之一就是,它是第一个直接把跨平台线程模型和正规的内存模型集成到语言中的主流语言。核心类库包含一个T...

20860
来自专栏linxu shell指南

软件构件、中间件、面向对象

    1、构件定义:组(构)件是软件系统可替换的、物理的组成部分,它封装了实现体(实现某个职能)并提供了一组接口的实现方法。可以认为组件是一个封装的代码模块或...

29340
来自专栏Golang语言社区

从Baa开发中总结Go语言性能渐进优化

在Go生态已经有很多WEB框架,但感觉没有一个符合我们的想法,我们想要一个简洁高效的核心框架,提供路由,context,中间件和依赖注入,而且拒绝使用正则和反射...

57380

扫码关注云+社区

领取腾讯云代金券