设为首页 加入收藏

TOP

《CUDA编程:基础与实践》读书笔记(3):同步、协作组、原子函数(一)
2023-09-09 10:25:47 】 浏览:171
Tags:CUDA 编程 同步

1. 单指令多线程模式

从硬件上看,一个GPU被分为若干个SM。线程块在执行时将被分配到还没完全占满的SM中,一个线程块不会被分配到不同的SM中,一个SM可以有一个或多个线程块。不同线程块之间可以并发或顺序地执行。当某些线程块完成计算任务后,对应的SM会部分或完全地空闲,然后会有新的线程块被分配到空闲的SM。从更细的粒度看,一个SM以32个线程为单位产生、管理、调度、执行线程,这样的32个线程称为一个线程束,每个线程束包含32个具有连续线程号的线程。

在Volta架构之前,一个线程束中的线程拥有同一个程序计数器(program counter),但有各自不同的寄存器状态。在同一时刻,一个线程束中的线程只能执行一个共同的指令或者闲置,这称为单指令多线程(single instruction multiple thread, SIMT)模式。当一个线程束中的线程顺序地执行判断语句中的不同分支时,即发生了分支发散(branch divergence)。例如有如下语句:

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

首先,满足condition条件的线程会执行语句A,其它线程闲置;然后,不满足condition条件的线程执行语句B,其它线程闲置。如果语句A和语句B的指令数差不多,则整个线程束的执行效率就会降低一半,所以在编写代码时,应该尽量避免分支发散。需要注意的是,分支发散是针对同一个线程束内部线程的,不同线程束执行条件语句的不同分支则不属于分支发散。

从Volta架构开始,引入了独立线程调度机制,每个线程拥有自己的程序计数器。同时,这也使得假设了线程束同步的代码变得不再安全。如果要在Volta或者更高架构的GPU中运行一个使用了线程束同步假设的程序,可以在编译时将虚拟架构指定为低于Volta架构的计算能力,例如-arch=compute_60 -code=sm_70,这样在生成PTX代码时就使用了Pascal架构的线程调度机制,而忽略了Volta架构的独立线程调度机制。

2. 线程同步

线程块同步函数:

//保证一个线程块中的所有线程(或者说所有线程束)在执行该语句后面的语句之前都执行完了该语句前面的语句
void __syncthreads();

线程束同步函数:

//参数mask是一个代表掩码的无符号整数,默认32个比特位都为1,表示线程束中的所有线程都参与同步,如果要排除一些线程,可以把对应比特位置0,例如0xfffffffe表示排除第0号线程
void __syncwarp(unsigned mask=0xffffffff);

此外,还有一些线程束内的基本函数,它们都具有隐式的同步功能。其中线程束表决函数(warp vote functions)和线程束洗牌函数(warp shuffle functions)自Kepler架构开始就可以使用,但在CUDA 9版本中进行了更新,线程束匹配函数(warp match functions)和线程束矩阵函数(warp matrix functions)只能在Volta或更高架构的GPU中使用。

线程束内基本函数中的参数mask称为掩码,是一个32位的无符号整数,其二进制从右边数起刚好对应线程束内的32个线程。掩码用于指定要参与计算的线程,比特位等于1表示参与计算,比特位等于0表示忽略。各种函数返回的结果对于被掩码排除的线程来说没有定义,所以不要在被排除的线程中使用函数的返回值。

// ================ 线程束表决函数 ================

//如果线程束内第n个线程参与计算且pred值非0,则返回无符号整数的第n个比特位取1,否则取0。该函数相当于从一个旧的掩码产生一个新的掩码。
unsigned __ballot_sync(unsigned mask, int pred);

//线程束内所有参与线程的pred值都不为0时才返回1,否则返回0。该函数类似于这样一种选举操作,当所有参选人都同意时才通过。
int __all_sync(unsigned mask, int pred);

//线程束内所有参与线程的pred值至少有一个不为0时就返回1,否则返回0。该函数类似于这样一种选举操作,只要有一个参选人同意就通过。
int __any_sync(unsigned mask, int pred);

// ================ 线程束洗牌函数 ================
//对于所有洗牌函数,类型T可以是int、long、long long、unsigned、unsigned long、unsigned long long、float、double。
//最后一个参数width默认值为warpSize(即32),且只能取2、4、8、16、32其中的一个,它表示逻辑上线程束的大小。
//标号srcLane指的是当前线程在width范围内的位置,例如当width等于8时,srcLane的范围就是0~7。

//参与线程返回标号为srcLane的线程中变量var的值,即将一个线程的数据广播到所有(包括自己)线程。
T __shfl_sync(unsigned mask, T var, int srcLane, int width);

//标号为srcLane的参与线程返回标号为srcLane - delta的线程中变量var的值,标号srcLane < delta的线程返回自己的var值。形象地说,这是一种将数据向上平移的操作。
T __shfl_up_sync(unsigned mask, T var, unsigned int delta, int width);

//标号为srcLane的参与线程返回标号为srcLane + delta的线程中变量var的值,标号srcLane >= width - delta的线程返回自己的var值。形象地说,这是一种将数据向下平移的操作。
T __shfl_down_sync(unsigned mask, T var, unsigned int delta, int width);

//标号为srcLane的参与线程返回标号为srcLane ^ laneMask的线程中变量var的值,这里的^符号表示整数按位异或的操作。例如width等于8,laneMask等于2时,第0~7号线程分别返回标号为2、3、0、1、6、7、4、5的线程中变量var的值。
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width)

为了更好地理解上述函数,可以参考下面的测试程序,输出结果在注释中:

#include <cstdio>

const unsigned WIDTH = 8;
const unsigned BLOCK_SIZE = 16;
const unsigned FULL_MASK = 0xffffffff;

void __global__ test_warp_primitives(void);

int main(int argc, char** argv)
{
    test_warp_primitives<<<1, BLOCK_SIZE>>>();
    cudaDeviceSynchronize();
    return 0
首页 上一页 1 2 3 下一页 尾页 1/3/3
】【打印繁体】【投稿】【收藏】 【推荐】【举报】【评论】 【关闭】 【返回顶部
上一篇《CUDA编程:基础与实践》读书笔.. 下一篇1.1QT网络通信

最新文章

热门文章

Hot 文章

Python

C 语言

C++基础

大数据基础

linux编程基础

C/C++面试题目