将 int 与 threadIdx.x 进行比较时的 CUDA 内核无限循环

Question

我在使用更大的内核时遇到了问题，但它似乎可以提炼出以下代码，内核永远不会从中返回。 有人可以解释为什么会有无限循环吗？

__global__ void infinite_while_kernel(void)
{
    int index = 0;
    while (index >= threadIdx.x) {
        index--;
    }
    return;
}
int main(void) {
    infinite_while_kernel<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

此外，下面的内核也被卡住了：

__global__ void not_infinite_while_kernel(void)
{
    int index = 0;
    while (index >= (unsigned int) 0u*threadIdx.x) {
        index--;
    }
return;
}

按预期的那样，将原始内核中的threadIdx.x替换为0将返回。 我正在使用 v5.5 工具包，并使用 -arch=sm_20 -O0 标志进行编译。 在特斯拉M2090上运行。 我目前无法访问任何其他硬件，也无法访问工具包版本（这不是我的系统）。

Answer 1

这段代码也挂在普通C++中（试试看）：

int main(){
  int index = 0;
  while (index >= 0U){
    index--;
    }
  return 0;
}

将有符号值与无符号值

进行比较时，编译器会将有符号值转换为无符号值。

threadIdx.x是一个无符号值。 代码中未标记的0常量不是。

作为无符号比较，您的测试始终为真，因此 while 循环永远不会退出。

另请注意，您的__global__函数应用 void 进行装饰。

最后，在内核启动后，代码中没有cudaDeviceSynchronize()或其他屏障，即使内核挂起，程序无论如何都会"正常"退出。

因此，我认为您发布的代码实际上不会重现您所描述的问题，但是如果您添加cudaDeviceSynchronize()则会重现。