ブロックレベルアトミック書き込み

Question

ブロックレベルでアトミックな書き込みが可能ですか？一例として
は、以下の点を考慮：あなただけの原子の追加を必要とするよう

__global__ kernel (int atomic) 
{ 
    atomic+=blockid.x; //should be atomic for each block 
} 

Answer 1

あなたがブロック/ブロックレベルで何を意味するかは不明ながら、それが聞こえます。 ものは

__global__ kernel (int atomic) 
{ 
    atomic_add(blockid.x,&atomic); 
} 

atomicはタイプatomic_tであることと、int型をblockid.xしなければならないだろう#include <asm/atomic.h>あなたのコード内のカーネルで発見されています。

Answer 2

CUDAでいくつかのアトミック操作を行うことができます。 「CUDAプログラミングガイド」の「Auntix B.11 Atomic関数」を参照してください。すなわち：

__global__ void kernel (int *result) 
{ 
    atomicAdd(result, blockIdx.x); // 
} 

また、どちらの例は、グローバルメモリで動作変数

__global__ void kernel (int *result) 
{ 
    atomicExch(result, blockIdx.x); // 
} 

の値を交換することができます。

Atomic functions operating on shared memory and atomic functions operating on 64-bit words are only available for devices of compute capability 1.2 and above.

よろしく。

Answer 3

共有メモリ上でアトミック操作を実行することはできますが、コードスニペットで行ったのとは異なります。カーネルのintパラメータはスレッド固有の変数です。すべてのスレッドが起動時に指定した値と同じ値を取得しても、共有メモリには保存されません。また、それを原子的に操作することは意味がありません。

たとえば、int *をいくつかのバッファに渡した場合、それはグローバルメモリのバッファになります。 @ pQBのanswerで説明されているように、グローバルメモリ内のデータに対してデバイス全体のアトミック操作を実行できます。しかし、あなたはブロックレベルのアトミック操作について尋ねましたが、それはグローバルデータにはあまり意味がありません。それでも、あなたのスレッドの1つがあるグローバルアドレスに書き込む場合、この書き込みの効果がブロック内の他のすべてのスレッドに見えるまで、すべて__threadfence_block()がストールすることがあります。

適切なブロックレベルアトミックはCUDAではサポートされていますが、共有メモリでもサポートされています。 this Parallel4All blog entryまたは関連するsectionCUDA Programming Guideに共有メモリを使用する方法をお読みください。

__shared__ int xがある場合、実際にはグローバルアトミックと同じ構文でブロックレベルのアトミック操作を実行できます。atomicAdd(&x, 7)は、xの値に123をアトミックに追加します。しかし、ブロック内のすべてのスレッドが同じ処理を行い、一度に最大1024個のアトミック書き込みを試したくないことを覚えておいてください。複数のスレッドが同じ場所に書き込まれますが、必ずしもそうではないところおそらく通常は、

__shared__ some_buffer[NumFoosPerBar]; 

// ... 

if (check_condition()) { 
    int foo_index = get_thread_foo_index_for(threadIdx.x); 
    atomicAdd(&some_buffer[foo_index], 7); 
} 

のようなものを持っているでしょう。 の場合、は複数の書き込みを期待します。アトミックは使用せず、書き込む値をいくらか減らします。