0

フラット化された 2D マトリックスを共有メモリにロードし、x に沿ってデータをシフトし、y にも沿ってシフトしながらグローバル メモリに書き戻そうとしています。したがって、入力データは x と y に沿ってシフトされます。私が持っているもの： 

__global__ void test_shift(float *data_old, float *data_new)

{

uint glob_index = threadIdx.x + blockIdx.y*blockDim.x;

__shared__ float VAR;
__shared__ float VAR2[NUM_THREADS];

// load from global to shared

VAR = data_old[glob_index];

// do some stuff on VAR 

if (threadIdx.x < NUM_THREADS - 1)
{
VAR2[threadIdx.x + 1] = VAR; // shift (+1) along x
}

__syncthreads();

// write to global memory

if (threadIdx.y < ny - 1)
{
glob_index = threadIdx.x + (blockIdx.y + 1)*blockDim.x; // redefine glob_index to shift along y (+1)
data_new[glob_index] = VAR2[threadIdx.x];
}

カーネルへの呼び出し:

test_shift <<< grid, block >>> (data_old, data_new);

およびグリッドとブロック (blockDim.x はマトリックス幅、つまり 64 に等しい):

dim3 block(NUM_THREADS, 1);
dim3 grid(1, ny);

私はそれを達成することができません。誰かがこれのどこが悪いのか指摘してもらえますか? ストライド インデックスまたはオフセットを使用する必要がありますか?

4

2 に答える 2

1

VAR現在の形式では、グローバル メモリからロードするときにすべてのスレッドが互いのデータを走り書きするため、共有として宣言されるべきではありません: VAR = data_old[glob_index];.

また、 にアクセスするときに範囲外アクセスがVAR2[threadIdx.x + 1]あるため、カーネルが終了することはありません (デバイスの計算能力に応じて - 1.x デバイスは共有メモリ アクセスを厳密にチェックしませんでした)。

後者は、CUDA 関数へのすべての呼び出しのリターン コードでエラーをチェックすることで検出できたはずです。

于 2012-11-30T11:05:52.820 に答える
1

共有変数は、1 つのブロック内のすべてのスレッドで共有されます。これは、共有変数の blockDim.y 複合体がなく、ブロックごとに 1 つの複合体しかないことを意味します。

uint glob_index = threadIdx.x + blockIdx.y*blockDim.x;

__shared__ float VAR;
__shared__ float VAR2[NUM_THREADS];
VAR = data_old[glob_index];

if (threadIdx.x < NUM_THREADS - 1)
{
  VAR2[threadIdx.x + 1] = VAR; // shift (+1) along x
}

これは、ブロック内のすべてのスレッドに、データを 1 つの変数 (VAR) に書き込むように指示します。次に、同期は行わず、この変数を 2 番目の代入で使用します。最初のワープからのスレッドがこの変数から読み取り、2 番目のワープからのスレッドがまだそこに何かを書き込もうとしているため、これは未定義の結果になります。VAR をローカルに変更するか、ブロック内のすべてのスレッドに対して共有メモリ変数の配列を作成する必要があります。

if (threadIdx.y < ny - 1)
{
  glob_index = threadIdx.x + (blockIdx.y + 1)*blockDim.x; 
  data_new[glob_index] = VAR2[threadIdx.x];
}

VAR2[0] には、まだいくつかのゴミがあります (そこに書いたことはありません)。ブロック内の threadIdx.y は常にゼロです。

また、uint の使用は避けてください。いくつかのパフォーマンスの問題があります (または以前はありました)。

実際、このような単純なタスクでは、共有メモリを使用する必要はありません

__global__ void test_shift(float *data_old, float *data_new)
{

int glob_index = threadIdx.x + blockIdx.y*blockDim.x;

float VAR;

// load from global to local
VAR = data_old[glob_index];

int glob_index_new;
// calculate only if we are going to output something
if ( (blockIdx.y < gridDim.y - 1) && ( threadIdx.x < blockDim.x - 1 ))
{
  glob_index_new = threadIdx.x + 1 + (blockIdx.y + 1)*blockDim.x;

  // do some stuff on VAR 
} else // just write 0.0 to remove garbage
{
  glob_index_new = ( (blockIdx.y == gridDim.y - 1) && ( threadIdx.x == blockDim.x - 1 ) ) ? 0 : ((blockIdx.y == gridDim.y - 1) ? threadIdx.x : (blockIdx.y)*blockDim.x );
  VAR = 0.0;
} 

// write to global memory

data_new[glob_index_new] = VAR;
}
于 2012-11-30T11:17:50.763 に答える