CUDA：whileループの正確さ

Question

このカーネルは正しい結果を得ています。パフォーマンスを向上させたい場合、私の問題はwhileループの正確さです。私はブロックとスレッドのいくつかの設定を試みたが、私はそれらを変更するつもりなら、whileループは私に正しい結果を与えません。 カーネルの設定を変更して得た結果は、firstArrayとsecondArrayが完全には満たされない（セル内に0が入る）ということです。ループの場合、両方の配列はから取得されたcurValueで埋められる必要があります。

何かアドバイスは歓迎される:)

をあなたはいくつかの意味の最適化を行うことを妨げ、ここでデータ依存性の問題を持って事前に

#define N 65536 

__global__ void whileLoop(int* firstArray_device, int* secondArray_device) 
{ 
    int curValue = 0; 
    int curIndex = 1; 

    int i = (threadIdx.x)+2; 

    while(i < N) { 
     if (i % curIndex == 0) { 
      curValue = curValue + curIndex; 
      curIndex *= 2; 
     } 
     firstArray_device[i] = curValue; 
     secondArray_device[i] = curValue; 
     i += blockDim.x * gridDim.x; 
    } 
} 

int main(){ 

    firstArray_host[0] = 0; 
    firstArray_host[1] = 1; 

    secondArray_host[0] = 0; 
    secondArray_host[1] = 1; 


    // memory allocation + copy on GPU 

    // definition number of blocks and threads 
    dim3 dimBlock(1, 1); 
    dim3 dimGrid(1, 1); 

    whileLoop<<<dimGrid, dimBlock>>>(firstArray_device, secondArray_device); 

    // copy back to CPU + free memory 
} 

Answer 1

、ありがとうございました。変数curValueとcurIndexはwhileループ内で変更され、次の実行にフィードフォワードされます。ループを最適化しようとするとすぐに、この変数の状態が異なり、結果が変更される状況になります。

私はあなたが達成しようとしていることを本当に知っていませんが、依存関係を避けるためにループの前回の実行の値に依存してwhileループを作ろうとします。 indizeし、threadIdx、blockDim、gridDimなどの環境状態で値を計算する方法で、データをスレッドとデータのチャンクに分割してみてください。

条件付きループを回避してください。実行回数が一定のforループを使用する方が良いです。これはまた、最適化する方が簡単です。

Answer 2

いくつかのこと：

1. あなたが デバイス上のグローバル配列を宣言するために使用されるコードを残しました。この情報を入手すると便利です。
2. 複数のブロックが使用されている場合、アルゴリズムは でスレッドセーフではありません。つまり、複数の ブロックを実行している場合は、冗長な作業（つまり、 は得られない）だけでなく、同じグローバルメモリの場所に を書き込んでエラーを作成しようとする可能性もあります。
3. コードが の場合は1ブロックしか使用しませんが、これは意味がありません。パラレルデバイスでシリアルまたは軽度のスレッド操作を実行しています。
   
   ：あなたは現在、このコードを持つ2つの主要な問題は、平行な観点からある

   ...すべての利用可能なリソース（メモリの競合なしで複数のSMP上の複数のブロック（下記参照）上で実行することはできません

4. int i = (threadIdx.x)+2; ... シングルスレッドの2の開始インデックスを生成する;私はこれが最初の2つのよう望むもの ある疑うように単一ブロック内の2つのスレッドの2と3、と。ポジション（0,1）は決して になります。 （配列はCのインデックス0で始まることに注意してください。）

5. さらに、あなたは複数ブロック（2つのブロック つのスレッドでそれぞれ言う）2 BX 1トンのために、あなたは、複数の重複したインデックス （例えばを持っているでしょう含まれている場合 - >インデックスがb1t1：2を、b1t2： 2）、インデックス を使用してグローバルメモリに書き込むと、競合とエラーが発生します。 int i = threadIdx.x + blockDim.x * blockIdx.x;のようなことは、この問題を避けるためにインデックスを正しく計算する典型的な方法です。その は私にスレッドの合計＃に数相当を追加し、これ が追加ぶつかり合ったりオーバーラップを作成していないため

6. あなた 最後の式i += blockDim.x * gridDim.x;は、大丈夫です。

7. なぜGPUを使用してメモリをシャッフルし、簡単な計算を行うのですか？アレイをデバイスに入れたり離したりする時間を考慮すると、高速なCPUと比べてスピードアップのスピードが向上することはありません。

問題1と2の作業は、あなたが望むなら、それを超えて、あなたの全体的な目標と種類アルゴリズムのまさにあなたが最適化し、より平行に優しいソリューションを思い付くしようとしている考える - GPUかどうかを検討しますコンピューティングは本当にあなたの問題に理にかなっている

Answer 3

このアルゴリズムを並列化するには、配列内の特定のインデックスの値を直接計算できる数式を用意する必要があります。そのため、配列の範囲内でランダムなインデックスを選択し、その位置の値がどのようなものになるかを判断する要因が何であるかを検討します。数式を見つけたら、ランダムなインデックスの出力値とシリアルアルゴリズムの計算値を比較してテストします。それが正しい場合は、スレッドとブロック索引に基づいて一意の索引を選択することで始まるカーネルを作成します。そのインデックスの値を計算し、配列の対応するインデックスに格納します。

簡単な例：

シリアル：

__global__ void serial(int* array) 
{ 
    int j(0); 
    for (int i(0); i < 1024; ++i) { 
    array[i] = j; 
    j += 5; 
} 

int main() { 
    dim3 dimBlock(1); 
    dim3 dimGrid(1); 
    serial<<<dimGrid, dimBlock>>>(array); 
} 

パラレル：

__global__ void parallel(int* array) 
{ 
    int i(threadIdx.x + blockDim.x * blockIdx.x); 
    int j(i * 5); 
    array[i] = j; 
} 

int main(){ 
    dim3 dimBlock(256); 
    dim3 dimGrid(1024/256); 
    parallel<<<dimGrid, dimBlock>>>(array); 
}