1

これを尋ねる理由は、コードに奇妙なバグがあり、エイリアシングの問題である可能性があると思われるためです。

__shared__ float x[32];
__shared__ unsigned int xsum[32];

int idx=threadIdx.x;
unsigned char * xchar=(unsigned char *)x;
//...do something
 if (threadIdx.x<32)
 {
    xchar[4*idx]&=somestring[0];
    xchar[4*idx+1]&=somestring[1];
    xchar[4*idx+2]&=somestring[2];
    xchar[4*idx+3]&=somestring[3];

    xsum[idx]+=*((unsigned int *)(x+idx));//<-Looks like the compiler sometimes fail to recongize this as the aliasing of xchar;
 };
4

2 に答える 2

2

コンパイラは、互換性のある型間のエイリアシングのみを尊重する必要があります。charとは互換性がないためfloat、コンパイラはポインターが決してエイリアスしないと自由に想定できます。

float に対してビット単位の操作を行う場合は、まず ( を使用して__float_as_int()) 符号なし整数に変換し、次にそれを操作してから、最後に ( を使用して__int_as_float()) float に戻します。

于 2013-03-29T01:18:23.700 に答える
1

ここに競合状態があると思います。しかし、私は何であるかわかりませんsomestring。すべてのスレッドで同じ場合は、次のようにできます。

__shared__ float x[32];

unsigned char * xchar=(unsigned char *)x;

//...do something

if(threadIdx.x<4) {
     xchar[threadIdx.x]&=somestring[threadIdx.x];
}

__syncthreads();

unsigned int xsum+=*((unsigned int *)x);

これは、すべてのスレッドが同じ配列を共有することを意味するため、xsum はすべてのスレッド間で同じになります。各スレッドに独自の配列が必要な場合は32*number_of_threads_in_block、オフセットの配列を割り当てて使用する必要があります。

PS: 上記のコードは 1D ブロックでのみ機能します。2D または 3D では、独自の threadID を計算し、コードを実行するスレッドが 4 つだけであることを確認する必要があります。

于 2013-03-29T00:56:51.523 に答える