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

Question

MxNの2D行列から1Dの配列を変換し、各列を並列化していくつかの操作を行いたいとします。各列にスレッドを割り当てるにはどうすればよいですか？私は3x3の行列がある場合

たとえば、：

1 2 3 

4 5 6 

7 8 9 

をそして、私は列番号に応じて、列の各番号を追加したい（その1列目は1が追加されます、第二には、2を追加します。.. ..）、それは次のようになります。

1+1 2+1 3+1 

4+2 5+2 6+2 

7+3 8+3 9+3 

CUDAでこれを行うにはどうすればよいですか？私は配列のすべての要素にスレッドを割り当てる方法を知っていますが、各列にスレッドを割り当てる方法はわかりません。だから私が望むのは、各列（1、2、3）（4,5,6）（7,8,9）を送信して操作を行うことです。

Answer 1

例では、行に基づいて数値を追加しています。それでも行列の行/列の長さは分かります（MxNだと分かります）。

matrix[ gid ] += my_col_number_function(gid%N); 

Answer 2

これらの剰余演算を避けるために、より良いグリッドレイアウトを使用する：あなたは別の番号を追加したい場合は、あなたのような何かができる

__global__ void MyAddingKernel(int* matrix, int M, int N) 
{ 

    int gid = threadIdx.x + blockDim.x*blockIdx.x; 
    //Let's add the row number to each element 
    matrix[ gid ] += gid % M; 
    //Let's add the column number to each element 
    matrix[ gid ] += gid % N; 

} 

：何あなたができることのようなものです。

最新のCudaで64ビット範囲の行に一意のブロック索引を使用します。

スレッドがすべての要素をループして反復し、ユニークなスレッドインデックスを追加します。

計算されたデータがブロック（行）全体にわたって一意である場合、特に複雑な計算の場合、入力データのタイル化は一般的なアプローチです。

/* 
* @param tileCount 
*/ 
__global__ void addRowNumberToCells(int* inOutMat_g, 
    const unsigned long long int inColumnCount_s, 
    const int inTileCount_s) 
{ 

    //get unique block index 
    const unsigned long long int blockId = blockIdx.x //1D 
     + blockIdx.y * gridDim.x //2D 
     + gridDim.x * gridDim.y * blockIdx.z; //3D 

    /* 
    * check column ranges in case kernel is called 
    * with more blocks then columns 
    * (since its block wide following syncthreads are safe) 
    */ 
    if(blockId >= inColumnCount_s) 
     return; 

    //get unique thread index 
    const unsigned long long int threadId = blockId * blockDim.x + threadIdx.x; 

    /* 
    * calculate unique and 1 blockId 
    * maybe shared memory is overhead 
    * but it shows concept if calculation is more complex 
    */ 
    __shared__ unsigned long long int blockIdAnd1_s; 
    if(threadIdx.x == 0) 
     blockIdAnd1_s = blockId + 1; 
    __sycnthreads(); 


    unsigned long long int idx; 

    //loop over tiles 
    for(int i = 0; i < inTileCount_s) 
    { 
     //calculate new offset for sequence thread writes 
     idx = i * blockDim.x + threadIdx.x; 
     //check new index range in case column count is no multiple of blockDim.x 
     if(idx >= inColumnCount_s) 
      break; 
     inOutMat_g[idx] = blockIdAnd1_s; 
    } 

} 

例クーダ2.0：

マット[131000] [1000]

必要blockCount blockDim.y用65535分の131000 = 2切り上げ=！

inTileCount_s = 1000/192 = 6切り上げ！

（クーダ2.0に= 100占有ブロック当たり192スレッド）

< <（65535、2、1）、（192、1、1）>> addRowNumberToCells（マット、1000 6）