1

RGBA画像をグレースケールに変換するcudaチュートリアルに取り組んでいます。blockSizeしかし、 と を変更するとgridSizeX33 の時間が改善される理由がわかりませんでした。

__global__
void rgba_to_greyscale(const uchar4* const rgbaImage,
                       unsigned char* const greyImage,
                       int numRows, int numCols)
{
    int i = blockIdx.x*numCols + threadIdx.x;
    float channelSum = .299f * rgbaImage[i].x + .587f * rgbaImage[i].y + .114f * rgbaImage[i].z;
    greyImage[i]= channelSum;
}

void your_rgba_to_greyscale(const uchar4 * const h_rgbaImage, uchar4 * const d_rgbaImage,
                            unsigned char* const d_greyImage, size_t numRows, size_t numCols)
{
  const dim3 blockSize(numCols, 1, 1);
  const dim3 gridSize(numRows, 1 , 1);
  rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols);

  cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError());
}

上記のように設定すると：

const dim3 blockSize(numCols, 1, 1);
const dim3 gridSize(numRows, 1 , 1);

私は得るYour code executed in 0.030304 ms

私が設定したとき：

 const dim3 blockSize(1, 1, 1);
 const dim3 gridSize(numRows, numCols , 1);

新しいインデックスで動作するようにスレッド関数を更新します。

int i = blockIdx.x*numCols + blockIdx.y;

私は得るYour code executed in 0.995456 ms。

  1. GPUは2番目のグリッド分割ですべてのピクセルを個別に計算できるため、逆になると思いますキャッシュの一貫性の問題に関連していますか? なぜこれらの結果が得られるのですか?
  2. 理論上、この問題のグリッドとブロックのサイズはどれくらいが最適ですか? 実行時に計算することは可能ですか？

ご参考までに：

numRows = 313 numCols =557

技術的特性:

#uname -a && /usr/bin/nvidia-settings -v
    Linux ip-10-16-23-92 3.2.0-39-virtual #62-Ubuntu SMP Thu Feb 28 00:48:27 UTC 2013 x86_64 x86_64 x86_64 GNU/Linux

    nvidia-settings:  version 304.54  (buildmeister@swio-display-x86-rhel47-11)
4

1 に答える 1

6

グリッド/ブロック構成のどちらも推奨されません。最初のものはスケーラブルではありません。これは、ブロックあたりのスレッド数が GPU に対して制限されているためです。そのため、大きな画像サイズでは最終的に失敗します。GPU の占有率が非常に低くなるため、ブロックごとに 1 つのスレッドしかないため、2 番目のスレッドは推奨されません。CUDA Toolkit に含まれているGPU Occupancy Calculatorで確認できます。推奨されるブロック サイズは、GPU に応じて、GPU ワープ サイズの倍数 (16 または 32) にする必要があります。

あなたの場合の2Dグリッドとブロックサイズの一般的でスケーラブルなアプローチは次のようになります:

const dim3 blockSize(16, 16, 1);
const dim3 gridSize((numCols + blockSize.x - 1)/blockSize.x, (numRows + blockSize.y - 1)/blockSize.y , 1);

デバイスの制限内であれば、ブロック サイズを 16 x 16 から任意のサイズに変更できます。計算機能が 1.0 から 1.3 のデバイスでは、ブロックごとに最大 512 のスレッドが許可されます。コンピューティング機能 2.0 以降のデバイスの場合、この制限はブロックあたり 1024 スレッドです。

現在、グリッドとブロックは 2 次元であるため、カーネル内のインデックスは次のように変更されます。

int i = blockIdx.x * blockDim.x + threadIdx.x; //Column
int j = blockIdx.y * blockDim.y + threadIdx.y; //Row

int idx = j * numCols + i;

//Don't forget to perform bound checks
if(i>=numCols || j>=numRows) return;

float channelSum = .299f * rgbaImage[idx].x + .587f * rgbaImage[idx].y + .114f *     rgbaImage[idx].z;
greyImage[idx]= channelSum;
于 2013-06-15T14:19:36.240 に答える