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;
    }
}

如果我使用32x32的网格大小和16x16的块大小启动内核，并且矩阵为500x500，那么将遇到线程分歧的线程有多少个warp？

由于矩阵右侧的每个线程块都会出现线程分歧，所以具有线程分歧的warp数量应该是256吗？

- csnate

Answer 1

你的代码中有两个潜在的分歧点。第一个可以由if语句创建，第二个可以由for循环中的条件创建。从线程分歧的角度来看，第二个是无害的，因为输入n在各个线程上是均匀的。

对于第一个问题，不满足条件的线程将快速退出。如果n为500，似乎是这样，快速退出的线程数为(16*16)*(32*32)-(500*500)=12144。考虑到这个问题的答案，会出现250个面临分歧的warp，每个warp来自16*16最上方的两行块，通过右边缘。在这些warp中，具有ID 0、1、2、3、16、17、18和19的lane满足条件并进入if块，而其余的则被禁用。会有6*(512/16)=192个warp，他们所有lane的if条件都为false，因此他们没有面临分歧的问题。

下面的图片显示了在右下角的图块中发生的情况。

- Farzad