CUDA 扭曲和线程发散

Posted

技术标签:

【中文标题】CUDA 扭曲和线程发散【英文标题】:CUDA Warps and Thread Divergence 【发布时间】:2014-12-09 23:05:51 【问题描述】:

我正在尝试了解 CUDA 扭曲和线程分歧。假设我有一个简单的矩阵乘法内核来乘以 n x n 矩阵。

__global__ void matrix_multiply(float* a, float* b, float* c, int n)

    int row = blockIdx.y + blockDim.y + threadIdx.y;
    int col = blockIdx.x + blockDim.x + threadIdx.x;

    if(row < n && col < n) 
        float tmp = 0.0f;
        for(int i = 0; i < n; ++i)
            tmp += a[row * n + i] * b[i * n + col];
        c[row * n + col] = tmp;
    

如果我启动一个网格大小为 32 x 32 和块大小为 16 x 16 的内核,并且矩阵为 500 x 500,那么有多少线程会遇到线程发散?

既然矩阵右边缘的每个线程块都会有线程发散,那么线程发散的warp数量不应该是256吗?

【问题讨论】:

【参考方案1】:

您的代码中有两个潜在的分歧点。第一个可以通过if 语句创建,第二个可以通过for 循环中的条件创建。从 warp 分歧 的角度来看,第二个是无害的,因为输入 n 在线程之间是一致的。

对于第一个,那些不满足条件的线程将快速退出。如果n 是 500,这似乎是,快速存在的线程数是 (16*16)*(32*32)-(500*500)=12144。考虑到this question 的答案,有 250 条经线面临发散,每条线来自通过右边缘的 16*16 最顶部块中的两行。在每个通道中,ID 为 0、1、2、3、16、17、18 和 19 的通道满足条件并进入if 块,而其余通道被禁用。将有 6*(512/16)=192 经线,if 条件对于他们的所有车道都是错误的,因此他们不会面临分歧。

下图显示了最右下角的图块中发生的情况。

【讨论】:

以上是关于CUDA 扭曲和线程发散的主要内容,如果未能解决你的问题,请参考以下文章

CUDA 扭曲和每个块的最佳线程数

有没有办法将线程显式映射到 CUDA 中的特定扭曲?

使用多少个 CUDA 核心来处理一个 CUDA 扭曲?

CUDA - 为啥基于扭曲的并行减少速度较慢?

CUDA:关于活动扭曲(活动块)以及如何选择块大小的问题

Cuda 编程中的主动扭曲