CUDAプログラムの制御フロー

Question

以下は、CUDAで競合状態がどのように起こるかを調べるために書いた小さなプログラムですが、出力には驚きました。

#include<cutil.h> 
#include<iostream> 
__global__ void testLocal(int *something, int val[]){ 

*something = *something/2; 


val[threadIdx.x] = *something; 
} 

void main(){ 

    int *a, *c; 
    int r =16; 

    cudaMalloc((void**)&a, 4*sizeof(int)); 
    cudaMalloc((void**)&c, sizeof(int)); 
    cudaMemcpy(c, &r, sizeof(int) , cudaMemcpyHostToDevice); 
    testLocal<<<1,4>>>(c,a); 
    int *b = (int *)malloc(4 * sizeof(int)); 
    cudaMemcpy(b,a, 4 * sizeof(int), cudaMemcpyDeviceToHost); 

    for(int j =0 ; j< 4; j++){ 
     printf("%d\n",b[j]); 

    } 
    getchar(); 


} 

私は4つのスレッドを起動しているので、それぞれのスレッドが*何かを2回割り切れると期待していました。私は彼らが何かを分割する順序は固定されていないことを理解しています。したがって、値を印刷しようとすると、印刷された値の1つが8、1が4、1が2、1が1となることが予想されました。しかし、すべての印刷値は8でした。 ？すべてのスレッドが*何かを一度分割するべきではない。

Answer 1

あなたが見ているのは未定義の動作です。 4つのスレッドで1つのブロックを起動しているため、すべてのスレッドが同じワープで実行されています。これは、起動したすべてのスレッドが同時に実行していることを意味します。

*something = *something/2; 

CUDAプログラミングモデルは、同じワープからの複数のスレッドが同じメモリ位置に書き込もうとすると、書き込みの1つが成功することを保証します。どのスレッドが成功するか、そして「勝利しない」ワープの他のスレッドには何も起こりません。直列化されたメモリアクセスが必要な動作を得るには、これをサポートするアーキテクチャ上でアトミックメモリアクセスプリミティブを使用することによってのみ可能です。

Answer 2

は強い言葉です。何をしているのかは不明ですので、はは何もしません。

は今、どのような可能性が高いことがすると、同じワープ以内に、同じ計算ユニット上の4つのスレッドを実行しているん。 （「SIMT」モデルは、各スレッドをワープの一部として実行させる）。 somethingでの操作はアトミックではないため、ワープ内のすべてのスレッドはロック・ステップでメモリを読み書きします。したがって、4つのスレッドは*somethingを一緒に読み込み、結果を2で割り、すべてが8をメモリに書き込もうとします。そこには、原子格差はないか、CUDAで利用可能掛けてもアトミック

何が*somethingが読まれることを、期待していたと書かれては、アトミック操作によって達成されます。だから本当にが必要な場合は、自分で書く必要があります（atomicCASの助けを借りて）。並行して実行するのが難しいスレッドを強制的に強制的に実行させるので、パフォーマンスの低下が劇的に見えるようになります。