CUDA Atomics

Atomic Functions

• 8スレッドが++countを実行したらcountの値はどうなる？

__device__ unsigned int count = 0;
// ...
++count;

• Read-Modify-Writeアトミック処理
  • 他のスレッドから干渉されないことが保証される
  • 順序は保証しない
• 共有メモリかグローバルメモリ
• コンピュート互換性1.1が必要 (> G80)

• 8スレッドが以下のatomicAddを実行したらcountの値はどうなる？

__device__ unsigned int count = 0;
// ...
// atomic ++count
atomicAdd(&count, 1);

• どうやってatomicAddを実装するか？

__device__ int atomicAdd(int* address, int value);

__device__ int atomicAdd(int* address, int val) {
    // でっち上げキーワード: __lock
    int old;
    __lock (address) {
        old = *address;
        *address += val;
    }
    return old;
}

• どうやってロック せずに atomicAddを実装するか？
• アトミックなコンペア＆スワップだとしたらどうだろう？

int atomicCAS(int* address, int compare, int value);

• atomicCASの疑似実装

int atomicCAS(int* address, int compare, int val) {
    // でっち上げキーワード
    __lock(address) {
        int old = *address;
        if (old == compare) *address = val;
        return old;
    }
}

• では、どうやってatomicCASを仮定してatomicAddを実装するか？

__device__ int atomicAdd(int* address, int val) {
    int old = *address, assumed;
    do {
        assumed = old;
        old = atomicCAS(address, assumed, val + assumed);
    } while (assumed != old);
    return old;
}

• どうやって異なるブロックからのスレッドは一緒に動作することができるか？
• アトミックは控えめに使う。なぜ？