那些遇到过的问题

CUDA独立线程调度

Kin*_*son 5 cuda

Q1：编程指南 v11.6.0 指出以下代码模式在 Volta 及更高版本的 GPU 上有效：

if (tid % warpSize < 16) {
    ...
    float swapped = __shfl_xor_sync(0xffffffff, val, 16);
    ...
} else {
    ...
    float swapped = __shfl_xor_sync(0xffffffff, val, 16);
    ...
}
Run Code Online (Sandbox Code Playgroud)

为什么这样？

假设if分支首先执行，当线程 0~15 命中该__shfl_xor_sync语句时，它们变为非活动状态，线程 16~31 开始执行指令，直到命中相同的语句，其中前半部分和后半部分扭曲交换val。我的理解正确吗？

如果是这样，编程指南还指出“如果目标线程处于非活动状态，则检索到的值是未定义的”并且“线程可能由于多种原因而处于非活动状态，包括……采用了与当前分支路径不同的分支路径”被亚空间处决了。” 这不是意味着if和else分支都会得到未定义的值吗？

Q2：在当前实现独立线程调度（Volta~Ampere）的GPU上，当if执行分支时，不活动的线程是否仍在执行NOOP？也就是说，我是否仍应将扭曲执行视为同步执行？

Q3：同步（例如__shfl_sync、 ）是语句交错（分支中的语句 A 和 B 与分支中的 X 和 Y 交错）__ballot_sync的唯一原因吗？我很好奇当前的 ITS 与subwarp interleaving有何不同。ifelse

Rob*_*lla 9

问题一：

为什么这样？

这是一个特例。__shfl_sync()尽管编程指南中给出的陈述是正确的，但编程指南没有给出理解这种情况（据我所知）的详细行为的完整描述。要获得该指令的详细行为描述，我建议查看PTX 指南：

shfl.sync 将导致执行线程等待，直到与membermask 对应的所有未退出线程都使用相同的限定符和相同的membermask 值执行了shfl.sync，然后才恢复执行。

仔细研究该陈述可能足以理解。但我们可以稍微解压一下。

  • 如前所述，这不适用于低于 7.0 的计算能力。对于这些计算功能，成员掩码中命名的所有线程都必须参与确切的代码/指令行，并且为了使任何扭曲通道的结果有效，源通道必须在成员掩码中命名，并且不得因参与而被排除在外。在该行代码处强制发散
  • 在 cc7.0+ 情况下，我将其描述__shfl_sync()为“例外”，因为它会导致部分 warp 执行在指令的该点暂停，然后将控制/调度交给其他 warp 片段。那些其他扭曲片段将被允许继续进行（由于 Volta ITS），直到成员掩码中指定的所有线程都到达__shfl_sync()“匹配”的语句，即具有相同的成员掩码和限定符。然后执行 shuffle 语句。因此，尽管此时强制发散，但__shfl_sync()操作的行为就好像扭曲在该点充分会聚以匹配成员蒙版。

我将其描述为“不寻常”或“异常”行为。

如果是这样，编程指南还指出“如果目标线程处于非活动状态，则检索到的值是未定义的”并且“线程可能由于多种原因而处于非活动状态，包括……采用了与当前分支路径不同的分支路径”被亚空间处决了。”

在我看来，“如果目标线程处于非活动状态，则检索到的值是未定义的”这句话最直接适用于计算能力低于 7.0 的情况。如果其他地方没有相应/匹配的 shuffle 语句，它也适用于计算能力 7.0+，线程调度程序可以使用该语句来创建适当的 warp 范围（或成员掩码范围）shuffle 操作。if提供的代码示例仅给出合理的结果，因为在部分和部分 中都有匹配的操作else。如果我们将该else部分设为空语句，则代码不会为扭曲中的任何线程提供有趣的结果。

问题2：

在当前实现独立线程调度（Volta~Ampere）的GPU上，当执行if分支时，不活动的线程是否仍在执行NOOP？也就是说，我是否仍应将扭曲执行视为同步执行？

如果我们考虑一般情况，我建议考虑非活动线程的方法是它们是非活动的。如果您愿意，您可以将其称为 NOOP。由于强制发散（在我看来），此时的扭曲执行并不是在整个扭曲中“步调一致”。我不想在这里争论语义。如果您觉得准确的描述是“锁步执行，因为一些线程正在执行指令，而另一些则没有”，那是可以的。然而，我们现在已经看到，对于 shuffle 同步操作的特定情况，Volta+ 线程调度程序围绕强制发散工作，组合来自不同执行路径的操作，以满足对该特定指令的期望。

Q3：

同步（例如 __shfl_sync、__ballot_sync）是语句交错（if 分支中的语句 A 和 B 与 else 分支中的 X 和 Y 交错）的唯一原因吗？

我不相信是这样。任何时候只要有一个条件 if-else 结构导致除法内部扭曲，就有可能进行交错。我将 Volta+ 交错（图 12）定义为一个经线片段向前推进，随后另一个经线片段向前推进，可能在重新收敛之前持续交替。这种来回交替的能力不仅适用于同步操作。原子可以通过这种方式处理（这是 Volta ITS 模型的一个特定用例 - 例如在生产者/消费者算法中使用或用于锁的扭曲内协商 - 在之前链接的文章中称为“无饥饿”）我们还可以想象，扭曲片段可能会因多种原因而停滞（例如，数据依赖性，可能是由于加载指令），这会“暂时”阻止该扭曲片段的前进。我相信 Volta ITS 可以通过将前向进度调度从一个扭曲片段交替到另一个扭曲片段来处理各种可能的延迟。论文的简介（“加载使用”）中介绍了这个想法。抱歉，我无法在这里提供对该论文的详细讨论。

编辑：回应评论中的问题，解释为“在什么情况下调度程序可以使用后续的洗牌操作来满足正在等待洗牌操作完成的扭曲片段的需求？”

首先，让我们注意上面的 PTX 描述意味着某种同步。调度程序已停止执行遇到 shuffle 操作的 warp 片段，等待其他 warp 片段参与（以某种方式）。这是同步的描述。

其次，PTX 描述考虑到退出的线程。

这是什么意思呢？最简单的描述就是如果可能的话，后续的“匹配”洗牌操作可以/将“由调度程序找到”以满足洗牌操作。让我们考虑一些例子。

测试用例 1：正如编程指南中给出的，我们看到了预期的结果：

$ cat t1971.cu
#include <cstdio>
__global__ void k(){
    int tid = threadIdx.x;
    float swapped = 32;
    float val = threadIdx.x;
    if (tid % warpSize < 16) {
        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    } else {
        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    }
    printf("thread: %d, swp: %f\n", tid, swapped);
}

int main(){

    k<<<1,32>>>();
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_70 -o t1971 t1971.cu
$ ./t1971
thread: 0, swp: 16.000000
thread: 1, swp: 17.000000
thread: 2, swp: 18.000000
thread: 3, swp: 19.000000
thread: 4, swp: 20.000000
thread: 5, swp: 21.000000
thread: 6, swp: 22.000000
thread: 7, swp: 23.000000
thread: 8, swp: 24.000000
thread: 9, swp: 25.000000
thread: 10, swp: 26.000000
thread: 11, swp: 27.000000
thread: 12, swp: 28.000000
thread: 13, swp: 29.000000
thread: 14, swp: 30.000000
thread: 15, swp: 31.000000
thread: 16, swp: 0.000000
thread: 17, swp: 1.000000
thread: 18, swp: 2.000000
thread: 19, swp: 3.000000
thread: 20, swp: 4.000000
thread: 21, swp: 5.000000
thread: 22, swp: 6.000000
thread: 23, swp: 7.000000
thread: 24, swp: 8.000000
thread: 25, swp: 9.000000
thread: 26, swp: 10.000000
thread: 27, swp: 11.000000
thread: 28, swp: 12.000000
thread: 29, swp: 13.000000
thread: 30, swp: 14.000000
thread: 31, swp: 15.000000
$
Run Code Online (Sandbox Code Playgroud)

测试用例 2：删除 else 子句的主体。这仍然“有效”，因为允许退出线程满足同步点，但结果与之前的情况根本不匹配。没有一个洗牌操作是“成功的”：

$ cat t1971.cu
#include <cstdio>
__global__ void k(){
    int tid = threadIdx.x;
    float swapped = 32;
    float val = threadIdx.x;
    if (tid % warpSize < 16) {
        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    } else {
//        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    }
    printf("thread: %d, swp: %f\n", tid, swapped);
}

int main(){

    k<<<1,32>>>();
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_70 -o t1971 t1971.cu
$ ./t1971
thread: 16, swp: 32.000000
thread: 17, swp: 32.000000
thread: 18, swp: 32.000000
thread: 19, swp: 32.000000
thread: 20, swp: 32.000000
thread: 21, swp: 32.000000
thread: 22, swp: 32.000000
thread: 23, swp: 32.000000
thread: 24, swp: 32.000000
thread: 25, swp: 32.000000
thread: 26, swp: 32.000000
thread: 27, swp: 32.000000
thread: 28, swp: 32.000000
thread: 29, swp: 32.000000
thread: 30, swp: 32.000000
thread: 31, swp: 32.000000
thread: 0, swp: 0.000000
thread: 1, swp: 0.000000
thread: 2, swp: 0.000000
thread: 3, swp: 0.000000
thread: 4, swp: 0.000000
thread: 5, swp: 0.000000
thread: 6, swp: 0.000000
thread: 7, swp: 0.000000
thread: 8, swp: 0.000000
thread: 9, swp: 0.000000
thread: 10, swp: 0.000000
thread: 11, swp: 0.000000
thread: 12, swp: 0.000000
thread: 13, swp: 0.000000
thread: 14, swp: 0.000000
thread: 15, swp: 0.000000
$
Run Code Online (Sandbox Code Playgroud)

测试用例 3：使用测试用例 2，引入屏障，以防止线程退出。现在我们看到了 Volta 的悬念。这是因为与 shuffle op 关联的同步点永远无法满足：

$ cat t1971.cu
#include <cstdio>
__global__ void k(){
    int tid = threadIdx.x;
    float swapped = 32;
    float val = threadIdx.x;
    if (tid % warpSize < 16) {
        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    } else {
//        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    }
    __syncwarp();
    printf("thread: %d, swp: %f\n", tid, swapped);
}

int main(){

    k<<<1,32>>>();
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_70 -o t1971 t1971.cu
$ ./t1971
<hang>
Run Code Online (Sandbox Code Playgroud)

测试用例 4：从测试用例 2 开始，在条件区域后面引入一个额外的 shuffle 操作。在这种情况下，我们看到部分正确的结果。在条件区域中遇到洗牌操作的扭曲片段的同步点显然由在条件区域之外遇到洗牌操作的剩余扭曲片段满足。然而，正如我们将看到的，对部分正确结果的解释是，一个 warp 片段正在进行 2 次洗牌，另一个仅进行 1 次洗牌。进行两次洗牌的片段（下部片段）有第二个洗牌操作，其同步点得到满足由于退出线程条件，但其结果“不正确”，因为源通道此时不参与；他们已经退出：

$ cat t1971.cu
#include <cstdio>
__global__ void k(){
    int tid = threadIdx.x;
    float swapped = 32;
    float val = threadIdx.x;
    if (tid % warpSize < 16) {
        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    } else {
//        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    }
    swapped = __shfl_xor_sync(0xffffffff, val, 16);
    printf("thread: %d, swp: %f\n", tid, swapped);
}

int main(){

    k<<<1,32>>>();
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_70 -o t1971 t1971.cu
$ ./t1971
thread: 16, swp: 0.000000
thread: 17, swp: 1.000000
thread: 18, swp: 2.000000
thread: 19, swp: 3.000000
thread: 20, swp: 4.000000
thread: 21, swp: 5.000000
thread: 22, swp: 6.000000
thread: 23, swp: 7.000000
thread: 24, swp: 8.000000
thread: 25, swp: 9.000000
thread: 26, swp: 10.000000
thread: 27, swp: 11.000000
thread: 28, swp: 12.000000
thread: 29, swp: 13.000000
thread: 30, swp: 14.000000
thread: 31, swp: 15.000000
thread: 0, swp: 0.000000
thread: 1, swp: 0.000000
thread: 2, swp: 0.000000
thread: 3, swp: 0.000000
thread: 4, swp: 0.000000
thread: 5, swp: 0.000000
thread: 6, swp: 0.000000
thread: 7, swp: 0.000000
thread: 8, swp: 0.000000
thread: 9, swp: 0.000000
thread: 10, swp: 0.000000
thread: 11, swp: 0.000000
thread: 12, swp: 0.000000
thread: 13, swp: 0.000000
thread: 14, swp: 0.000000
thread: 15, swp: 0.000000
$
Run Code Online (Sandbox Code Playgroud)

测试用例 5：从测试用例 4 开始，最后引入同步。我们再次观察到挂起。正在执行 2 个 shuffle 操作的 warp 片段（下）没有满足其第二个 shuffle 操作同步点：

$ cat t1971.cu
#include <cstdio>
__global__ void k(){
    int tid = threadIdx.x;
    float swapped = 32;
    float val = threadIdx.x;
    if (tid % warpSize < 16) {
        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    } else {
//        swapped = __shfl_xor_sync(0xffffffff, val, 16);
    }
    swapped = __shfl_xor_sync(0xffffffff, val, 16);
    printf("thread: %d, swp: %f\n", tid, swapped);
    __syncwarp();
}

int main(){

    k<<<1,32>>>();
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_70 -o t1971 t1971.cu
$ ./t1971
thread: 16, swp: 0.000000
thread: 17, swp: 1.000000
thread: 18, swp: 2.000000
thread: 19, swp: 3.000000
thread: 20, swp: 4.000000
thread: 21, swp: 5.000000
thread: 22, swp: 6.000000
thread: 23, swp: 7.000000
thread: 24, swp: 8.000000
thread: 25, swp: 9.000000
thread: 26, swp: 10.000000
thread: 27, swp: 11.000000
thread: 28, swp: 12.000000
thread: 29, swp: 13.000000
thread: 30, swp: 14.000000
thread: 31, swp: 15.000000
<hang>
Run Code Online (Sandbox Code Playgroud)

此时挂起之前的部分打印输出是预期的。这是留给读者解释的练习：

  • 为什么我们会看到任何打印输出？
  • 为什么它是这样的（只有上面的片段，但显然有正确的洗牌结果）？

归档时间：

查看次数：

677 次

最近记录：

2 年，11 月 前
相关归档
我是否应该使用'if'语句统一两个类似的内核,从而有降低性能的风险？ 34
AMD的OpenCL是否提供类似于CUDA的GPUDirect的东西？ 10
编译时需要CUDA硬件吗？ 9
是否可以在多个gpus上运行cuda内核 7
如何诊断由于资源不足导致的CUDA启动失败？ 5
CUDA的逐元素向量乘法 5
在CUDA中共享内存中的并行扫描 4
如何获得设备的GPU架构？ 3
块中的线程和warp(32个线程)之间有什么区别？ 2
无法分配CUDA共享内存 1
难疑归档
为什么Android模拟器这么慢？我们如何加快Android模拟器的速度？ 3356
使用jQuery禁用/启用输入？ 2216
什么是三法则？ 2067
在Vim中复制整行 1555
为什么将0.1f改为0会使性能降低10倍？ 1491
传输安全性阻止了明文HTTP 1425
使用*args和**kwargs 1367
显示屏上的转换:属性 1322
是否有"以前的兄弟"CSS选择器？ 1253
删除目录的符号链接 1063