CUDA 扭曲和线程发散

Question

我正在尝试了解 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吗？

Answer 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 条件对于他们的所有车道都是错误的，因此他们不会面临分歧。

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