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

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

0 голосов
/ 20 мая 2019

Я пытаюсь реализовать алгоритм свертки в 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
     * }
     */
}

Ответы [ 2 ]

0 голосов
/ 23 мая 2019

Вы можете использовать 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);
0 голосов
/ 21 мая 2019

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

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...
...