0

cuda プログラムを実行したいのですが、初心者です。ヒストグラム用のプログラムを作成する必要があります。しかし、バケツで。maxValue (例では 40) に応じて、数値が適切なバケットに追加されます。4 つのバケットがある場合:

ヒスト: | 1 | 10 | 30 | 39 | 32 | 2 | 4 | 5 | 1 |

0 ～ 9 (1 番目のバケット)

10-19 (2 番目のバケット)

20-29(3番目のバケット)

30～39(第4バケット)

私の GPU にはCompute Capability 1.1 があります。

各スレッドが自分の一時テーブルに自分の値を追加しているブロックの共有 temp[] を持つようなことをしようとしていました:

__global__ void histo_kernel_optimized5( unsigned char *buffer, long size,
                               unsigned int *histo )
{
     extern __shared__ unsigned int temp[];
     temp[threadIdx.x] = 0;
     __syncthreads();

     int i = threadIdx.x + blockIdx.x * blockDim.x;
     int offset = blockDim.x * gridDim.x;
     int bucketID;
     while (i < size)
     {
              bucketID = array[i]/Bwidth;
              atomicAdd( &temp[bucketID], 1);
              i += offset;
     }
     __syncthreads();


    atomicAdd( &(histo[threadIdx.x]), temp[threadIdx.x] );
}

histo_kernel_optimized <<<array_size/buckets, buckets,buckets*sizeof(unsigned int)>>>(buffer,SIZE, histogram)

しかし、コンパイラは次のように言います: 命令 '{atom,red}.shared' は を必要とします。対象sm_12以上

また、作成されたスレッドごとに一時テーブルを作成しようとしました:

__global__ void histo_kernel_optimized5( unsigned char *buffer, long size,
                               unsigned int *histo )
{
    unsigned int temp[buckets];
     int j;
    for (j=0;j<buckets;j++){
        temp[j]=0;
    }

    int bucketID;

    int i = threadIdx.x + blockIdx.x * blockDim.x;
    int offset = blockDim.x * gridDim.x;
    while (i < size)
    {
        bucketID = array[i]/Bwidth;
        temp[bucketID]++;
        i += offset;
    }


    for (j=0;j<buckets;j++){
        histo[j] += temp[j];    
    }
 }

しかし、一時テーブルを作成するには定数が必要なため、コンパイラはそれを許可しません。しかし問題は、バケットがコマンドラインに対して動的に与えられることです。

それを行う別の方法はありますか？やり方がわかりません。私は混乱しています。

4

3 に答える 3

8

アトミックを使用する場合、より少ないブロック間で調整する必要がないため、より少ないブロックを起動すると競合が減少します (したがって、パフォーマンスが向上します)。起動するブロックの数を減らし、各ブロックがより多くの入力要素をループするようにします。

for (unsigned tid = blockIdx.x*blockDim.x+threadIdx.x; 
              tid < size; tid += gridDim.x*blockDim.x) {
    unsigned char value = array[tid]; // borrowing notation from another answer here
    int bin = value % buckets;
    atomicAdd(&histo[bin],1);
}
于 2013-04-04T16:06:58.273 に答える
0

アトミック操作のないデバイス用のソリューションがあり、CUDA のヒストグラム計算でPodlozhnyuk によって提案されたワープへの細分化を使用して、オンチップ メモリの衝突を最小限に抑えるアプローチを示しています。

コードは CUDASamples\3_Imaging\histogram にあります (CUDA サンプルから)

于 2016-05-29T16:25:54.090 に答える