如何在 OpenCL 中使用 __local 变量？答案

【问题标题】：How to use __local variables in OpenCL?如何在 OpenCL 中使用 __local 变量？
【发布时间】：2019-05-20 15:23:47
【问题描述】：

我正在尝试在 OpenCL 中实现卷积算法（使用 Vivado HLS）。我正在尝试在执行工作组之前将部分图像加载到本地内存中（例如，如果工作组为 128*128，卷积滤波器为 5*5，我将加载 132*132 像素）。如何编写内核，使本地内存在工作组启动时只加载一次？

伪代码：

#define WKGRP_W 128
#define WKGRP_H 128

#define FILTER_SIZE 5

#define BUFFER_W WKGRP_W+FILTER_SIZE-1
#define BUFFER_H WKGRP_H+FILTER_SIZE-1


__kernel void  __attribute__ ((reqd_work_group_size(WKGRP_W, WKGRP_H, 1)))
convolve(
    const __global data_t* input,
    __global data_t* output,
    __constant data_t* filter_params
){
    __local data_t img_buffer[BUFFER_H][BUFFER_W];
    __local data_t output_buffer[WKGRP_H][WKGRP_W];

    /**
     * if (the workgroup is starting) {
     *     load data from input into img_buffer
     * }
     */

    filter(img_buffer, filter_params, get_local_id(0), get_local_id(1), output_buffer);

    /**
     * if (the workgroup is finished) {
     *     load data from output_buffer into output
     * }
     */
}

【问题讨论】：

标签： opencl vivado-hls

【解决方案1】：

假设每个输入像素有一个工作项，您可以让每个线程将一个像素从全局内存加载到本地内存。

int x = get_local_id(0);
int y = get_local_id(1);
img_buffer[x][y] = input[...];
barrier(CLK_LOCAL_MEM_FENCE);

// filter here or whatever computation you need to perform

如果您的数据是对齐的，您可以通过转换为更大的数据类型来加快速度。详情请见this link。

编辑：让每个线程获取 2 个值，直到您到达 132。当然，有些工作项可能最终在这个过程中什么都不做

int x = get_local_id(0);
int y = get_local_id(1);

if (2*x < 132 && 2*y < 132) {
    img_buffer[2*x][2*y] = input[...];
    img_buffer[2*x][2*y + 1] = input[...];
    img_buffer[2*x + 1][2*y] = input[...];
    img_buffer[2*x + 1][2*y + 1] = input[...];
}
barrier(CLK_LOCAL_MEM_FENCE);

// filter here...

【讨论】：

缓冲区比工作组稍大。例如，如果工作组为 128*128，则缓冲区将为 132*132。这种方式还能加载吗？
@LimingXu 肯定是，你可以看看更新的答案。这可能不是最有效的方法，但你明白了......

【解决方案2】：

您可以使用async_work_group_copy 执行复制到本地。它将返回一个事件对象，您可以立即使用wait_group_events 等待它。

或者在此期间进行一些不需要复制数据的其他处理。（即：计算过滤器的一些因素）。让工作组保持忙碌。

//Copy input from global to local
event_t global2local = async_work_group_copy(img_buffer, input, size, 0);
wait_group_events(1, &global2local);

//Copy the buffer from local to global
event_t local2global = async_work_group_copy(output_buffer, output, size2, 0);
wait_group_events(1, &local2global);

【讨论】：