CUDA编程---线程束洗牌指令

从Kepler系列的GPU(计算能力为3.0或更高)开始,洗牌指令(shuffle instruction)作为一种机制被加入其中,只要两个线程在相同的线程束中,那么就允许这两个线程直接读取另一个线程的寄存器。
洗牌指令使得线程束中的线程彼此之间可以直接交换数据,而不是通过共享内存或全局内存来进行的。洗牌指令比共享内存有更低的延迟,并且该指令在执行数据交换时不消耗额外的内存。因此,洗牌指令为应用程序快速交换线程束中线程间的数据提供了一个有吸引力的方法。

束内线程

首先介绍一下束内线程lane)的概念。简单来说,一个束内线程指的是线程束内的单一线程。线程束中的每个束内线程是[0,31]范围内束内线程索引(lane index)的唯一标识。线程束中的每个线程都有一个唯一的束内线程索引,并且同一线程块中的多个线程可以有相同的束内线程索引(就像同一网格中的多个线程可以有相同的threadIdx.x值一样)。然而,束内线程索引没有内置变量。在一维线程块中,对于一个给定线程的束内线程索引线程束索引可以按以下公式进行计算:
在这里插入图片描述
例如,线程块中的线程1和线程33都有束内线程ID 1,但它们有不同的线程束ID。对于二维线程块,可以将二维线程坐标转换为一维线程索引,并应用前面的公式来确定束内线程和线程束的索引。

线程束洗牌指令的不同形式

有两组洗牌指令:一组用于整型变量,另一组用于浮点型变量。每组有4种形式的洗牌指令。在线程束内交换整型变量,其基本函数标记如下:
在这里插入图片描述
内部指令__shfl返回值是var,var通过由srcLane确定的同一线程束中的线程传递给__shfl。srcLane的含义变化取决于宽度值。这个函数能使线程束中的每个线程都可以直接从一个特定的线程中获取某个值。线程束内所有活跃的线程都同时产生此操作,这将导致每个线程中有4字节数据的移动。
变量width可被设置为2~32之间2的任意整数次幂(包括2和32),这是可选的。当设置为默认的warpSize(即32)时,洗牌指令跨整个线程束来执行,并且srcLane指定源线程的束内线程索引。然而,设置width允许将线程束细分为,使每段包含有width个线程,并且在每个段上执行独立的洗牌操作。对于不是32的其他width值,线程的束内线程ID和其在洗牌操作中的ID不一定相同。在这种情况下,一维线程块中的线程洗牌ID可以按以下公式进行计算:
在这里插入图片描述
例如,如果shfl被线程束中的每个线程通过以下参数调用:
在这里插入图片描述
那么线程0~15将从线程3接收x的值,线程16~31将从线程19接收x的值(在线程束的前16个线程中其偏移量为3)。为了简单起见,srcLane将被称为束内线程索引
__shfl指令从特定的束内线程到线程束中所有线程执行线程束广播操作,如下图所示:
在这里插入图片描述
洗牌操作的另一种形式是从与调用线程相关的线程中复制数据:
在这里插入图片描述
__shfl_up通过给调用的束内线程索引减去delta来计算源束内线程索引。返回由源线程所持有的值。因此,这一指令通过束内线程delta将var右移到线程束中。__shfl_up周围没有线程束,所以线程束中最低的delta个线程将保持不变,如图所示。
在这里插入图片描述
相反,洗牌指令的第三种形式是从相对于调用线程而言具有高索引值的线程中复制:
在这里插入图片描述
__shfl_down通过给调用的束内线程索引增加delta来计算源束内线程索引。返回由源线程持有的值。因此,该指令通过束内线程delta将var的值左移到线程束中。使用__shfl_down时周围没有线程束,所以线程束中最大的delta个束内线程将保持不变,如图所示。
在这里插入图片描述
洗牌指令的最后一种形式是根据调用束内线程索引自身的按位异或来传输束内线程中的数据:
在这里插入图片描述
通过使用laneMask执行调用束内线程索引的按位异或,内部指令可计算源束内线程索引。返回由源线程持有的值。该指令适合于蝴蝶寻址模式(Butterfly Addressing Pattern),如图所示。
在这里插入图片描述
洗牌函数还支持单精度浮点值。浮点洗牌函数采用浮点型的var参数,并返回一个浮点数。

线程束内的共享数据

跨线程束值的广播

下面的内核实现了线程束级的广播操作。每个线程都有一个寄存器变量value。源束内线程由变量srcLane指定,它等同于跨所有线程。每个线程都直接从源线程复制数据。
在这里插入图片描述
为了简单起见,使用有16个线程的一维线程块:
在这里插入图片描述
调用内核的方法如下。通过第三个参数test_shfl_broadcast将源束内线程设置为每个线程束内的第三个线程。
在这里插入图片描述
调用后的结果如下:
在这里插入图片描述

线程束内上移

下面的内核实现了洗牌上移的操作。线程束中每个线程的源束内线程都是独一无二的,并由它自身的线程索引减去delta来确定。
在这里插入图片描述
通过指定delta为2调用核函数:
在这里插入图片描述
其结果是,每个线程的值向右移动两个束内线程,结果如下所示。最左边的两个束内线程值保持不变
在这里插入图片描述

线程束内下移

下面的内核实现了下移操作。线程束中每个线程的源束内线程都是独一无二的,并由它自身的线程索引加上delta来确定。
在这里插入图片描述
通过指定delta为2调用核函数:
在这里插入图片描述
每个线程的值向左移动两个束内线程,结果如下所示。最右边的两个束内线程值保持不变。
在这里插入图片描述

线程束内环绕移动

下面的核函数实现了跨线程束的环绕移动操作。每个线程的源束内线程是不同的,并由它自身的束内线程索引加上偏移量来确定。偏移量可为正数也可为负数
在这里插入图片描述
通过指定一个正偏移量来调用内核,代码如下:
在这里插入图片描述
这个内核实现了环绕式左移操作,如下所示。不同于由test_shfl_down产生的结果,最右边的两个束内线程的值也变化了。
在这里插入图片描述

跨线程束的蝴蝶交换

下面的内核实现了两个线程之间的蝴蝶寻址模式,这是通过调用线程和线程掩码确定的。

调用掩码值为1的内核将导致相邻的线程交换它们的值
在这里插入图片描述
这个内核启动的输出如下:
在这里插入图片描述

使用线程束洗牌指令的并行归约

一个线程块中可能有几个线程束。对于线程束级归约来说,每个线程束执行自己的归约。每个线程不使用共享内存,而是使用寄存器存储一个从全局内存中读取的数据元素:
在这里插入图片描述
线程束级归约作为一个内联函数实现,如下所示:
在这里插入图片描述
在这个函数返回之后,每个线程束的总和保存到基于线程索引和线程束大小的共享内存中,如下所示:
在这里插入图片描述
对于线程块级归约,先同步块,然后使用相同的线程束归约函数将每个线程束的总和进行相加。之后,由线程块产生的最终输出由块中的第一个线程保存到全局内存中,如下所示:
在这里插入图片描述
对于网格级归约,g_odata被复制回到执行最终归约的主机中。下面是完整的reduceShfl核函数:
在这里插入图片描述

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/pingmian/1796.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

清华大学:序列推荐模型稳定性飙升,STDP框架惊艳登场

获取本文论文原文PDF,请公众号留言:论文解读 引言:在线平台推荐系统的挑战与机遇 在线平台已成为我们日常生活中不可或缺的一部分,它们提供了丰富多样的商品和服务。然而,如何为用户推荐感兴趣的项目仍然是一个挑战。…

【笔记】Telephony SIM SPN及运营商名称显示数据来源介绍

来源介绍 网络名称显示 来源及优先级(高到低) SourceCommentEnhanced Operator Name String(Eons) 名称信息存放: EF_PNN(PLMN Network Name, fid: 6FC5) :LAC和EF_PNN中的Record Identifier EF_OPL(Operator PLMN List, fid: 6FC…

67条tips实战案例渗透测试大佬的技巧总结

67条tips实战案例渗透测试大佬的技巧总结。 Tips 1. 手动端口探测 nmap的-sV可以探测出服务版本,但有些情况下必须手动探测去验证 使用Wireshark获取响应包未免大材小用,可通过nc简单判断 eg. 对于8001端口,nc连接上去,随便输…

GlobalFilter全局过滤器

这个跟跟刚才那个GatewatFilert默认全局配置的效果是一样的,但是那个是配置,只能使用已有的进行配置,GlobalFilter全局过滤器是通过类实现的 可以自己用代码实现拦截后要处理的逻辑。 定义方式: 先实现GlobalFilter接口&#xf…

深入C语言,发现多样的数据之枚举和联合体

一、枚举 枚举 是列出某些有穷序列集的所有成员的程序,或者是一种特定类型对象的计数。这两种类型经常(但不总是)重叠。是一个被命名的整型常数的集合。简单来说就将某种特定类型的对象一一进行列举,一一列举特定类型可能的取值。…

探索RadSystems:低代码开发的新选择(二)

系列文章目录 探索RadSystems:低代码开发的新选择(一)🚪 文章目录 系列文章目录前言一、RadSystems Studio是什么?二、用户认证三、系统角色许可四、用户记录管理五、时间戳记录总结 前言 在数字化时代,低…

【做一名健康的CSDNer】程序员哪几种行为最伤肾(程序员必看)

虽然没有专门针对程序员这一职业群体特有的伤肾行为的研究报道,但根据一般人群的健康风险和生活习惯,程序员由于其特殊的工作模式和环境,可能更容易出现如下伤肾的行为: 熬夜加班: 程序员由于项目进度、bug修复等原因&…

函数的创建和调用及删除

Oracle从入门到总裁:​​​​​​https://blog.csdn.net/weixin_67859959/article/details/135209645 函数和存储过程非常类似,也是可以存储在 Oracle 数据库中的 PL/SQL代码块,但是有返回值。 可以把经常使用的功能定义为一个函数,就像系统…

数仓建模—逻辑数据模型

数仓建模—逻辑数据模型 数据模型是数据元素及其基于现实世界对象之间的关系的可视化表示。数据模型揭示并定义数据在业务流程中的连接方式,并支持创建高效的信息系统或应用程序。例如,在商业智能中,数据模型定义用户可以在其分析中使用哪种数据。 逻辑数据模型 (LDM Logi…

【C++ STL序列容器】array 数组

文章目录 【 1. 基本原理 】【 2. array 的创建 】2.1 不赋初值2.2 赋默认值2.3 赋指定值 【 3. array 的成员函数 】实例 【 1. 基本原理 】 array 是在 C 普通数组的基础上添加了一些成员函数和全局函数。在使用上,它 比普通数组更 安全,且效率并没…

以太网帧格式解析

以太网的正式标准是IEEE802.3,它规定了以太网传输的帧结构。 以太网帧格式如下图所示: 以太网传输数据时,是按照上图的格式,自左到右依次传输的。需要注意的是前导码和SFD不属于以太网协议的内容,应该是属于物理层数据…

学习ArkTS -- 状态管理

装饰器 State 在声明式UI中,是以状态驱动试图更新: 状态(State):指驱动视图更新的数据(被装饰器标记的变量) 视图(View):基于UI描述渲染得到用户界面 说明…

病理验证mIF和TMA路线(自学)

目录 技术 使用配对病理切片 mIF验证 单基因使用TMA验证 技术 多重荧光免疫组化技术 (Multiplex immunohistochemical,mIHC) 也称作酪氨酸信号放大 (Tyramide dignal amplification,TSA) 技术,是一类利用辣根过氧化酶 (Horseradish Pero…

weblogic反序列化漏洞(CVE-2017-10271)复现

直接用vuluhub搭建现成的靶场做 访问靶场 打开是这样表示成功 想反弹shell 就先开启kali1的nc监听,这就监听2233端口吧 linux:nc -l -p 2233 抓包修改为攻击数据包 ip和端口可以任意修改 反弹的shell 还可以写入文件shell 只需要把提供的poc POS…

修复Windows搜索不工作的几种方法,总有一种适合你

序言 Windows搜索是Windows 10中一个非常有用的功能,它允许你搜索特定的程序、应用程序、文档、图片、文件、设置等,以便快速访问它们。但有时Windows搜索找不到我们预期的结果,甚至没有响应,这会给Windows用户带来很多不便。如果Windows 10中的搜索栏不工作,该怎么办?你…

【MySQL】SQL优化

SQL优化 插入数据 insert 一次插入数据和批量插入数据 insert into tb_test (id, name) values (1,Tom); insert into tb_test (id, name) values (1,Tom),(2,Jack),(3,Jerry);优化方案: 手动控制事务,且按主键顺序插入。start transaction; insert …

机器人实验室LAAS-CNRS介绍

一、LAAS-CNRS介绍 1、缩写介绍 同样的,给出英文缩写的全称,以便理解。这里的LAAS(Laboratory for Analysis and Architecture of Systems)指法国的系统分析与架构实验室,CNRS(Centre National de la Rec…

OpenHarmony实战开发-提升应用响应速度。

应用对用户的输入需要快速反馈,以提升交互体验,因此本文提供了以下方法来提升应用响应速度。 避免主线程被非UI任务阻塞减少组件刷新的数量 避免主线程被非UI任务阻塞 在应用响应用户输入期间,应用主线程应尽可能只执行UI任务(…

【电控笔记6.3】采样-Z转换-零阶保持器

本质 数字转模拟:零阶保持器 采样 z-1所描述的物理意义即为延迟T时间的拉氏转换e-sT 信号采样延时

Python --- 新手小白自己动手安装Anaconda+Jupyter Notebook全记录(Windows平台)

新手小白自己动手安装AnacondaJupyter Notebook全记录 这两天在家学Pythonmathine learning,在我刚刚入手python的时候,我写了一篇新手的入手文章,是基于Vs code编译器的入手指南,里面包括如何安装python,以及如何在Vs…