0

私はCUDAを初めて使用し、テストプログラムを作成しようとしています。GeForce GT 520 カードでアプリケーションを実行していますが、パフォーマンスが非常に低下します。

アプリケーションを使用して画像を処理し、各行を個別のスレッドで処理します。以下は、アプリケーションの簡略化されたバージョンです。実際のアプリケーションでは、すべての定数は実際には変数であることに注意してください。

以下のコードを実行すると、実行が完了するまでに20 秒以上かかります。

ただし、 malloc/freeを使用するのとは対照的に、 がローカル配列として定義されている場合(コメント行に表示されているように)、実行が完了するまでに1 秒もl_SrcIntegralかかりません。

配列の実際のサイズは (1700 ではなく) 動的であるため、このローカル配列は実際のアプリケーションでは使用できません。

このかなり単純なコードのパフォーマンスを改善する方法についてアドバイスをいただければ幸いです。

#include "cuda_runtime.h"
#include <stdio.h>

#define d_MaxParallelRows 320
#define d_MinTreatedRow   5
#define d_MaxTreatedRow   915
#define d_RowsResolution  1
#define k_ThreadsPerBlock 64

__global__ void myKernel(int Xi_FirstTreatedRow)
{
  int l_ThreadIndex = blockDim.x * blockIdx.x + threadIdx.x;
  if (l_ThreadIndex >= d_MaxParallelRows)
    return;
  int l_Row = Xi_FirstTreatedRow + (l_ThreadIndex * d_RowsResolution);
  if (l_Row <= d_MaxTreatedRow) {

    //float l_SrcIntegral[1700];
    float* l_SrcIntegral = (float*)malloc(1700 * sizeof(float));

    for (int x=185; x<1407; x++) {
      for (int i=0; i<1700; i++)
        l_SrcIntegral[i] = i;
    }

    free(l_SrcIntegral);
  }
}

int main()
{
  cudaError_t cudaStatus;

  cudaStatus = cudaSetDevice(0);

  int l_ThreadsPerBlock = k_ThreadsPerBlock;
  int l_BlocksPerGrid = (d_MaxParallelRows + l_ThreadsPerBlock - 1) / l_ThreadsPerBlock;

  int l_FirstRow = d_MinTreatedRow;
  while (l_FirstRow <= d_MaxTreatedRow) {
    printf("CUDA: FirstRow=%d\n", l_FirstRow);
    fflush(stdout);

    myKernel<<<l_BlocksPerGrid, l_ThreadsPerBlock>>>(l_FirstRow);

    cudaDeviceSynchronize();

    l_FirstRow += (d_MaxParallelRows * d_RowsResolution);
  }

  printf("CUDA: Done\n");

  return 0;
}
4

2 に答える 2

1

1.

@alandが言ったように、各カーネル呼び出しで1行だけを計算すると、パフォーマンスが低下する可能性さえあります。

大規模な並列処理の能力を理論的に使用するには、入力全体を処理することを考える必要があります。

1 つの行を計算するためだけに、320 スレッドだけで複数のカーネルを開始するのはなぜですか? 行があるのと同じ数のブロックを使用し、ブロックごとのスレッドが 1 つの行を処理するようにします。

(ブロックあたり 320 スレッドは適切な選択ではありません。占有率を高める方法を確認してください)

2.

レジスタや共有メモリなどの高速リソースが十分でない場合は、GPGPU プログラミングを使用する基本の 1 つであるタイル アプローチを使用する必要があります。

入力データを同じサイズのタイルに分割し、スレッド内のループで処理します。

ここでは、そのようなタイル アプローチの例を投稿しました。

CUDA での並列化、各列へのスレッドの割り当て

そのタイル アプローチでの範囲チェックに注意してください。

あなたにアイデアを与える例：

任意のサイズの行列の列ベクトルのすべての要素の合計を計算します。

各ブロックは 1 つの列を処理し、そのブロックのスレッドは共有メモリ配列の要素をタイル ループに格納します。終了したら、並列削減を使用して合計を計算し、次の反復を開始します。
最後に、各ブロックはそのベクトルの合計を計算しました。

于 2012-05-08T02:11:47.353 に答える
0

共有メモリを使用して動的配列サイズを引き続き使用できます。<<<...>>>カーネル呼び出しの に3 番目の引数を渡すだけです。これは、ブロックごとの共有メモリのサイズになります。

そこに着いたら、関連するすべてのデータを共有配列に取り込みます（結合アクセスを維持するようにしてください）。スレッドごとに1つまたは複数の（結合アクセスを維持することが適切な場合）要素をもたらします。スレッドが持ち込まれた後にスレッドを同期します (競合状態を停止する必要がある場合のみ、計算が完了する前に配列全体が共有メモリにあることを確認します)。

また、ループではなく、ブロックとスレッドを使用してテッセレーションする必要があります。これはローカル配列を使用した単なる例であることは理解していますが、それでも、ブロック/スレッドをテッセレーションすることは可能で、for ループをネストすることはできません (これはパフォーマンスにとって非常に悪いことです!) 1 つのブロックのみを使用してサンプル コードを実行していることを願っています。それ以外の場合はあまり意味がありません。

于 2012-05-03T14:43:56.993 に答える