CUDA扭曲和线程分歧

Question

我试图了解CUDA扭曲和线程分歧.假设我有一个朴素矩阵乘法内核来乘以nxn矩阵.

__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乘32且块大小为16乘16并且矩阵为500乘500的内核,那么有多少warp会遇到会遇到线程分歧的线程？

由于矩阵右边缘的每个线程块都有线程发散,因此线程发散的经线数不应该是256吗？

Answer 1

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

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

下图显示了最右下方瓷砖的情况.