Как использовать переменные __local в OpenCL?

Я пытаюсь реализовать алгоритм свертки в 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
person Liming Xu    schedule 20.05.2019    source источник

Ответы (2)


arrow_upward
0
arrow_downward

Предполагая, что у вас есть один рабочий элемент на входной пиксель, вы можете заставить каждый поток загружать один пиксель из глобальной памяти в локальную память.

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...
person AlexG    schedule 21.05.2019
comment
Буфер немного больше рабочей группы. Например, если рабочая группа 128 * 128, буфер будет 132 * 132. Можно ли еще так загрузить? - person Liming Xu; 22.05.2019
comment
@LimingXu уверен, что это так, вы можете взглянуть на обновленный ответ. Вероятно, это не самый эффективный способ сделать это, но вы поняли ... - person AlexG; 22.05.2019

arrow_upward
0
arrow_downward

Для выполнения копия на локальный. Он вернет вам объект события, который вы можете немедленно дождаться с помощью _ 2_.

Или тем временем выполните другую обработку, которая не требует копирования данных. (например: вычисление некоторых факторов для фильтра). Чтобы рабочая группа была занята.

//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);
person DarkZeros    schedule 23.05.2019