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 扭曲和线程发散的主要内容,如果未能解决你的问题,请参考以下文章