Замените CUDA cudaMemcpy2DToArray() на OpenCL clEnqueueCopyBufferRect()

Я уже некоторое время работаю с портированием CUDA-> OpenCL. У меня есть вопрос, на который я пока нигде не нашел ответа. Таким образом, я думаю, это будет интересный вопрос для портировщиков CUDA->OpenCL.

У меня все шло гладко, ожидая «перевода» одной функции. Это функция cudaMemcpy2DToArray(). Простой пример CUDA с использованием этой функции будет выглядеть так:

uint8_t*    srcBuffer;
cudaArray*  dstArray;

size_t      pitch;
size_t      x = 700;
size_t      y = 300;

cudaMallocPitch((void **)&srcBuffer,
                &pitch,                     // pitch = 1024
                x*sizeof(uint8_t),
                y);

struct cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uint8_t>();

cudaMallocArray(&dstArray,
                &channelDesc,
                x,
                y,
                cudaArrayDefault);

/* Fill srcBuffer with some valuable data via a kernel,
   USING pitch when writing to srcBuffer:

   (uint8_t*)((char*)srcBuffer + Row * pitch) + Column = someuint8Value; */

cudaMemcpy2DToArray(dstArray,
                    0,
                    0,
                    (void *) srcBuffer,
                    pitch,
                    x,
                    y,
                    cudaMemcpyDeviceToDevice);

Версия OpenCL будет выглядеть так:

cl_mem      srcBuffer;
cl_mem      dstImage;

size_t      pitch;
size_t      x = 700;
size_t      y = 300;

В OpenCL нет функции cudaMallocPitch(). Чтобы добиться того же поведения, что и cudaMallocPitch(), необходимо сначала запросить предпочтительный размер выравнивания устройства, затем вычислить шаг и, наконец, выделить исходный буфер с помощью clCreateBuffer(). функция:

uint32_t    alignBits;
uint32_t    alignBytes;

clGetDeviceInfo (deviceID, /* Definition not shown here */
                 CL_DEVICE_MEM_BASE_ADDR_ALIGN,
                 sizeof(uint32_t),
                 &alignBits,                        // alignBits = 4096
                 NULL);

alignBytes = alignBits/8;                           // alignBytes = 512
pitch = (x*sizeof(uint8_t) + alignBytes - 1)/       // pitch = 1024
                             alignBytes*alignBytes;

srcBuffer = clCreateBuffer(context, /* Definition not shown here */
                           CL_MEM_READ_WRITE,
                           y*pitchBytes,
                           NULL,
                           NULL);

В OpenCL есть функция, которая ведет себя как cudaMallocArray(). Это clCreateImage2D():

cl_image_format channelDesc = {CL_A, CL_UNSIGNED_INT8};

dstImage = clCreateImage2D(context, /* Definition not shown here */
                           CL_MEM_READ_ONLY,
                           &channelDesc,
                           x,
                           y,
                           0, NULL, NULL);

Для согласованности добавлен следующий блок кода:

/* Fill srcBuffer with some valuable data via a kernel,
   USING pitch when writing to srcBuffer:

   (uchar*)((char*)srcBuffer + Row * pitch) + Column = someucharValue; */

Теперь мы в интересной части. Функция OpenCL с поведением, наиболее близким к поведению cudaMemcpy2DToArray(), которое я впервые смог найти, была clEnqueueCopyBufferToImage(). Однако у этой функции не было параметра pitch. Без параметра шага было бы невозможно правильно сделать копию буфера->изображения, так как количество элементов различается. Я понял, что могу использовать функцию OpenCL clEnqueueCopyBufferRect(), чтобы сначала скопировать данные из srcBuffer в другой буфер с тем же количеством элементов, что и < em>dstImage. И, наконец, используйте clEnqueueCopyBufferToImage(). Вот как это выглядит:

srcBuffer_tmp = clCreateBuffer(context, /* Definition not shown here */
                               CL_MEM_READ_WRITE,
                               x*y*sizeof(uint8_t),
                               NULL,
                               NULL);

size_t src_origin = {0, 0, 0};
size_t dst_origin = {0, 0, 0};
size_t region     = {x, y, 1};

clEnqueueCopyBufferRect(commandQueue, /* Definition not shown here */
                        srcBuffer,
                        srcBuffer_tmp,
                        src_origin,
                        dst_origin,
                        region,
                        pitch,
                        0,
                        x*sizeof(uint8_t),
                        0,
                        0, NULL, NULL);

clEnqueueCopyBufferToImage(commandQueue,
                           srcBuffer_tmp,
                           dstImage,
                           0,
                           dst_origin,
                           region,
                           0, NULL, NULL);

clReleaseMemObject(srcBuffer_tmp);

Вы заметите, что это НЕ дает такого же поведения. Мы выделили больше памяти, буфер srcBuffer_tmp. Немного поэкспериментировав, я обнаружил, что приведенный выше код можно просто заменить на:

size_t src_origin = {0, 0, 0};
size_t dst_origin = {0, 0, 0};
size_t region     = {x, y, 1};

clEnqueueCopyBufferRect(commandQueue, /* Definition not shown here */
                        srcBuffer,
                        dstImage,     /* DIRECT TO IMAGE */
                        src_origin,
                        dst_origin,
                        region,
                        pitch,
                        0,
                        x*sizeof(uint8_t),
                        0,
                        0, NULL, NULL);

Это дало мне поведение при сохранении как cudaMemcpy2DToArray(). Однако использование функции clEnqueueCopyBufferRect() действительно выглядит НЕПРАВИЛЬНЫМ. Отсутствие функции Rect-copy для изображений очевидно, поскольку все уже обрабатывается в 2D.

Мой вопрос:

Можно ли предположить, что clEnqueueCopyBufferRect() будет вести себя как функция копирования буфера->изображения, если дескриптор формата изображения (cl_image_format) инициализирован с порядком каналов изображения < strong>CL_A и тип канала изображения, равный типу элемента, используемого в исходном буфере?