1

我目前正在做 CUDA 编程的作业，我发现自己被迫同步内核中的所有线程。我已经实现了本文所述的简单同步机制。但是我遇到了奇怪的行为，所以我决定为这种锁编写一个测试程序：

#include <stdio.h>

__device__ int g_mutex = 0;
__device__ void __gpu_sync(int goalVal) {
    int tid_in_block = threadIdx.x * blockDim.y + threadIdx.y;
    if(tid_in_block == 0) {
        atomicAdd(&g_mutex, 1);
        while(g_mutex != goalVal) {}
    }
    __syncthreads();
}

__global__ void deadlock(int *out) {
    __gpu_sync(1000);
    *out = 42;
}

int main() {
    int *dev, local;
    cudaMalloc((void**)&dev, sizeof(int));
    deadlock<<<1,1>>>(dev);
    cudaMemcpy(&local, dev, sizeof(int), cudaMemcpyDeviceToHost);
    printf("%d\n", local);
    return 0;
}

我希望这个应用程序永远不会终止（因为互斥锁永远不会达到 1000 的值）。然而，应用程序运行时好像没有涉及锁定，并立即打印出 42。你能告诉我我缺少什么吗？我正在运行 CC1.3 系统（GTX 260）、64 位 Windows 7、CUDA 5.5。编译由nvcc -arch compute_12 main.cu.

谢谢！

4

1 回答 1

3

编译器正在优化会死锁的代码，因为从该线程的角度来看，它没有任何作用（该代码不会修改任何可识别的状态）。

如果您想查看代码死锁，请在-G编译（或在 Visual Studio 中编译调试项目）时添加开关，这将禁用许多编译器优化。

您的 cc1.3 设备与我最初观察到的（在我的 cc2.0 设备上）以及 JackOLantern 在他的 cc2.1 设备上观察到的行为之间存在差异的原因是，nvcc对于 cc1.x 设备使用不同的设备编译器作为与所有 cc2.x 和更新的设备相比，确切的优化行为可能会有所不同。

我在 cc1.3 设备（在 linux 下）上尝试了您的代码，并且能够重现您的观察结果，即使使用-arch=sm_13

对您的设备代码进行以下修改将阻止编译器优化 while 循环（即使您没有指定-G）：

__device__ int __gpu_sync(int goalVal) {
    int test;
    int tid_in_block = threadIdx.x * blockDim.y + threadIdx.y;
    if(tid_in_block == 0) {
        atomicAdd(&g_mutex, 1);
        while(g_mutex != goalVal) {test++;}

    }
    __syncthreads();
    return test;
}

__global__ void deadlock(int *out) {
    *out = __gpu_sync(1000);
}
于 2013-11-12T18:53:52.320 回答