0

フラット化された 2D 配列を受信するカーネルがあり、共有メモリを使用するたびに配列の 1 行をコピーしたいと考えています。私のカーネルは次のようになります。

__global__ void searchKMP(char *test,size_t pitch_test,int ittNbr){
    int  tid = blockDim.x * blockIdx.x + threadIdx.x;
    int strideId = tid * 50;

    int m = 50;

    __shared__ char s_test[m];

    int j;
               //this loops over the number of lines in my 2D array           
                   for(int k=0; k<ittNbr; k++){

                   //this loops to store my flattened (basically threats 1 line at a time) array into shared memory     
                   if(threadIdx.x==0){
                     for(int n =0; n<50; ++n){
                    s_test[n] = *(((char*)test + k * pitch_test) + n);

                }
             }
            __syncthreads();


             j=0;

            //this is loop to process my shared memory array against another 1D array
             for(int i=strideID; i<(strideID+50); i++{
             ...dosomething...
             (increment x if a condition is met) 
             ...dosomething...
             }
             __syncthreads();
             if(x!=0)
                cache[0]+=x;

            ...dosomething...

}

ただし、x の値を確認すると、x の値は常に変化するか、スレッドの数によって変化します。たとえば、250 スレッドの 20 ブロックが実行に応じて値 7 または 6 を返す場合、500 スレッドの 10 ブロックは 9 を返します。問題が共有メモリにコピーされた 2D 平坦化配列にあるのか、それともこのコードで何かが間違っているのか疑問に思います。

4

2 に答える 2

1

共有メモリ内の配列には 20 個の要素があるようです:

   int m = 20;
   __shared__ char s_test[m];

しかし、内側のループでは、50 個の要素を書き込もうとしています。

   for(int n =0; n<50; ++n){
      s_test[n] = *(((char*)test + k * pitch_test) + n);

これがあなたが探していた問題であるかどうかはわかりませんが、うまくいかないようです。

于 2013-03-22T20:08:09.280 に答える
0

共有メモリは、同じブロック内のすべてのスレッドで共有されます

なぜ共有メモリが必要なのか、何をしているのかはあまり明確ではありません。

コードでは、ブロック内のすべてのスレッドが同じ値を共有メモリに何度も書き込みますが、冗長です

共有メモリを操作する一般的な方法は、次のようなものです。

if(threadIdx.x < m)
  s_test[threadIdx.x] = *(global_mem_pointer + threadIdx.x);

__syncthreads();

ブロック内のすべてのスレッドは、「同時に」独自の値を書き込み、__syncthreads();メモリが必要なものでいっぱいになり、ブロック内のすべてのスレッドに表示されます

于 2013-03-22T13:47:42.550 に答える