0

これは、いくつかの配列を並行して計算するために起動しているカーネルです。

__device__ bool mult(int colsize,int rowsize,int *Aj,int *Bi)
    {       
        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 *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))
                {                                                
                     Cjc[tid+1] += 1;
                     //atomicAdd(Cjc+tid+1,1);           
                }
            }                
        }               
    }

私の起動構成とカーネル呼び出しは次のとおりです。

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);

認めざるを得ないのは、このカーネルの実行速度がかなり遅いということです。アレイをホスト側に戻したらthrust::inclusive_scan、結果のアレイを見つけるために使用します。私の質問は、カーネルの改善/最適化の余地はありますか? 共有メモリを使用しようとしましたが、間違った答えが生成されるか、実行時例外がスローされます。

また、動的に割り当てられた共有メモリ (kernel launch の 3 番目のパラメータによって割り当てられる) は、ブロック間でどのように分散されますか?

ヘルプ/ヒント/ほのめかしをいただければ幸いです。前もって感謝します。

4

1 に答える 1

1

memを使用して割り当てられた共有メモリkernel<<<blocks,threads,mem>>>は、各ブロックに割り当てられたメモリの量です。したがって、各ブロックはmem大量のメモリを取得します。

あなたのコードでは、mult 関数に 2 つの for ループがある理由がわかりません。各スレッドがこれら 2 つの for ループを実行することを指摘しておきます。さらに、kernel関数に for ループもあるということは、各スレッドがformult 関数で 2 つのループを数回実行することを意味します。これは遅いです。さらに、やって

int beg = Bptr[tid]; 
int end = Bptr[tid+1];

完全に結合されたアクセスではありません。合体していないアクセスは遅いです。

于 2012-08-25T13:56:18.637 に答える