Как сделать пошаговую копию из глобальной в локальную память?

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

Я знаю о async_work_group_copy, и это красиво (точнее, неуклюже и раздражающе, но работает). Однако мои данные не являются непрерывными - они разделены, то есть между каждыми двумя последовательными байтами Y, которые я хочу скопировать, может быть X байтов.

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

threadId = get_local_id(0);
if (threadId < length) {
    unsigned offset = threadId * stride;
    localData[threadId] = globalData[offset];
}

memory stride opencl memory-model gpgpu
person einpoklum    schedule 18.07.2013    source источник
comment
@AndreasNiedermair: правильный комментарий, но я задавал этот вопрос уже 3 года назад.   -  person einpoklum    schedule 13.09.2016

Ответы (1)


arrow_upward
4
arrow_downward

Вы можете использовать async_work_group_strided_copy() OpenCL API вызов.

Вот небольшой пример в pyopencl благодаря комментарию @DarkZeros. Предположим, это небольшая полоса изображения RGB, например, 4 на 1:

img = np.array([58, 83, 39, 157, 190, 199, 64, 61, 5, 214, 141, 6])

и вы хотите получить доступ к четырем красным каналам, т.е. [58 157 64 214], вы должны сделать:

def test_asyc_copy_stride_to_local(self):
    #Create context, queue, program first
     ....
    #number of R channels
    nb_of_el = 4
    img = np.array([58, 83, 39, 157, 190, 199, 64, 61, 5, 214, 141, 6])
    cl_input = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=img)
    #buffer used to check if the copy is correct
    cl_output = cl.Buffer(ctx, mf.WRITE_ONLY, size=nb_of_el * np.dtype('int32').itemsize)
    lcl_buf = cl.LocalMemory(nb_of_el * np.dtype('int32').itemsize)
    prog.asynCopyToLocalWithStride(queue, (nb_of_el,), None, cl_input, cl_output, lcl_buf)
    result = np.zeros(nb_of_el, dtype=np.int32)
    cl.enqueue_copy(queue, result, cl_output).wait()
    print result

Ядро:

kernel void asynCopyToLocalWithStride(global int *in, global int *out, local int *localBuf){
    const int idx = get_global_id(0);
    localBuf[idx] = 0;
    //copy 4 elements, the stride = 3 (RGB)
    event_t ev = async_work_group_strided_copy(localBuf, in, 4, 3, 0);
    wait_group_events (1, &ev);
    out[idx] = localBuf[idx];
}
person CaptainObvious    schedule 18.07.2013
comment
@DarkZeros Обновлено с примером - person CaptainObvious; 25.07.2013
comment
Спасибо, подробный пример, подобный вашему, отсутствует в официальном документе функции. - person DarkZeros; 05.08.2013