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

2024-04-19 04:04
文章标签 线程 编程 指令 cuda 洗牌

本文主要是介绍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执行调用束内线程索引的按位异或,内部指令可计算源束内线程索引。返回由源线程持有的值。该指令适合于蝴蝶寻址模式(a butterfly addressing pattern),如图所示。
在这里插入图片描述
洗牌函数还支持单精度浮点值。浮点洗牌函数采用浮点型的var参数,并返回一个浮点数。

线程束内的共享数据

跨线程束值的广播

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

线程束内上移

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

线程束内下移

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

线程束内环绕移动

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

跨线程束的蝴蝶交换

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

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

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

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

这篇关于CUDA编程---线程束洗牌指令的文章就介绍到这儿,希望我们推荐的文章对编程师们有所帮助!



http://www.chinasem.cn/article/916585

相关文章

Python包管理工具核心指令uvx举例详细解析

《Python包管理工具核心指令uvx举例详细解析》:本文主要介绍Python包管理工具核心指令uvx的相关资料,uvx是uv工具链中用于临时运行Python命令行工具的高效执行器,依托Rust实... 目录一、uvx 的定位与核心功能二、uvx 的典型应用场景三、uvx 与传统工具对比四、uvx 的技术实

Java中实现线程的创建和启动的方法

《Java中实现线程的创建和启动的方法》在Java中,实现线程的创建和启动是两个不同但紧密相关的概念,理解为什么要启动线程(调用start()方法)而非直接调用run()方法,是掌握多线程编程的关键,... 目录1. 线程的生命周期2. start() vs run() 的本质区别3. 为什么必须通过 st

Linux实现线程同步的多种方式汇总

《Linux实现线程同步的多种方式汇总》本文详细介绍了Linux下线程同步的多种方法,包括互斥锁、自旋锁、信号量以及它们的使用示例,通过这些同步机制,可以解决线程安全问题,防止资源竞争导致的错误,示例... 目录什么是线程同步?一、互斥锁(单人洗手间规则)适用场景:特点:二、条件变量(咖啡厅取餐系统)工作流

Java中常见队列举例详解(非线程安全)

《Java中常见队列举例详解(非线程安全)》队列用于模拟队列这种数据结构,队列通常是指先进先出的容器,:本文主要介绍Java中常见队列(非线程安全)的相关资料,文中通过代码介绍的非常详细,需要的朋... 目录一.队列定义 二.常见接口 三.常见实现类3.1 ArrayDeque3.1.1 实现原理3.1.2

SpringBoot3中使用虚拟线程的完整步骤

《SpringBoot3中使用虚拟线程的完整步骤》在SpringBoot3中使用Java21+的虚拟线程(VirtualThreads)可以显著提升I/O密集型应用的并发能力,这篇文章为大家介绍了详细... 目录1. 环境准备2. 配置虚拟线程方式一:全局启用虚拟线程(Tomcat/Jetty)方式二:异步

如何解决Druid线程池Cause:java.sql.SQLRecoverableException:IO错误:Socket read timed out的问题

《如何解决Druid线程池Cause:java.sql.SQLRecoverableException:IO错误:Socketreadtimedout的问题》:本文主要介绍解决Druid线程... 目录异常信息触发场景找到版本发布更新的说明从版本更新信息可以看到该默认逻辑已经去除总结异常信息触发场景复

Python 异步编程 asyncio简介及基本用法

《Python异步编程asyncio简介及基本用法》asyncio是Python的一个库,用于编写并发代码,使用协程、任务和Futures来处理I/O密集型和高延迟操作,本文给大家介绍Python... 目录1、asyncio是什么IO密集型任务特征2、怎么用1、基本用法2、关键字 async1、async

Java并发编程之如何优雅关闭钩子Shutdown Hook

《Java并发编程之如何优雅关闭钩子ShutdownHook》这篇文章主要为大家详细介绍了Java如何实现优雅关闭钩子ShutdownHook,文中的示例代码讲解详细,感兴趣的小伙伴可以跟随小编一起... 目录关闭钩子简介关闭钩子应用场景数据库连接实战演示使用关闭钩子的注意事项开源框架中的关闭钩子机制1.

JAVA保证HashMap线程安全的几种方式

《JAVA保证HashMap线程安全的几种方式》HashMap是线程不安全的,这意味着如果多个线程并发地访问和修改同一个HashMap实例,可能会导致数据不一致和其他线程安全问题,本文主要介绍了JAV... 目录1. 使用 Collections.synchronizedMap2. 使用 Concurren

shell编程之函数与数组的使用详解

《shell编程之函数与数组的使用详解》:本文主要介绍shell编程之函数与数组的使用,具有很好的参考价值,希望对大家有所帮助,如有错误或未考虑完全的地方,望不吝赐教... 目录shell函数函数的用法俩个数求和系统资源监控并报警函数函数变量的作用范围函数的参数递归函数shell数组获取数组的长度读取某下的