opencl - 如何在 OpenCL 中使用 __local 变量？

Question

我正在尝试在 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
     * }
     */
}

Answer 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

如果您的数据是对齐的，您可以通过转换为更大的数据类型来加快速度。有关更多详细信息，请参阅此链接。

编辑：让每个线程获取 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...

Answer 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);