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

最新推荐文章于 2024-08-20 14:45:21 发布
最新推荐文章于 2024-08-20 14:45:21 发布
阅读量9.3k 收藏 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.rar_cuda_cuda 例子
09-20
这个“CUDA.rar_cuda_cuda 例子”压缩包提供了一些关于CUDA编程的基础教程和实例，帮助我们更深入地理解和掌握CUDA编程的核心概念和技术。 《深入浅出谈CUDA.pdf》这本书或者文档，可能是压缩包的主要内容，它将...
GPU编程 CUDA C++ 分别使用“线程束内同步函数 __syncwarp( )“、“线程束内洗牌函数 __shfl_down_sync( )”、“线程协作组” 三种方式进行数组归约计算
行子的博客
04-02 801
【代码】GPU编程 CUDA C++ 分别使用“线程束内同步函数 __syncwarp( )“、“线程束内洗牌函数 __shfl_down_sync( )”、“线程协作组” 三种方式进行数组归约计算。
CUDA中的Warp Shuffle
专注于人工智能领域的小何尚
06-27 6427
在 warp 内的线程之间交换变量。由计算能力 3.x 或更高版本的设备支持。弃用通知： 在 CUDA 9.0 中已针对所有设备弃用。删除通知：当面向具有 7.x 或更高计算能力的设备时， 不再可用，而应使用它们的同步变体。 可以是 。 包含 头文件后， 也可以是 。 同样，包含 cuda_bf16.h 头文件后，T 也可以是 。 内在函数允许在 warp 内的线程之间交换变量，而无需使用共享内存。 交换同时发生在 warp 中的所有活动线程（并以命名），根据类型移动每个线程 4 或 8 个字节的数据。
11.理解线程束执行的本质（一）
最新发布
sunhjb的博客
08-20 864
前面已经大概的介绍了CUDA执行模型的大概过程，硬件的大概结构，例如SM的大概结构，而对于硬件来说，CUDA执行的实质是线程束的执行。从软件的角度看，cuda执行中似乎所有的线程都是并行地运行的。在逻辑上这是正确的，但从硬件的角度来看，实际上硬件资源是有限的，不是所有线程在物理上都可以同时并行地执行。现在从硬件的角度来介绍，了解执行的硬件过程是有必要的。
cuda __shfl_xor
gonaYet的博客
06-04 5334
__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
CUDA C编程（十六）线程束洗牌指令
AAAA202012的博客
01-06 2068
   线 程 束 洗 牌 指 令 的 不 同 形 式    线 程 束 内 的 共 享 数 据    使 用 线 程 束 洗 牌 指 令 的 并 行 归 约    总 结   
最简示例 简介洗牌函数 之 __shfl_up_sync() cuda 之 shuffle
eloudy的专栏
02-07 1445
__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
CUDA编程---线程束洗牌指令
qq_40268672的博客
04-18 1311
从Kepler系列的GPU（计算能力为3.0或更高）开始，（shuffle instruction）作为一种机制被加入其中，只要两个线程在相同的线程束中，那么就允许这两个线程直接读取另一个线程的寄存器。。洗牌指令比共享内存有更低的延迟，并且该指令在执行数据交换时不消耗额外的内存。因此，洗牌指令为应用程序快速交换线程束中线程间的数据提供了一个有吸引力的方法。
MT.rar_进度条 线程
09-23
标题中的“MT.rar_进度条 线程”暗示了这是一个关于使用多线程来更新进度条的示例项目，而“VS2008”则表明是使用Visual Studio 2008开发的。描述进一步说明了这是一个通过工作者线程（Worker Thread）来精确控制...
DuoXiancheng.rar_多线程 进度条
09-23
在IT领域，多线程是一种常见的编程技术，用于在单个进程中同时执行多个任务，以提高应用程序的效率和响应性。在本项目"DuoXiancheng.rar"中，重点是实现两个独立的进度条，这通常涉及到并发处理和用户界面的更新。...
pb_thread多线程
02-24
3. 线程同步：为了确保线程间的正确通信和避免数据竞争，源码中可能会包含互斥锁、条件变量等同步原语的使用。例如，当两个线程试图访问同一块共享数据时，互斥锁可确保一次只有一个线程能访问。 4. 线程控制：可能...
duoxiancheng.zip_vb 多线程_vb多线程
09-24
2. **线程同步**：为了避免线程间的竞态条件和资源冲突，可以使用Monitor类的Enter和Exit方法进行互斥访问，或者使用Mutex、Semaphore等同步原语。 3. **线程池**：VB.NET支持线程池，通过ThreadPool类可以更高效地...
深入理解warp shuffle
IsayIwant的博客
04-06 1238
最后，warp 中的最后一个线程将最终的前缀和值存储在 prefix_sum 变量中。
CUDA编程与硬件架构理解
simple的博客
05-19 866
Nvidia GPU的CUDA编程模型 预览版，内容有待完善。 1. Kernel kernel函数是CUDA单个线程所执行的函数。其调用方式如下所述： kernel_name<<< Dg, Db, Ns, S >>>([kernel arguments]); Dg is of type dim3 and specifies the dimensions and size of the grid Db is of type dim3 and specifies t
cuda warp内原语——vote instructions&warp instructions
a12638915的博客
01-07 748
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
CUDA编程10 线程束基本函数与协作组
qq_43616651的博客
10-27 989
线程束（warp）一个线程块分成两个线程束，每个线程束都是连续的32个线程。从硬件上来看，一个GPU被分为若干个流多处理器（SM）。核函数中定义的线程块在执行时将被分配到还没有完全沾满的SM中。一个线程块不会被分配到不同的SM中，而总是在一个SM中，但一个SM可以有一个或多个线程块。不同的线程块之间可以并发或顺序地执行，一般来说不能同步（即使利用协作组，也只能在一些特殊的情况下进行线程块的同步）。当某些线程块完成计算任务后。对应的SM会部分或完全地空闲，然后会有新的线程块被分配到空闲的SM。
线程间怎么交换数据_在LLVM后端实现跨通道数据搬移
weixin_39536427的博客
12-04 529
作者：汪岩AMD GPU的每个CU有一个64kB的存储空间，称为本地数据共享（Local Data Share，LDS），用于同一计算单元中的work group内各个work item之间的低延迟通信和数据共享。LDS配置为32个bank，每个bank有512个4字节的条目（entry）。尽管有LDS，大多数实际的计算指令仍对寄存器中的数据进行操作。从峰值性能的角度看，global memory...
(只此一篇便绝b能懂的)cuda reduction操作优化
qq_42058018的博客
03-06 1366
aaaaaaaaaaaaaaaaaaaaaaaaaaaaqqqqqqqqqqqqqqqqqqqq2ssqeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeeee ...
CUDA10.0官方文档的翻译与学习之CUDA合作组
四国棋客
01-09 1907
目录 背景 介绍 块内组 线程组和线程块 分片（Tiled Partitions） 线程块分片（Thread Block Tiles） 伪线程混洗函数 伪线程投票函数 伪线程匹配函数 合并组 块内合作组的使用 发现模板 伪线程同步代码模板 组合 网格同步 多设备同步 结语 背景 今天我们翻译一下CUDA10.0官方文档中最后一个值得我们注意的部分——合作组 介绍 合作组(Cooperative Groups)是在CUDA 9中为了组织通信线程组而引进的对CUDA编程
评论 4
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值