0

OpenCL カーネルのフィルタリングのために、3D で多くの作業を開始しています。2D または 3D サブセットをグローバル メモリからローカルまたはプライベート メモリにコピーする最適な方法はありますか?

これは、3D データセットを取得して 3D カーネルを適用する (または 3D カーネルが占めるスペースを操作する) ために使用できます。各スレッドは 1 つのピクセルを見て、カーネルのサイズ (1、3、5 など) である 3 次元でピクセルの周囲のデータをトリミングし、このデータのサブセットをローカルまたはプライベート メモリにコピーしてから、たとえば、データのサブセットの標準偏差を計算します。

最も簡単で最も効率の悪い方法は、力ずくで行うことです。

__kernel void Filter_3D_StdDev(__global float *Data_3D_In,
                               int KernelSize){
//Note: KernelSize is always ODD

int k = get_global_id(0); //also z
int j = get_global_id(1); //also y
int i = get_global_id(2); //also x

//Convert 3D to 1D
int linear_coord = i + get_global_size(0)*j + get_global_size(0)*get_global_size(1)*k;

//private memory
float Subset[KernelSize*KernelSize*KernelSize];

int HalfKernel = (KernelSize - 1)/2; //compute the pixel radius

for(int z = -HalfKernel ; z < HalfKernel; z++){
     for(int y = -HalfKernel ; y < HalfKernel; y++){
          for(int x = -HalfKernel ; z < HalfKernel; x++){
               int index = (i + x) + get_global_size(0)*(j + y) + \            
                                  get_global_size(0)*get_global_size(1)*(k + z);
               Subset[x + HalfKernel + (y + HalfKernel)*KernelSize + (z + HalfKernel)*KernelSize*KernelSize] = Data_3D_In[index];
          }

     }
}

//Filter subset here

}

非常に多くの呼び出しがグローバル メモリに対して行われるため、これは恐ろしく非効率的です。これを改善する方法はありますか？

私の最初の考えはvload、次のようなループの数を減らすために使用することです。

__kernel void Filter_3D_StdDev(__global float *Data_3D_In,
                               int KernelSize){
//Note: KernelSize is always ODD

int k = get_global_id(0); //also z
int j = get_global_id(1); //also y
int i = get_global_id(2); //also x

//Convert 3D to 1D
int linear_coord = i + get_global_size(0)*j + get_global_size(0)*get_global_size(1)*k;

//private memory
float Subset[KernelSize*KernelSize];

int HalfKernel = (KernelSize - 1)/2; //compute the pixel radius

for(int z = -HalfKernel ; z < HalfKernel; z++){
     for(int y = -HalfKernel ; y < HalfKernel; y++){
          //##TODO##
          //Automatically determine which vload to use based on Kernel Size
          //for now, use vload3
               int index = (i + -HalfKernel) + get_global_size(0)*(j + y) + \            
                                  get_global_size(0)*get_global_size(1)*(k + z);
               int subset_index = (z + HalfKernel)*KernelSize*KernelSize
               float3 temp = vload3(index, Data_3D_In);
               vstore3(temp, subset_index, Subset);

     }
}

//Filter subset here

}

もっと良い方法はありますか？

前もって感謝します！

4

1 に答える 1

1

まず、これらのループを展開する必要があります。関数のコピーをいくつか作成するか、コンパイルする前に文字列を置換するか、最初にループを展開する必要がありますが、テストと同じように:

#define HALF_KERNEL_SIZE = 2
#pragma unroll HALF_KERNEL_SIZE * 2 + 1
for(int z = -HALF_KERNEL_SIZE ; z < HALF_KERNEL_SIZE ; z++){
    #pragma unroll HALF_KERNEL_SIZE * 2 + 1
    for(int y = -HALF_KERNEL_SIZE ; y < HALF_KERNEL_SIZE ; y++){

GPU の場合、ローカル メモリに読み込む必要があります (特に 5x5x5 の場合は、既にデータがあり、それを取得するために戻りたくない場合に、グローバル メモリに大量に読み戻すためです。(これは、 GPU) CPU の場合、それほど大きな問題ではありません。

したがって、これを畳み込みの場合とまったく同じように行いますが、追加の次元を使用します。

1. Read in a block (or cube) of memory into local memory for a number of threads.
2. Create a barrier to make sure all data is read before you continue.
3. Sample into your local memory using your local id as an offset.
4. Test various local workgroup sizes until you get best performance

他のすべては同じです。より大きなオーバーラップを持つより大きなカーネルの場合、これは manatudes の桁数より高速になります。

于 2013-08-09T23:24:56.133 に答える