Асинхронная передача данных, управляемая ядрами в OpenCL? - PullRequest
Новая
       6

Асинхронная передача данных, управляемая ядрами в OpenCL?

0 голосов
/ 05 октября 2018

В моей конкретной проблеме мне нужно оценить функцию с независимой переменной x в ядрах с помощью интерполяционной формулы, которая использует данные, которые хранятся в массиве в памяти хоста.Аргумент функции становится доступным для каждого ядра только к моменту его запуска. Как только аргумент становится доступным, я бы хотел прекратить выполнение ядра и переместить соответствующий кусок данных, а конкретный double8от хоста до домена памяти на устройстве, которое предлагает лучшее преимущество в производительности (я думаю, что в идеале local Memory), асинхронно.Мне интересно, могу ли я использовать и как эффективно использовать объекты памяти для достижения этой цели?Я рассматриваю следующие конструкции OpenCL: 

cl_mem clCreateBuffer( cl_context ctx,
                       cl_mem     flags,
                       syse_t     size,
                       void*      host_ptr,
                       cl_int     errorcode_ret)

err = clEnqueueWriteBuffer(
     command_queue,    // command queue managing the transaction
     output,           // buffer object to write to, could it be in local memory?
     CL_TRUE,          // indicating a blocking transfer
     0,                // offset in the output to start writing the data
     size,             // size of the data transfer  
     host_ptr,         // pointer to the buffer in host memory holding the data
     0,                // number of event, what could I do with this?
     NULL,             // number of events that predate the current one? the previous argument, I guess? 
     NULL,             // event object to return after successful completion
     );

Следуя OpenCL Best Practices Guide Я хотел бы использовать рекомендованный шаблон.

1) Объявление cl_mem объектов буфера длязакрепленная память хоста и графическое устройство GMEM, соответственно, и стандартные указатели для ссылки на закрепленную память хоста. 

cl_context cxGPUContext;         // computational context for the current arena
cl_mem cmPinnedBufIn = NULL;     // memory buffer on the host-side 
cl_mem cmDevBufIn = NULL;        // memory buffer on the device-side 
unsigned char* cDataIn = NULL;   // holder for the data buffer

2) Выделение объектов буфера cl_mem для закрепленной памяти хоста и устройства графического процессораGMEM, соответственно при подготовке к транзакции: 

cmPinnedBufIn = clCreateBuffer(cxGPUContext,                              // computational context for the current arena
                               CL_MEM_READ_ONLY | CL_MEM_ALLOC_HOST_PTR,  // zero-copy I guess??
                               memSize,                                   // size of the stack of memory to hold the data transfers
                               NULL,                                      // initializing the host pointer to NULL
                               NULL);                                     // error code

cmDevBufIn = clCreateBuffer(cxGPUContext,     // computational context for the current arena
                            CL_MEM_READ_ONLY, // read-only mode in the device side 
                            memSize,          // size of the stack of memory to hold the transaction, it is in the Global memory I guess? 
                            NULL,             // initializing the device pointer to NULL
                            NULL);            // error code

3) Сопоставить стандартный указатель для ссылки на буферизованные буферы ввода и вывода памяти хоста со стандартными указателями. 

cDataIn = (double8 *)clEnqueueMapBuffer(cqCommandQue,  //command queue to manage the data transfer
                                        cmPinnedBufIn, // pinned buffer instantiated avobe
                                        CL_TRUE,       // is this a blocking map?
                                        CL_MAP_WRITE,  // what we are doing? 
                                        0,             // offset in the data 
                                        memSize,       // size of the transfer 
                                        0,             // number of events in the waiting list, howt to capitalize on this?
                                        NULL,          // waiting list  
                                        NULL,          // event
                                        NULL);         // error code

4) Инициализируйте или обновите содержимое закрепленной памяти, используя стандартный указатель узла и стандартный код узла. здесь я определил функцию для извлечения из таблицы интерполяции double8, соответствующегонезависимая переменная: 

cDataIn = get_data(x);

5) Запись данных из закрепленной памяти хоста в устройство GPEM GMEM в любое время в приложении, которое 'свежие данные были записаны в закрепленную память хоста. 

err = clEnqueueWriteBuffer(
     cqCommandQue,     // command queue managing the transaction
     cmDevBufIn,       // buffer object to write to, could it be in local memory?
     CL_FALSE,         // indicating that this is not a blocking transfer
     0,                // offset in the output to start writing the data
     sizeof(double8),  // size of the data transfer  
     cDataIn,          // pointer to the buffer in host memory holding the data
     0,                // number of event, what could I do with this?
     NULL,             // number of events that predate the current one? the previous argument, I guess? 
     NULL,             // event object to return after successful completion
     );

6) Запуск вычислительного ядра на устройстве с графическим процессором. На этом шаге мое приложение ломаетшаблон более стандартных вычислений в том смысле, что блок данных, который необходимо передать для завершения вычисления, зависит от значения x здесь. например, допустим, что ядро ​​похоже на то, что в этом вопрос : 

//d_kernel.cl

__kernel void distance_kernel(__global double *pixelInfo,
                                __global double *clusterCentres,
                                __global double *distanceFromClusterCentre)
{
    int index = get_global_id(0);

    int d, dl, da, db, dx, dy;

    dl = pixelInfo[5 * index] - clusterCentres[0];
    dl = dl * dl;

    da = pixelInfo[5 * index + 1] - clusterCentres[1];
    da = da * da;

    db = pixelInfo[5 * index + 2] - clusterCentres[2];
    db = db * db;

    dx = pixelInfo[5 * index + 3] - clusterCentres[3];
    dx = dx * dx;

    dy = pixelInfo[5 * index + 4] - clusterCentres[4];
    dy = dy * dy;

    double x = dx + dy + dl + da + db;
    // how could I grab from the host the data corresponding to
    // the value taken by x above?
    // let ussupose that the  inlined function get_data(x) does it
    // 'transparently' :)
    double8 point = get_data(x);
    // use point to compute y by interpolation
    double y = interp(x,point);
    istanceFromClusterCentre[index] = y;

}

Меня интересует только проведение Host-to-Device переводов максимально эффективно.Как сказано в Best practices Guide:

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

То, что я наивно полагаю, - лучший способ справиться с этим, не так ли?Моя главная проблема заключается в том, как можно реализовать подпрограмму get_data(x), которая вызывается из ядер для эффективного управления передачей данных Host-to-Device?Я также обеспокоен тем, что выделение закрепленной памяти может занять значительное количество времени, как можно уменьшить этот штраф, выделяя стек памяти в начале выполнения?Каково лучшее предположение для хорошего размера стека?

Это утверждение в спецификации OpenCL заставляет меня думать, что то, что я спрашиваю, возможно:

Глобальная память.Эта область памяти разрешает доступ на чтение / запись ко всем рабочим элементам во всех рабочих группах.Рабочие элементы могут читать или записывать в любой элемент объекта памяти.Чтение и запись в глобальную память могут кэшироваться в зависимости от возможностей устройства.

Но я пока не выясняю, как инициировать передачу данных из ядра.это, как get_data(double x) должен быть реализован.

...