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，那么有多少经线具有会遇到线程分歧的线程？

既然矩阵右边缘的每个线程块都会有线程发散，那么线程发散的经纱数不应该是256吗？

Answer 1

您的代码中有两个潜在的分歧点。第一个可以由 if 创建语句和第二个条件在 for环形。由于输入 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条件对于他们所有的车道都是假的，因此他们不会面临分歧。

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