Staskaer

Hi, this is Staskaer

0%

CUDA高性能并行计算(10)--线程束基本函数与协作组

Posted on Edited on In Valine:

CUDA线程束基本函数与协作组

由于GPU的每次调度计算的最小单元就是一个拥有32个线程的线程束，所以我们首先再来回顾一下线程束和GPU的处理模式。

单指令多线程的执行模式

一个GPU在硬件上是由多个SM组成的，核函数中定义的线程块在执行时会被调度到一个资源还没有被完全占满的SM上。一个线程块只会被一个SM调度，但是一个SM上可以拥有很多线程块。

不同的线程块之间可以并发或顺序的执行，但是线程块之间一般不能同步，当某个线程块完成计算后，SM会进入部分或完全空闲状态并调度下一个线程块。

一个SM以线程束作为最小单元进行调度，一个线程束为32个线程。

线程块分支发散

在伏特架构之前，每个线程束只有一个程序计数器，导致所有的线程必须执行相同的指令。如果一个线程束内的线程顺序地执行分支语句中的不同分支时，就会导致分支发散。

1
2
3
4
5
6
7
8
9
10

if(condition)
{
    A;
}
else
{
    B;
}

例如上述代码，在一个线程束内，首先满足condition的线程会执行A，其他线程闲置，然后不满足condition的线程再去执行B，其他线程闲置。因此，如果A,B的运算量相差不多，就会导致性能下降到原来的一半。而如果存在很多分支(switch-case)时，就会导致分支发散的很严重，性能严重下降。

而在伏特架构开始，引入了线程独立调度，每个线程都有自己的程序计数器。因此分支发散的问题没有那么严重了，但是又引入了新的问题：

首先是由于CUDA核非常多，所以程序计数器也得非常多，导致空间浪费。另外，由于每个线程有自己的程序计数器，所以线程束也不再是完全同步执行的了。要引入线程束内同步，可以通过__syncwarp()来完成。

线程块的基本函数

线程束内线程同步函数

当需要同步的所有线程都位于一个线程块时，可以使用更加廉价的__syncwarp(unsigned mask = 0xffffffff)去代替线程块同步，其内部的掩模表示需要参与同步的线程块内的线程id。

但使用的时候需要十分注意，要保证范围内的线程同属一个线程块。

线程束表决函数

表决函数名字感觉有些奇怪，其函数族大致形如__xxx_sync(mask, predicate)，效果是线程块内所有线程求predicate的值，并根据所有线程块的计算结果来返回值。

函数 功能
unsigned __ballot_sync(unsigned mask, int predicate) mask掩模为1的线程参与计算。返回值也是一个掩模，其中位1的位表示对应的线程计算结果非0。相当于从旧的掩模中产生新的掩模
int __all_sync(unsigned mask, int predicate) make掩模为1的线程参与计算。只有所有参与计算的线程的计算结果均为非0，返回值为1，否则为0.所有人都同意才执行
int __any_sync(unsigned mask, int predicate) make掩模为1的线程参与计算。只要参与计算的线程中有一个计算结果非0，那么就返回1，否则返回0.只要有一个人同意就执行

这些函数结尾都带有_sync，因此其都具有隐式同步功能。

下面给出一个例子

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>
#include <stdlib.h>

#define TARGET 1
#define WIDTH 8
#define BLOCK_SIZE 16
#define FULL_MASK 0Xffffffff

__global__ void test_warp_primitives(void)
{
	int tid = threadIdx.x;
	int lane_id = tid % WIDTH;

	// 输出线程号
	if (tid == 0) 
		printf("threadIdx.x: ");
	printf("%2d ", tid);
	if (tid == 0)
		printf("\n ");

	// 输出定义的线程束内线程号
	if (tid == 0)
		printf("lane id: ");
	printf("%2d ", lane_id);
	if (tid == 0)
		printf("\n ");

	// 使用表决函数计算mask并打印
	unsigned mask1 = __ballot_sync(FULL_MASK, tid > 0);
	unsigned mask2 = __ballot_sync(FULL_MASK, tid == 0);
	if (tid == 0)
		printf("fullmask = %x\n", FULL_MASK);
	if (tid == 1)
		printf("mask1 = %x\n", mask1);
	if (tid == 0)
		printf("mask2 = %x\n", mask2);

	// 下面也是表决函数的使用
	int result = __all_sync(FULL_MASK, tid);
	if (tid == 0)
		printf("all_sync(fullmask) = %d\n", result);

	result = __any_sync(mask1, tid);
	if (tid == 1)
		printf("any_sync(mask1) = %d\n", result);

	result = __any_sync(FULL_MASK, tid);
	if (tid == 0)
		printf("any_sync(fullmask) = %d\n", result);

	result = __any_sync(mask2, tid);
	if (tid == 0)
		printf("any_sync(mask2) = %d\n", result);
}

int main()
{
	test_warp_primitives << <1, BLOCK_SIZE >> > ();
	return 0;
}

这个例子介绍了表决函数的使用，非常简单，不再赘述。

线程束洗牌函数

下面介绍洗牌函数的功能

函数 功能
T __shfl_sync(unsigned mask,T v,int srcLane,int w = warpSize) 参与线程返回标号为scrLane的线程中变量v的值。即广播数据交互，将一个线程的数据广播到线程束中所有的线程上(包括自身)。
T __shfl_up_sync(unsigned mask,T v,unsigned d,int w = warpSize) 标号为t的参与线程中会返回t-d的线程中变量的值。标号满足t-d<0的线程返回原来v的值，也就是数据向标号高的线程传递
T __shfl_down_sync(unsigned mask,T v,unsigned d,int w = warpSize) 标号为t的参与线程中会返回t+d的线程中变量的值。标号满足t+d>w的线程返回原来v的值，也就是数据向标号低的线程传递
T __shfl_xor_sync(unsigned mask,T v,int laneMask,int w = warpSize) 标号为t的参与线程返回标号为t^laneMask的线程中的变量为v的值

洗牌函数可以用于将数据进行线程内传递，同样是同步的。注意其中的参数w表示作用的范围，比如其值取8的时候，表示把线程束再按照8个线程一组分成更小的'线程束'

可能会觉得这几个函数令人费解，下面来具体讲述一下。

首先线程是从核函数创建的，所以每个线程内部的代码都是一样的(类似c语言的fork)，这也就是说，假如在核函数中写了一个洗牌函数，那么实际上所有的线程内都会有这个洗牌函数。假如在核函数内有这么一句

1
2
val = __shfl_up_sync(0xffffffff, tid, 2);
// 其中tid = threadIdx.x

由于所有线程都参与了，就不用考虑参与线程的问题。

下表给出了16个线程中val的值，其中tid_{i}表示这个值来自于线程i

线程编号线程内val值线程编号线程内val值
00,即tid_0，下同86,即tid_6，下同
1197
20108
31119
421210
531311
641412
751513
861614

同样给出一个例子

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include <stdio.h>
#include <stdlib.h>

#define TARGET 2
#define WIDTH 8
#define BLOCK_SIZE 16
#define FULL_MASK 0Xffffffff

__global__ void test_warp_primitives(void)
{
	int tid = threadIdx.x;
	int lane_id = tid % WIDTH;

	// 输出线程号
	if (tid == 0)
		printf("threadIdx.x: ");
	printf("%2d ", tid);
	if (tid == 0)
		printf("\n ");

	// 输出定义的线程束内线程号
	if (tid == 0)
		printf("lane id: ");
	printf("%2d ", lane_id);
	if (tid == 0)
		printf("\n ");

	// 把2号线程的值广播到0-7号中，把10号线程的值广播到8-15中
	int value = __shfl_sync(FULL_MASK, tid, 2, WIDTH);
	if (tid == 0)
		printf("shfl : ");
	printf("%2d ", value);
	if (tid == 0)
		printf("\n");

	// 在每个小的线程束内把值向上传递
	value = __shfl_up_sync(FULL_MASK, tid, 1, WIDTH);
	if (tid == 0)
		printf("shfl_up : ");
	printf("%2d ", value);
	if (tid == 0)
		printf("\n");

	// 在每个小的线程束内把值向下传递
	value = __shfl_down_sync(FULL_MASK, tid, 1, WIDTH);
	if (tid == 0)
		printf("shfl_down : ");
	printf("%2d ", value);
	if (tid == 0)
		printf("\n");

	// 类似上面，线程号异或后判断传递方向
	value = __shfl_xor_sync(FULL_MASK, tid, 1, WIDTH);
	if (tid == 0)
		printf("shfl_xor : ");
	printf("%2d ", value);
	if (tid == 0)
		printf("\n");

}

int main()
{
	test_warp_primitives << <1, BLOCK_SIZE >> > ();
	return 0;
}

主要使用了洗牌函数并查看其在线程束内的工作情况。

协作组

在有些算法中，需要若干线程之间进行协作，而协作，就需要用到某种同步机制，否则乱序是没有办法进行很好的协作的。

协作组可以看成是线程块与线程束的同步机制的推广，它提供了更加灵活的线程协作方式，包括线程块内部同步、线程块之间(计算网格)同步和设备与设备之间的同步。

使用协作组的相关函数时需要包含头文件#include <cooperative_groups.h>，且声明命名空间using namespace cooperative_groups;

线程块级别的协作组

协作组编程中最基本的类型就是线程组thread_group类，我们把其头文件中源码放上来

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35

class thread_group
{
    friend _CG_QUALIFIER thread_group this_thread();
    friend _CG_QUALIFIER thread_group tiled_partition(const thread_group& parent, unsigned int tilesz);
    friend class thread_block;

 protected:
    union __align__(8) {
        unsigned int type : 8;
        struct {
            unsigned int type : 8;
            unsigned int size : 24;
            unsigned int mask;
        } coalesced;
        struct {
            void* ptr[2];
        } buffer;
    } _data;

    _CG_QUALIFIER thread_group operator=(const thread_group& src);
    _CG_QUALIFIER thread_group(__internal::groupType type) {
        _data.type = type;
    }

#if __cplusplus >= 201103L
    static_assert(sizeof(_data) == 16, "Failed size check");
#endif

public: // 接口
    _CG_QUALIFIER unsigned int size() const;
    _CG_QUALIFIER unsigned int thread_rank() const;
    _CG_QUALIFIER void sync() const;
};

我们可以看到其接口只有三个函数

  • size() : 返回组的大小
  • sync() : 同步
  • thread_rank() : 返回当前调用该函数的线程在组内标号(从0开始)

另外我们注意到其有一个友元类thread_block，这个类公有继承自thread_group，且提供了额外的公有接口，我们介绍两个

  • dim3 group_index() : 等价于blockIdx
  • dim3 thread_index() : 等价于threadIdx

我们可以使用如下方式定义一个thread_block对象
thread_block g = this_thread_block();
其中this_thread_block()也是一个常量，这样定义的g就变成了我们熟悉的线程块。调用g.sync()等价于__syncthreads()函数。而g.group_index()和g.thread_index()也完全等于内建变量。

可以使用函数tiled_partition()来讲一个线程块进行分割

1
2
3
4
5
6
7

// 将线程块分成32一组的线程束
thread_block g32 = tiled_partition(this_thread_block(), 32);

// 也可以分割成更细的线程束
thread_block g4 = tiled_partition(g32), 4);

当线程组大小在编译时就能确定，可以使用静态结构进行定义，更高效。

1
2
3
4

thread_block_tile<32> g32 = tiled_partition<32>(this_thread_block());
thread_block_tile<4> g4 = tiled_partition<4>(g32);

线程组内也可以使用表决函数和洗牌函数等，但是其不能有掩模和宽度，也就是所有的线程必须参与进去。

总结

  • __syncwarp()可以用于线程束内同步，__syncthreads()则用于线程块同步。
  • 表决函数可以用于线程束内线程计算表达式并返回计算结果，可以用于将一个掩模映射成另一个掩模。表决函数时同步的。
  • 洗牌函数可以用于将数据进行线程内传递，同样是同步的。洗牌函数可以指定作用在更小的"线程束"上。
  • 使用协作组的相关函数时需要包含头文件#include <cooperative_groups.h>，且声明命名空间using namespace cooperative_groups;。
  • 协作组使用tiled_partition()进行分割，thread_block类用于确定线程块协作组。