0

CUDAプログラミングモデルとその機能を理解しようとしています。演習として、関数呼び出しを使用して次のループ構造を効率的なCUDAカーネルに変換しようとしています。

//function call
bool gmul(int rowsize,int *Ai,int *Bj,int colsize)
{
    for(int i = 0;i < rowsize;i++)
    {
        for(int j = 0;j < colsize;j++)
        {
            if(Ai[i] == Bj[j])
            {
                return true;
            }
        }
    }
    return false;
}

//Some for loop in main function is as follows

for(i = 0;i < q ;i++)
    {
        cbeg = Bjc[i];
        cend = Bjc[i+1];        
        for(j = 0;j < m;j++)
        {
            beg = Aptr[j];
            end = Aptr[j+1];            
            if(gmul(end - beg,Acol + beg,Bir + cbeg,cend - cbeg))
            {   
                temp++;             
            }                       
        }
        Cjc1[i+1] = temp ;              
    }

そして、関数呼び出しのある私のカーネルは次のとおりです。

    __device__ bool mult(int colsize,int rowsize,int *Aj,int *Bi,int *val)
    {       
        for(int j = 0; j < rowsize;j++)
        {           
           for(int k = 0;k < colsize;k++)
            {   
              if(Aj[j] == Bi[k])
               {    
                return true;
                }                               
            }           
        }
            return false;       
    }


__global__ void kernel(int *Aptr,int *Aj,int *Bptr,int *Bi,int rows,int cols,int *count,int *Cjc)
    {
        int tid = threadIdx.x + blockIdx.x * blockDim.x;
        int i;
        if(tid < cols)
        {
            int beg = Bptr[tid];
            int end = Bptr[tid+1];
            for(i = 0;i < rows;i++)
            {
                int cbeg = Aptr[i];
                int cend = Aptr[i+1];
                if(mult(end - beg,cend - cbeg,Aj+cbeg,Bi+beg,count))
                {
                    //atomicAdd(count,1);
                                    //Changes made are in next line
                              atomicAdd(Cjc+tid+1,1);           
                }
            }
            //atomicAdd(Cjc+tid+1,*count);              
        }               
    }

私が欲しいのは、値__device__ multとともに返されるときはいつでもtrue、私のグローバルカーネル関数はその特定のスレッドのカウンターをインクリメントし、forループ（カーネル関数内）が終了すると、値をCjc配列に格納し、カウントが他のスレッドに渡されることですインクリメント操作用。しかし、期待値が得られません。このCjc配列で取得するのは、すべてのスレッドの実行が終了した後の最終カウントだけです。

CC2.0でGTX480を使用しています

なぜこのCUDAカーネルに対して間違った答えや最適化が得られるのかについての提案/ヒントをいただければ幸いです。前もって感謝します。 * *******解決済み**** ** * **** _ _ _

現在、4000以上のサイズに達すると、配列内のすべての要素の値がとして取得されるという問題に直面しています0。これが私がカーネルを起動する方法です。

    int numBlocks,numThreads;

        if(q % 32 == 0)
        {
            numBlocks = q/32;
            numThreads = 32;
        }
        else
        {
            numBlocks = (q+31)/32;
            numThreads = 32;
        }
findkernel<<<numBlocks,numThreads>>>(devAptr,devAcol,devBjc,devBir,m,q,d_Cjc);

ブロックまたはグリッドの寸法の制限を超えているのではないかと思っていましたが、CC 2.0の場合、制限を超えない十分なブロックとスレッドを起動するのが適切だと思います。なぜまだすべての答えがとして出ているのだろうか0。

4

1 に答える 1

1

count同期なしで読み取りと書き込みを行う並列スレッドを作成しました。スレッドは予測できない順序で同時に実行されるため、スレッドはアトミックに変更され、予測できない順序countで読み取らcountれます。式*countは、正確な実行順序に応じて異なる結果を生成します。

forループ（カーネル関数内）が終了すると、値がCjc配列に格納され、カウントがインクリメント操作のために他のスレッドに渡されます。

同期がないため、別のスレッドがループを終了するのを待つスレッドはありません。すべてのスレッドに同じストレージを共有させる代わりに、count各スレッドに異なるストレージを与えてみませんか？そうすれば、スレッドはお互いの結果に影響を与えません。この後、スキャンカーネルを実行して、結果を組み合わせることができます。

于 2012-07-30T04:54:53.797 に答える