ガウス消去の並列処理

Question

ガウス消去のためのCUDAでシングルスレッドプログラムを正常に実装しており、並列処理を実現したいと考えています。我々はGPU上のコードは、下三角行列に1Dアレイ（行列）を形質転換し、その後CPU上で、私は戻ってまいります見ることができるように

__global__ void ParallelGaussian(double* A) 
{ 
    int index = threadIdx.x; 
    int stride = blockDim.x; 

    if (index < ROWS) //Skip additional threads 
    { 
     for (unsigned int r = index; r < ROWS; r += stride) 
     { 
      //Forward elimination to reduce to row echelon form 
      for (unsigned int k = r + 1; k < ROWS; ++k) 
      { 
       double c = -A[(ROWS + 1) * k + r]/A[(ROWS + 1) * r + r]; 
       for (unsigned int j = r; j < ROWS + 1; ++j) 
       { 
        if (r == j) 
         A[(ROWS + 1) * k + j] = 0.0; 
        else 
         A[(ROWS + 1) * k + j] += c * A[(ROWS + 1) * r + j]; 
       } 
      } 
     } 
    } 
} 

：この時点まで並列コードは次のようになります最後の結果を得るための置換。このアプローチでは、完全には必要ではないが、実際にはアルゴリズムの数値的安定性が向上するため、ピボット操作は行われません。

シングルスレッドとブロックの動作とカーネルを起動し、行階段形に行列を変換：

ParallelGaussian << < 1, 1 >> >(dev_a); 

しかし、私は

ParallelGaussian << < 1, 32 >> >(dev_a); 

ように、スレッドの数を増やしたい場合

下三角行列を生成できません。ブロック内のスレッドを同期させるために__syncthreads（）呼び出しをコードに追加しても、これまでの状況が改善されず、理由がわかりません。

Answer 1

内部ループを考慮してください。すべてのスレッドはAにアクセスし、kとjはrから行列の最後まで実行されるため、複数のスレッドが同じA[(ROWS + 1) * k + j]値を変更する可能性があります。

他のスレッドがその値を更新している間に、A[(ROWS + 1) * r + j]にアクセスするスレッドがある可能性もあります。

解決策の1つは、各スレッドを個々の結果配列に累積させ、それらを最後に組み合わせることです。これはメモリ集中です。

もう1つは、特定の値に書き込むスレッドが1つだけで、新しいスレッドに必要な値を変更しないようにその値を新しい行列に格納するように再構成することです。