使用 const global <type>* 限制参数与 OpenCL 卤化？

Question

在 OpenCL 中，当我们将输入参数指定为 const global * restrict 时，我们将获得有效的输入参数硬件路径，如（对于一段手写的 OpenCL 代码）：

__kernel void oclConvolveGlobalMem(const global   float* restrict input,
                                         constant float* restrict filterWeights,
                                         global   float* restrict output)

然而，正如 HL_DEBUG_CODEGEN=1 卤化物生成的那样：

// Address spaces for kernel_conv_70_s0_y___block_id_y
#define __address_space__conv__70 __global
#define __address_space__input __global
#define __address_space__kernel __global
__kernel void kernel_conv_70_s0_y___block_id_y(
 const int _conv__70_extent_0,
 const int _conv__70_extent_1,
 const int _conv__70_min_0,
 const int _conv__70_min_1,
 const int _conv__70_stride_1,
 const int _input_min_0,
 const int _input_min_1,
 const int _input_stride_1,
 const int _kernel_min_0,
 const int _kernel_min_1,
 const int _kernel_stride_1,
 __address_space__conv__70 float *_conv__70,
 __address_space__input const float *_input,
 __address_space__kernel const float *_kernel,
 __address_space___shared int16* __shared)

未声明输入参数的地方restrict。我希望这会真正限制性能。我确实让 Halide 添加了指针受限的概念（它们使用的缓冲区没有别名。）

Answer 1

您上次更新 Halide 是什么时候？ Halide 最近（2016 年 10 月）对缓冲区参数添加了限制：https://github.com/halide/Halide/pull/1550。最新的二进制版本确实有这个变化，几乎没有。

使用 const global <type>* 限制参数与 OpenCL 卤化？

Halide with OpenCL using const global <type>* restrict arguments?

halide