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コンパイル時に (またはビジュアル スタジオでデバッグ プロジェクトをコンパイルして) ​​スイッチを追加します。これにより、多くのコンパイラの最適化が無効になります。

あなたのcc1.3デバイスと私が最初に観察したもの（私のcc2.0デバイスで）とJackOLanternがcc2.1デバイスで観察したものとの間に動作の違いがある理由は、cc1.xデバイスnvccに対して異なるデバイスコンパイラを使用しているためです。すべての 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 に答える