cudaカーネルのブロックとスレッドの数と使い方

Question

2次元行列の値の平均を取る単純なカーネルと思ったことをどうやって作ろうとしているのですが、私は思考過程を真っ直ぐにするいくつかの問題を抱えています。

私のdeviceQuery出力によると、私のGPUは16MP、32cores/mp、最大ブロックは1024x1024x64で、最大スレッド/ブロック= 1024があります。

私はいくつかの大きな画像を処理しています。多分5000px x 3500pxかそのようなものです。私のカーネルの1つは、画像内のすべてのピクセルにわたって平均値をとっています。

既存のコードには、2D配列[rows] [cols]として保存されたイメージがあります。だからカーネルは、Cの中で、あなたが期待していたように見えます。中間の計算では、行の上のループと列の上のループです。

CUDAでこのコードの寸法計算部分を設定するにはどうすればよいですか？私はSDKのリダクションコードを見てきましたが、これは1次元配列のものです。それはsoemthing 2Dを持っているときのブロックとスレッドの数を設定する方法の任意の言及を持っていません。これは、セットアップのために意味をなすように見えるん

num_threads=1024; 
blocksX = num_cols/sqrt(num_threads); 
blocksY = num_rows/sqrt(num_threads); 
num_blocks = (num_rows*num_cols)/(blocksX*blocksY); 

dim3 dimBlock(blocksX, blocksY, 1); 
dim3 dimGrid(num_blocks, 1, 1); 

：私は実際にそうようにそれを設定する必要があるだろう考えていて、これは私がチャイムに誰か、助けが欲しいところです

？

そしてカーネル内で、特定の行または列上で動作するように、私は

rowidx =（blockIdx.x * blockDim.x）+ threadId.x colidx =（blockIdxを使用する必要があるだろう。 y * blockDim.y）+ threadId.y

少なくとも、私はそれが行と列を取得するために働くと思う。

カーネルの特定の行rと列cにどうすればアクセスできますか？ CUDAプログラミングガイドでは、私は次のコードが見つかりました：

// Host code int width = 64, height = 64; 
float* devPtr; size_t pitch; 
cudaMallocPitch(&devPtr, &pitch, width * sizeof(float), height); 
MyKernel<<<100, 512>>>(devPtr, pitch, width, height); 
// Device code __global__ void MyKernel(float* devPtr, size_t pitch, int width, int height) 
{ 
for (int r = 0; r < height; ++r) 
{ 
float* row = (float*)((char*)devPtr + r * pitch); 
for (int c = 0; c < width; ++c) 
{ 
float element = row[c]; 
} 
} 
} 

あなたは2次元配列を宣言するためにC言語でmalloc関数を使用したい方法に似ていますが、それは自分自身のカーネルでその配列にアクセスする一切の言及を持っているdoesntの。私のコードでは、私はそのcudaMallocPitch呼び出しを使用し、デバイス上の2D配列に私のデータを取得するmemcpyを実行するでしょうか？

ヒントありがとうございます！ありがとう！

Answer 1

以下は、自分のコードの簡単なカーネルを使った短いスニペットです。 floatポインタはすべてデバイスポインタです。これが参考になることを願っています。

定義し、ヘルプ機能：

#define BLOCK_SIZE 16 

int iDivUp(int a, int b){ 
    return (a % b != 0) ? (a/b + 1) : (a/b); 
} 

ブロックサイズの計算：

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
dim3 dimGridProj(iDivUp(width,BLOCK_SIZE), iDivUp(height,BLOCK_SIZE)); 

ホストコール：

calc_residual<<<dimGridProj, dimBlock>>>(d_image1, d_proj1, d_raynorm1, d_resid1, width, height); 

カーネル：

__global__ void calc_residual(float *d_imagep, float *d_projp, float *d_raysump, float *d_residualp, int width, int height) 
{ 
    int iy = blockDim.y * blockIdx.y + threadIdx.y; 
if (iy >= height) { 
    return; 
} 
int ix = blockDim.x * blockIdx.x + threadIdx.x; 
if (ix >= width) { 
    return; 
} 
int idx = iy * width + ix; 
float raysumv = d_raysump[idx]; 
if (raysumv > 0.001) { 
    d_residualp[idx] = (d_projp[idx]-d_imagep[idx])/raysumv; 
} 
else{ 
    d_residualp[idx] = 0; 
} 
} 

Answer 2

このようなパフォーマンスアプリケーションでは、2D行列情報をメモリ内に単一の配列として格納する必要があります。だからM×Nの行列があれば、それを長さM * Nの単一の配列に格納することができます。

あなたは2x2の行列

(1 , 2) 
(3 , 4) 

を保存したいのであれば、あなたは、単一の配列を作成し、AN行の要素Iを初期化し、以下を使用して列j。

int rows=2; 
int cols=2; 
float* matrix = malloc(sizeof(float)*rows*cols); 
matrix[i*cols+j]=yourValue; 
//element 0,0 
matrix[0*cols+0]=1.0; 
//element 0,1 
matrix[0*cols+1]=2.0; 
//element 1,0 
matrix[1*cols+0]=3.0; 
//element 1,1 
matrix[1*cols+1]=4.0; 

2Dアレイを取得し、それを行優先順序でデータを記憶する呼ばれ、このようにメモリの単一の連続片を格納するこの方法。ウィキペディアの記事hereを参照してください。データのレイアウトをこのような形式に変更すると、SDKに表示された縮小を使用することができ、GPUカーネルコードでより多くの合体読み取りを実行できるようになります。

Answer 3

最近、私はこの質問を次のように考えました。

// Grid and block size 
const dim3 blockSize(16,16,1); 
const dim3 gridSize(numRows, numCols, 1); 
// kernel call 
rgba_to_greyscale<<<gridSize, blockSize>>>(d_rgbaImage, d_greyImage, numRows, numCols 

ここでブロック

あたりのブロック
のグリッド寸法=数 ブロックサイズ=スレッドは、対応するカーネル

__global__ void rgba_to_greyscale(const uchar4* const rgbaImage, 
         unsigned char* const greyImage, 
         int numRows, int numCols) 
{ 
    int idx = blockIdx.x + blockIdx.y * numRows; 
    uchar4 pixel  = rgbaImage[idx]; 
    float intensity = 0.299f * pixel.x + 0.587f * pixel.y + 0.114f * pixel.z; 
    greyImage[idx] = static_cast<unsigned char>(intensity); 
} 

幸運です！