title: 【CUDA 基础】5.6 线程束洗牌指令
categories:
- CUDA
- Freshman
tags:
- 线程束洗牌指令
toc: true
date: 2018-06-06 19:53:12
Abstract: 本文介绍线程束洗牌指令的用法
Keywords: 线程束洗牌指令
开篇废话
赶紧写博客，少说废话。
前面介绍了共享内存，常量内存，只读内存的使用，今天我们来研究一个比较特殊的机制，名字也很特殊，叫做线程束洗牌指令。
支持线程束洗牌指令的设备最低也要3.0以上，
洗牌指令，shuffle instruction作用在线程束内，允许两个线程见相互访问对方的寄存器。这就给线程束内的线程相互交换信息提供了了一种新的渠道，我们知道，核函数内部的变量都在寄存器中，一个线程束可以看做是32个内核并行执行，换句话说这32个核函数中寄存器变量在硬件上其实都是邻居，这样就为相互访问提供了物理基础，线程束内线程相互访问数据不通过共享内存或者全局内存，使得通信效率高很多，线程束洗牌指令传递数据，延迟极低，切不消耗内存
线程束洗牌指令是线程束内线程通讯的极佳方式。
我们先提出一个叫做束内线程的概念，英文名lane，简单的说，就是一个线程束内的索引，所以束内线程的ID在 KaTeX parse error: Expected 'EOF', got '\[' at position 1: \̲[̲0,31\] 内，且唯一，唯一是指线程束内唯一，一个线程块可能有很多个束内线程的索引，就像一个网格中有很多相同的threadIdx.x 一样，同时还有一个线程束的ID，可以通过以下方式计算线程在当前线程块内的束内索引，和线程束ID：
unsigned int LaneID=threadIdx.x%32;
unsigned int warpID=threadIdx.x/32;
根据上面的计算公式，一个线程块内的threadIdx.x=1,33,65等对应的laneID都是1
线程束洗牌指令的不同形式
线程束洗牌指令有两组：一组用于整形变量，另一种用于浮点型变量。一共有四种形式的洗牌指令。
在线程束内交换整形变量，其基本函数如下：
int __shfl(int var,int srcLane,int width=warpSize);
这个指令必须好好研究一下，因为这里的输入非常之乱，谁乱？var乱，一个int值，这个变量很明显是当前线程中的一个变量，作为输入，其传递的给函数的并不是这个变量存储的值，而是这个变量名，换句话说，当前线程中有var这个变量，比如我们说1号线程的var值是1，那么2号线程中的var值不一定是1，所以，这个__shfl返回的就是var值，哪个线程var值呢？srcLane这个线程的，srcLane并不是当前线程的束内线程，而是结合with计算出来的相对线程位置，比如我想得到3号线程内存的var值，而且width=16，那么就是，015的束内线程接收0+3位置处的var值，也就是3号束内线程的var值，1632的束内线程接收16+3=19位置处的var变量。
这个是非常重要的，虽然有些困难，但是却相当灵活。width的默认参数是32.srcLane我们后面简单的叫他束内线程，注意我们必须心理明白只有width是默认值的时候，他才是真正的束内线程。
图示如下
接着是另一个指令，其主要特点是从与调用线程相关的线程中复制数据。
int __shfl_up(int var,unsigned int delta,int with=warpSize);
这个函数的作用是调用线程得到当前束内线程编号减去delta的编号的线程内的var值，with和__shfl中都一样，默认是32，作用效果如下：
如果是with其他值，我们可以根据前面的讲解，把线程束再分成若干个大小为with的块，进行上图的操作。
最左边两个元素没有前面的delta号线程，所以不做任何操作，保持原值。
同样下一个指令是上面的反转版本：
int __shfl_down(int var,unsigned int delta,int with=warpSize);
作用和参数和up一模一样，图示如下：
最后一个洗牌指令比较夸张，也是很灵活的一个指令
int __shfl_xor(int var,int laneMask,int with=warpSize);
xor是异或操作，这个指令如果学过硬件或者c语言学的比较扎实的人可能知道，这是电路里面最最最最最重要的操作，没有之一，什么是异或？逻辑中我们假设只有0，1两种信号，用 “^”表示异或:
0^0=0;
1^0=1;
0^1=1;
1^1=0;
二元操作，只要两个不同就会得到真，否则为假
那么我们的__shfl_xor操作就是包含抑或操作在内的洗牌指令，怎么算呢？
如果我们输入的laneMask是1，其对应的二进制是 000⋯001000\cdots001000⋯001 ,当前线程的索引是031之间的一个数，那么我们用laneMask与当前线程索引进行抑或操作得到的就是目标线程的编号了，这里laneMask是1，那么我们把1与031分别抑或就会得到：
000001^000000=000001;
000001^000001=000000;
000001^000010=000011;
000001^000011=000010;
000001^000100=000101;
000001^000101=000100;
.
.
.
000001^011110=011111;
000001^011111=011110;
这就是当前线程的束内线程编号和目标线程束内县城编号之间的对应关系，画成图会非常犀利：
这就是4个线程束洗牌指令对整形的操作了。对应的浮点型不需要该函数名，而是只要把var改成float就行了，函数就会自动重载了。