Cudaカーネルの時間実行を最適化する

Question

私はCudaの学習者です。カーネル関数の実行時間を最適化したいと思います。その結果、私は2つの写真の違いを計算する短いプログラムを実現しました。だから私はCで古典的なCPUの実行、およびクーダC.

におけるGPUの実行間の実行時間を比較してここでは、私が話しているコードを見つけることができます。

int *imgresult_data = (int *) malloc(width*height*sizeof(int)); 
int size = width*height; 

switch(computing_type) 
{ 

    case GPU: 

    HANDLE_ERROR(cudaMalloc((void**)&dev_data1, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data2, size*sizeof(unsigned char))); 
    HANDLE_ERROR(cudaMalloc((void**)&dev_data_res, size*sizeof(int))); 

    HANDLE_ERROR(cudaMemcpy(dev_data1, img1_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data2, img2_data, size*sizeof(unsigned char), cudaMemcpyHostToDevice)); 
    HANDLE_ERROR(cudaMemcpy(dev_data_res, imgresult_data, size*sizeof(int), cudaMemcpyHostToDevice)); 

    float time; 
    cudaEvent_t start, stop; 

    HANDLE_ERROR(cudaEventCreate(&start)); 
    HANDLE_ERROR(cudaEventCreate(&stop)); 
    HANDLE_ERROR(cudaEventRecord(start, 0)); 

    for(int m = 0; m < nb_loops ; m++) 
    { 
     diff<<<height, width>>>(dev_data1, dev_data2, dev_data_res); 
    } 

    HANDLE_ERROR(cudaEventRecord(stop, 0)); 
    HANDLE_ERROR(cudaEventSynchronize(stop)); 
    HANDLE_ERROR(cudaEventElapsedTime(&time, start, stop)); 

    HANDLE_ERROR(cudaMemcpy(imgresult_data, dev_data_res, size*sizeof(int), cudaMemcpyDeviceToHost)); 

    printf("Time to generate: %4.4f ms \n", time/nb_loops); 

    break; 

    case CPU: 

    clock_t begin = clock(), diff; 

    for (int z=0; z<nb_loops; z++) 
    { 
     // Apply the difference between 2 images 
     for (int i = 0; i < height; i++) 
     { 
      tmp = i*imgresult_pitch; 
      for (int j = 0; j < width; j++) 
      { 
       imgresult_data[j + tmp] = (int) img2_data[j + tmp] - (int) img1_data[j + tmp]; 
      } 
     } 
    } 
    diff = clock() - begin; 

    float msec = diff*1000/CLOCKS_PER_SEC; 
    msec = msec/nb_loops; 
    printf("Time taken %4.4f milliseconds", msec); 

    break; 
} 

そして、ここでは私のカーネルです機能：私は各1

• CPUのためにこれらの実行時間を得

  __global__ void diff(unsigned char *data1 ,unsigned char *data2, int *data_res) 
  { 
      int row = blockIdx.x; 
      int col = threadIdx.x; 
      int v = col + row*blockDim.x; 
  
      if (row < MAX_H && col < MAX_W) 
      { 
       data_res[v] = (int) data2[v] - (int) data1[v]; 
      } 
  } 
  

  ：1,3210ms
• GPU：GPUの結果は、それがあるべきように、下のようにできない理由0,3229ms

私は疑問に思います。私はCudaの初心者ですので、いくつかの古典的なエラーがある場合包括してください。

EDIT1： フィードバックいただきありがとうございます。私はカーネルから 'if'条件を削除しようとしましたが、プログラムの実行時間はそれほど変わっていませんでした。

しかし、Cudaプロファイラをインストールした後、スレッドは同時に実行されていないと言いました。なぜ私はこのようなメッセージを持っているのか理解できませんが、CPUよりもGPUの方が5倍または6倍高速です。この比率は、各スレッドが1つのピクセルを他のすべてのスレッドと同時に処理することになっているため、より大きくなければなりません。私が間違っていることを知っていれば、それはひどいでしょう...

フロー。

Answer 1

おそらくコードには他にも問題がありますが、ここには私の見ているものがあります。 __global__ void diffの次の行には最適ではないと考えられている。

if (row < MAX_H && col < MAX_W) 
{ 
    data_res[v] = (int) data2[v] - (int) data1[v]; 
} 

条件演算子をワープ発散内のカーネル結果内。これは、ワープ内のifとelse部分が、並列ではなく順番に実行されることを意味します。また、実現したように、ifは、境界でのみfalseと評価されます。 row < MAX_H && col < MAX_Wは常にtrueある

1. 中央部分：発散や不要な計算を回避するために、二つの部分に自分の画像を分割します。この領域に追加のカーネルを作成します。ここではifは不要です。

2. diffカーネルを使用するボーダーエリア。

明らかに、カーネルを呼び出すコードを変更するでしょう。

---

、別のノートで：

1. GPUは、スループット指向のアーキテクチャを採用していますが、CPUとして、待ち時間指向ではありません。これは、少量のデータを処理する場合、CPUがCUDAよりも高速である可能性があることを意味します。大きなデータセットを試しましたか？

2. CUDAプロファイラは、コードに最適ではないことを示す非常に便利なツールです。

Answer 2

時間を正しく測定しているとは思わないが、メモリコピーはGPUで時間がかかるので、時間を測定する際に考慮する必要があります。

は、私はあなたがテストすることができますいくつかの詳細を参照してください。

1. 私はあなたが定数としてMAX_HとMAX_Hを使用していると仮定し、あなたは（）cudaMemcpyToSymbolを使用してそうすることを検討してください。

2. __syncthreads（）を使用してスレッドを同期することを忘れないでください。ループの繰り返しごとに問題は発生しません。

3. CUDAはワープで動作するため、ブロックあたりのブロック数とスレッド数はハードウェアでサポートされていない限り、ブロック数が8の倍数で、ブロックあたり512スレッド以下であると効果的です。ブロックあたり128スレッドを使用する例を次に示します。< < <（cols * rows + 127）/ 128,128 >>>

4. 割り当てられたメモリをGPUで解放し、作成されたタイムイベントを破棄することも忘れないでください。

5. カーネル関数では、1つの変数int v = threadIdx.x + blockIdx.x * blockDim.xを持つことができます。

6. 結果が正しいことを実行時間の横にテストしましたか？私はあなたがパディングのために配列を操作している間、あなたはcudaMallocPitch（）とcudaMemcpy2D（）を使うべきだと思います。