cuda线程束原语 __shfl_xor、__shfl、__shfl_up()、__shfl_down()

最新推荐文章于 2024-08-20 14:45:21 发布
最新推荐文章于 2024-08-20 14:45:21 发布
阅读量9.5k 收藏 18
点赞数 8
分类专栏： cuda
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/jqw11/article/details/103071556
版权

在CC3.0以上，支持了shuffle指令，允许thread直接读其他thread的寄存器值，只要两个thread在 同一个warp中，这种比通过shared Memory进行thread间的通讯效果更好，latency更低，同时也不消耗额外的内存资源来执行数据交换。

这里介绍warp中的一个概念lane，一个lane就是一个warp中的一个thread，每个lane在同一个warp中由lane索引唯一确定，因此其范围为[0,31]。在一个一维的block中，可以通过下面两个公式计算索引：

laneID = threadIdx.x % 32

warpID = threadIdx.x / 32

例如，在同一个block中的thread1和33拥有相同的lane索引1。

Variants of the Warp Shuffle Instruction

有两种设置shuffle的指令：一种针对整型变量，另一种针对浮点型变量。每种设置都包含四种shuffle指令变量。为了交换整型变量，使用过如下函数：

参考书籍：《cuda专家手册|GPU编程权威》

1：_shfl_xor

首先介绍__shfl_xor，因为最先用到它。

__shfl_xor（var,laneMask）：Copy from a lane based on bitwise XOR of own lane ID

意思就是从当前的线程id与laneMak异或运算的值作为线程号的，把这个线程号的var值取出来。

演示图：

举例：

tid =0

laneMask =16

tid xor laneMask(0000 xor 1000)=0111=15

所有取到的值为15号线程的var

那我们看下完成测试代码：

__global__ void test_shfl_xor(int A[], int B[])
{
    int tid = threadIdx.x;
    int best = B[tid];
    //best = subgroup_min<32>(best, 0xffffffffu);
    best = __shfl_xor(best, 8);
    A[tid] = best;
}

int main()
{
  
    int *A,*Ad, *B, *Bd;
    int n = 32;
    int size = n * sizeof(int);

    // CPU端分配内存
    A = (int*)malloc(size);
    B = (int*)malloc(size);

    for (int i = 0; i < n; i++)
    {
      
      B[i] = rand()%101;
      std::cout << B[i] << std::endl;
    }
   
    std::cout <<"----------------------------" << std::endl;
   
    // GPU端分配内存
    cudaMalloc((void**)&Ad, size);
    cudaMalloc((void**)&Bd, size);
    cudaMemcpy(Bd, B, size, cudaMemcpyHostToDevice);
  

    // 定义kernel执行配置，（1024*1024/512）个block，每个block里面有512个线程
    dim3 dimBlock(128);
    dim3 dimGrid(1000);

    // 执行kernel
    const auto t1 = std::chrono::system_clock::now();

    test__shfl_xor << <1, 32 >> > (Ad,Bd);
   
    cudaMemcpy(A, Ad, size, cudaMemcpyDeviceToHost);

    // 校验误差
    float max_error = 0.0;
    for (int i = 0; i <     32; i++)
    {
       
            std::cout << A[i] << std::endl;
    }

   
    // 释放CPU端、GPU端的内存
    free(A);    
    cudaFree(Ad);
    free(B);
    cudaFree(Bd);  
    return 0;
}

运行结果：

41 85 72 38 80 69 65 68 96 22 49 67 51 61 63 87 66 24 80 83 71 60 64 52 90 60 49 31 23 99 94 11
----------------------------
66 24 80 83 71 60 64 52 90 60 49 31 23 99 94 11 41 85 72 38 80 69 65 68 96 22 49 67 51 61 63 87

--------------------------------

再高阶的用法，求取线程束中最大值：

template <typename T, unsigned int GROUP_SIZE, unsigned int STEP>
struct subgroup_min_impl {
    static __device__ T call(T x, uint32_t mask) {
#if CUDA_VERSION >= 9000
        x = min(x, __shfl_xor_sync(mask, x, STEP / 2, GROUP_SIZE));
#else
        x = min(x, __shfl_xor(x, STEP / 2, GROUP_SIZE));
#endif
        return subgroup_min_impl<T, GROUP_SIZE, STEP / 2>::call(x, mask);
    }
};
template <typename T, unsigned int GROUP_SIZE>
struct subgroup_min_impl<T, GROUP_SIZE, 1u> {
    static __device__ T call(T x, uint32_t) {
        return x;
    }
};


template <unsigned int GROUP_SIZE, typename T>
__device__ inline T subgroup_min(T x, uint32_t mask) {
    return subgroup_min_impl<T, GROUP_SIZE, GROUP_SIZE>::call(x, mask);
}




__global__ void test__shfl_xor(int A[], int B[])
{
    int tid = threadIdx.x;
    int best = B[tid];
    best = subgroup_min<32>(best, 0xffffffffu);
    //best = __shfl_xor(best, 16);
    A[tid] = best;
}

int main()
{
    int *A,*Ad, *B, *Bd;
    int n = 32;
    int size = n * sizeof(int);

    // CPU端分配内存
    A = (int*)malloc(size);
    B = (int*)malloc(size);

    for (int i = 0; i < n; i++)
   {   
      B[i] = rand()%101;
      std::cout << B[i] << std::endl;
   }
   
    std::cout <<"----------------------------" << std::endl;
   
    // GPU端分配内存
    cudaMalloc((void**)&Ad, size);
    cudaMalloc((void**)&Bd, size);
    cudaMemcpy(Bd, B, size, cudaMemcpyHostToDevice); 

    // 定义kernel执行配置，（1024*1024/512）个block，每个block里面有512个线程
    dim3 dimBlock(128);
    dim3 dimGrid(1000);

    // 执行kernel
    const auto t1 = std::chrono::system_clock::now();

    test_shfl_xor << <1, 32 >> > (Ad,Bd);
   
    cudaMemcpy(A, Ad, size, cudaMemcpyDeviceToHost);

    // 校验误差
    float max_error = 0.0;
    for (int i = 0; i <     32; i++)
    {
       
            std::cout << A[i] << std::endl;
    }

    cout << "max error is " << max_error << endl;

    // 释放CPU端、GPU端的内存
    free(A);
    free(B);   
    cudaFree(Ad);
    cudaFree(Bd);
 
    return 0;
}

运行结果

41 85 72 38 80 69 65 68 96 22 49 67 51 61 63 87 66 24 80 83 71 60 64 52 90 60 49 31 23 99 94 11
----------------------------
11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11 11

_shfl_xor介绍完毕

--------------------------------------

2.__shfl()

 

Direct copy from indexed lane：复制lane id数据

__shfl(int var,int srclane,int width =32)

 

这个就是比较简单，咱们直接上代码：

__global__ void test_shfl(int A[], int B[])
{
    int tid = threadIdx.x;
    int best = B[tid];
   
    best = __shfl(best, 3);
    A[tid] = best;
}

int main()
{
    int *A,*Ad, *B, *Bd;
    int n = 32;
    int size = n * sizeof(int);

    // CPU端分配内存
    A = (int*)malloc(size);
    B = (int*)malloc(size);

    for (int i = 0; i < n; i++)
   {   
      B[i] = rand()%101;
      std::cout << B[i] << std::endl;
   }
   
    std::cout <<"----------------------------" << std::endl;
   
    // GPU端分配内存
    cudaMalloc((void**)&Ad, size);
    cudaMalloc((void**)&Bd, size);
    cudaMemcpy(Bd, B, size, cudaMemcpyHostToDevice); 

    // 定义kernel执行配置，（1024*1024/512）个block，每个block里面有512个线程
    dim3 dimBlock(128);
    dim3 dimGrid(1000);

    // 执行kernel
    const auto t1 = std::chrono::system_clock::now();

    test_shfl << <1, 32 >> > (Ad,Bd);
   
    cudaMemcpy(A, Ad, size, cudaMemcpyDeviceToHost);

    // 校验误差
    float max_error = 0.0;
    for (int i = 0; i <     32; i++)
    {
       
            std::cout << A[i] << std::endl;
    }

    cout << "max error is " << max_error << endl;

    // 释放CPU端、GPU端的内存
    free(A);
    free(B);   
    cudaFree(Ad);
    cudaFree(Bd);
 
    return 0;
}

按以上代码逻辑，取得数据全是第3号线程的数：

运行结果：

41 85 72 38 80 69 65 68 96 22 49 67 51 61 63 87 66 24 80 83 71 60 64 52 90 60 49 31 23 99 94 11
----------------------------
38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38 38

 

--------------------------------------------------------------------------------------------------------------------------------

3.__shfl_up()

__shfl_up(int var,unsigned int delta,int width =32)：Copy from a lane with lower ID relative to caller

 

把tid-delta的线程好的var复制给tid的 var，如果tid-delta<0，var保持原来的值

见代码：

__global__ void test_shfl_up(int A[], int B[])
{
    int tid = threadIdx.x;
    int best = B[tid];

    best = __shfl_up(best, 3);
    A[tid] = best;
}

运行结果：

41 85 72 38 80 69 65 68 96 22 49 67 51 61 63 87 66 24 80 83 71 60 64 52 90 60 49 31 23 99 94 11
----------------------------
41 85 72 41 85 72 38 80 69 65 68 96 22 49 67 51 61 63 87 66 24 80 83 71 60 64 52 90 60 49 31 23

--------------------------------------------------------------------------------------

4.__shfl_down

__shfl_down(int var,unsigned int delta,int width =32)

把tid+delta的线程好的var复制给tid的 var，如果tid+delta>32，var保持原来的值

测试代码:

__global__ void test_shfl_down(int A[], int B[])
{
    int tid = threadIdx.x;
    int best = B[tid];

    best = __shfl_down(best, 3);
    A[tid] = best;
}

运行结果：

41 85 72 38 80 69 65 68 96 22 49 67 51 61 63 87 66 24 80 83 71 60 64 52 90 60 49 31 23 99 94 11
----------------------------
38 80 69 65 68 96 22 49 67 51 61 63 87 66 24 80 83 71 60 64 52 90 60 49 31 23 99 94 11 99 94 11

 

8BitCat
关注
专栏目录
使用CUDA Warp-Level级原语
吴建明wujianming_110117
01-29 1560
使用CUDA Warp-Level级原语 NVIDIA GPU以SIMT（单指令，多线程）的方式执行称为warps 的线程组。许多CUDA程序通过利用warp执行来实现高性能。本文将展示如何使用cuda9中引入的原语，使warp级编程安全有效。 Figure 1: The Tesla V100 Accelerator with Volta GV100 GPU. SXM2 Form Factor. 图1:Volta GV100 GPU的特斯拉V100加速器。SXM2外形尺寸。 Warp-level Prim
cuda-example.rar_cuda examples_cuda 例子_cuda学习
09-19
5. **同步与同步原语**：在CUDA编程中，线程间的同步是必要的，以确保数据的一致性。可以使用`__syncthreads()`函数来同步一个线程块内的所有线程，或者使用流（Stream）进行异步计算和内存操作。 6. **CUDA错误...
cuda __shfl_xor
gonaYet的博客
06-04 5363
__shfl是cuda提供的wrap(线程束)级别的方法,一般线程束为32。相似的方法有几个。 __shfl() Direct copy from indexed lane __shfl_up() Copy from a lane with lower ID relative to caller __shfl_down() Copy from a lane with higher ID re
11.理解线程束执行的本质（一）
最新发布
sunhjb的博客
08-20 910
前面已经大概的介绍了CUDA执行模型的大概过程，硬件的大概结构，例如SM的大概结构，而对于硬件来说，CUDA执行的实质是线程束的执行。从软件的角度看，cuda执行中似乎所有的线程都是并行地运行的。在逻辑上这是正确的，但从硬件的角度来看，实际上硬件资源是有限的，不是所有线程在物理上都可以同时并行地执行。现在从硬件的角度来介绍，了解执行的硬件过程是有必要的。
cuda warp内原语——vote instructions&warp instructions
a12638915的博客
01-07 802
warp内原语 ——shfl_sync, shfl_up_sync, shfl_down_sync, and shfl_xor_sync 可在warp级别交换线程的变量。 写在开头： **·**shfl, shfl_up, shfl_down, and __shfl_xor在cuda9.0以上的设备已经被禁止使用，所以最新的cuda版本都需要使用sync版本 **·**内存延迟：global memory ~500 cycles /shared memory ~5 cycles/registers ~1 c
最简示例 简介洗牌函数 之 __shfl_up_sync() cuda 之 shuffle
eloudy的专栏
02-07 1480
__shfl_up_sync() up 是指 从取值的线程对齐。 #include <cuda_runtime.h> #include <stdio.h> #define warpSize 32 __global__ void scan4(float* a, float* b) { int laneId = threadIdx.x & 0x1f; float value; value = a[laneId]; value
xor
daojin505的专栏
08-07 682
x ^ y == (~x & y) | (x & ~y)
CUDA.rar_cuda_cuda 例子
09-20
这个“CUDA.rar_cuda_cuda 例子”压缩包提供了一些关于CUDA编程的基础教程和实例，帮助我们更深入地理解和掌握CUDA编程的核心概念和技术。 《深入浅出谈CUDA.pdf》这本书或者文档，可能是压缩包的主要内容，它将...
bank_service.rar_C线程模拟排队_界面_界面 多线程_银行 排队_银行排队
09-21
这可以通过线程同步原语（如互斥锁、信号量或条件变量）来实现，以确保VIP线程能够优先访问资源。 多线程在银行排队系统中的应用可以提高服务效率，因为不同的线程可以同时处理多个客户。线程之间的通信和协作是...
Posix.rar_POSIX多线程_posix_posix_queue
09-20
在IT行业中，多线程编程是提升程序性能和并发能力的重要技术之一，特别是在POSIX（Portable Operating System Interface）这样的开放标准平台上。POSIX是为操作系统提供接口的一系列标准，它定义了多线程编程的规范...
XOR
ACdream
06-08 1447
https://www.lijinma.com/blog/2014/05/29/amazing-xor/
GPU编程 CUDA C++ 分别使用“线程束内同步函数 __syncwarp( )“、“线程束内洗牌函数 __shfl_down_sync( )”、“线程协作组” 三种方式进行数组归约计算
行子的博客
04-02 825
【代码】GPU编程 CUDA C++ 分别使用“线程束内同步函数 __syncwarp( )“、“线程束内洗牌函数 __shfl_down_sync( )”、“线程协作组” 三种方式进行数组归约计算。
CUDA中的Warp Shuffle
专注于人工智能领域的小何尚
06-27 6870
在 warp 内的线程之间交换变量。由计算能力 3.x 或更高版本的设备支持。弃用通知： 在 CUDA 9.0 中已针对所有设备弃用。删除通知：当面向具有 7.x 或更高计算能力的设备时， 不再可用，而应使用它们的同步变体。 可以是 。 包含 头文件后， 也可以是 。 同样，包含 cuda_bf16.h 头文件后，T 也可以是 。 内在函数允许在 warp 内的线程之间交换变量，而无需使用共享内存。 交换同时发生在 warp 中的所有活动线程（并以命名），根据类型移动每个线程 4 或 8 个字节的数据。
CUDA编程---线程束洗牌指令
qq_40268672的博客
04-18 1388
从Kepler系列的GPU（计算能力为3.0或更高）开始，（shuffle instruction）作为一种机制被加入其中，只要两个线程在相同的线程束中，那么就允许这两个线程直接读取另一个线程的寄存器。。洗牌指令比共享内存有更低的延迟，并且该指令在执行数据交换时不消耗额外的内存。因此，洗牌指令为应用程序快速交换线程束中线程间的数据提供了一个有吸引力的方法。
线程间怎么交换数据_在LLVM后端实现跨通道数据搬移
weixin_39536427的博客
12-04 561
作者：汪岩AMD GPU的每个CU有一个64kB的存储空间，称为本地数据共享（Local Data Share，LDS），用于同一计算单元中的work group内各个work item之间的低延迟通信和数据共享。LDS配置为32个bank，每个bank有512个4字节的条目（entry）。尽管有LDS，大多数实际的计算指令仍对寄存器中的数据进行操作。从峰值性能的角度看，global memory...
CUDA10.0官方文档的翻译与学习之CUDA合作组
四国棋客
01-09 1989
目录 背景 介绍 块内组 线程组和线程块 分片（Tiled Partitions） 线程块分片（Thread Block Tiles） 伪线程混洗函数 伪线程投票函数 伪线程匹配函数 合并组 块内合作组的使用 发现模板 伪线程同步代码模板 组合 网格同步 多设备同步 结语 背景 今天我们翻译一下CUDA10.0官方文档中最后一个值得我们注意的部分——合作组 介绍 合作组(Cooperative Groups)是在CUDA 9中为了组织通信线程组而引进的对CUDA编程
【CUDA 基础】5.6 线程束洗牌指令
Tony的博客
06-26 1074
Abstract: 本文介绍线程束洗牌指令的用法 Keywords: 线程束洗牌指令 开篇废话 赶紧写博客，少说废话。 前面介绍了共享内存，常量内存，只读内存的使用，今天我们来研究一个比较特殊的机制，名字也很特殊，叫做线程束洗牌指令。 支持线程束洗牌指令的设备最低也要3.0以上， 洗牌指令，shuffle instruction作用在线程束内，允许两个线程见相互访问对方的寄存器...
评论 4
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值