0

我开始为我的 OpenCL 内核做大量的 3D 工作以进行过滤。是否有将 2D 或 3D 子集从全局内存复制到本地或私有内存的最佳方法？

其用途可能是获取 3D 数据集并应用 3D 内核（或对 3D 内核占用的空间进行操作）。每个线程将查看一个像素，将像素周围的数据裁剪为 3 维，即内核大小（例如 1、3、5 等），将此数据子集复制到本地或私有内存，然后例如，计算数据子集的标准偏差。

最简单且效率最低的方法是蛮力：

__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 回答