CUDA - обрабатывать данные (массив) буфера одного пикселя на нескольких одновременных ядрах, возможно ли это?

В настоящее время у меня есть один пиксельный буфер, и я обрабатываю данные в нем одним вызовом ядра:

dim3 threadsPerBlock(32, 32)
dim3 blocks(screenWidth / threadsPerBlock.x, screenHeight / threadsPerBlock.y);
kernel<<<blocks, threadsPerBlock>>>();

Буфер пикселей содержит все пиксели окна с размерами screenWidth x screenHeight.

Моя идея состоит в том, чтобы разделить окно на 2 или 4 части и одновременно обрабатывать данные о пикселях.

Можно ли это сделать, и если можно, то как?

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

Редактировать: Моя видеокарта с вычислительными возможностями 3.0

Изменить 2: я использую SDL для рисования, у меня есть один графический процессор, и я использую пользовательский массив данных:

main.ru

 Color vfb_linear[VFB_MAX_SIZE * VFB_MAX_SIZE]; // array on the Host
 Color vfb[VFB_MAX_SIZE][VFB_MAX_SIZE] // 2D array used for SDL
 extern "C" void callKernels(Color* dev_vfb);

int main()
{
    Color* dev_vfb; // pixel array used on the GPU
    // allocate memory for dev_vfb on the GPU
    cudaMalloc((void**)&dev_vfb, sizeof(Color) * RES_X * RES_Y);
    // memcpy HostToDevice
    cudaMemcpy(dev_vfb, vfb_linear, sizeof(Color) * RES_X * RES_Y, cudaMemcpyHostToDevice);

    callKernels(dev_vfb); // wrapper function that calls the kernels

    // memcpy DeviceToHost
    cudaMemcpy(vfb_linear, dev_vfb, sizeof(Color) * RES_X * RES_Y, cudaMemcpyDeviceToHost);

    // convert vfb_linear into 2D array so it can be handled by SDL
    convertDeviceToHostBuffer();    

    display(vfb); // render pixels on screen with SDL

}

cudaRenderer.cu

__global__ void kernel(Color* dev_vfb)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    if (offset < RES_X * RES_Y)
    {
        dev_vfb[offset] = getColorForPixel();
    }
}

extern "C" callKernels(Color* dev_vfb)
{
    dim3 threadsPerBlock(32, 32)
    dim3 blocks(screenWidth / threadsPerBlock.x, screenHeight / threadsPerBlock.y);
    kernel<<<blocks, threadsPerBlock>>>(dev_vfb);
}

содержимое дисплея (vfb):

void display(Color vfb[VFB_MAX_SIZE][VFB_MAX_SIZE])
{
    // screen is pointer to SDL_Surface
    int rs = screen->format->Rshift;
    int gs = screen->format->Gshift;
    int bs = screen->format->Bshift;

    for (int y = 0; y < screen->h; ++y)
    {
        Uint32* row = (Uint32*) ((Uint8*) screen->pixels + y * screen->pitch);
        for (int x = 0; x < screen->w; ++x)
            row[x] = vfb[y][x].toRGB32(rs, gs, bs);
    }
    SDL_Flip(screen);
}

Это простой пример того, что я делаю в своем проекте. Это трассировщик лучей, и, возможно, SDL — худший выбор для взаимодействия с CUDA, но я не знаю, успею ли я его изменить.


c++ cuda cuda-streams
person Geto    schedule 04.02.2014    source источник
comment
Это возможно. Вы имеете в виду OpenGL или DX PBO или просто простой массив данных, в котором есть пиксельные данные? Код, который вы показали, не имеет значения для ответа на ваш вопрос. Из него нельзя даже составить представление о том, на что похож ваш массив данных. Похоже, у вас есть один графический процессор. В этом случае маловероятно, что вы сможете повысить производительность, запустив два ядра одновременно, а не одно хорошо написанное ядро.   -  person Robert Crovella    schedule 04.02.2014
comment
Я добавлю больше информации в свой вопрос.   -  person Geto    schedule 04.02.2014
comment
Я добавил больше информации, и я понимаю, что я могу не получить производительности от одновременных ядер, но стоит попробовать, если это можно сделать.   -  person Geto    schedule 04.02.2014

Ответы (1)


arrow_upward
1
arrow_downward

Ничто не мешает двум потокам работать с одним и тем же фрагментом данных в глобальной памяти одного устройства.

Как я сказал в комментариях, я не думаю, что это разумный подход для ускорения работы. Однако изменения в вашем коде будут примерно такими (закодировано в браузере, не проверено):

__global__ void kernel(Color* dev_vfb, int slices)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    if (offset < (RES_X * RES_Y/slices)
    {
        dev_vfb[offset] = getColorForPixel();
    }
}

extern "C" callKernels(Color* dev_vfb)
{
    int num_slices=2;
    cudaStream_t streams[num_slices];
    for (int i = 0; i < num_slices; i++)
      cudaStreamCreate(&(streams[i]));
    dim3 threadsPerBlock(32, 32)
    dim3 blocks(screenWidth / threadsPerBlock.x, screenHeight / (num_slices*threadsPerBlock.y));
    for (int i = 0; i < num_slices; i++){
      int off = i * (screenWidth*screenHeight/num_slices);
      kernel<<<blocks, threadsPerBlock, 0, streams[i]>>>(dev_vfb+off, num_slices); }
}
person Robert Crovella    schedule 08.02.2014
comment
Это работает, и, как вы сказали, это не ускорило программу. Я отмечу это как ответ, потому что это то, что я хотел увидеть, как это делается. Есть ли у вас какие-либо предложения о том, как ускорить это конкретное ядро, или мне следует заглянуть в внутренности моей функции getColorForPixel() и поискать там улучшения? - person Geto; 08.02.2014
comment
Да, пожалуй, основное внимание следует уделить getColorForPixel(). - person Robert Crovella; 08.02.2014