CUDA версия медленнее, чем версия процессора?

Question

Я пишу подсимплекс изображения в CUDA и использую потоки для выполнения операции усреднения. Однако, если я это сделаю без вызова ядра, он работает намного быстрее по сравнению с тем, когда я на самом деле вызываю ядро ​​CUDA. Размер элемента сейчас составляет 1280x1024. Является ли вызов ядра обычно занимающим значительное время или что-то не так с моей реализацией?

P.S Я попытался вызвать только ядро ​​(с удаленным кодом), и это почти в то же время, что и ядро ​​с кодом. Также мой код без вызова ядра работает примерно 350 мс, тогда как вызов ядра работает до 1000 мс.

__global__ void subsampler(int *r_d,int *g_d,int *b_d, int height,int width,int *f_r,int*f_g,int*f_b){ 
     int id=blockIdx.x * blockDim.x*blockDim.y+ threadIdx.y*blockDim.x+threadIdx.x+blockIdx.y*gridDim.x*blockDim.x*blockDim.y; 
     if (id<height*width/4){ 
     f_r[id]=(r_d[4*id]+r_d[4*id+1]+r_d[4*id+2]+r_d[4*id+3])/4; 
     f_g[id]=(g_d[4*id]+g_d[4*id+1]+g_d[4*id+2]+g_d[4*id+3])/4; 
     f_b[id]=(b_d[4*id]+b_d[4*id+1]+b_d[4*id+2]+b_d[4*id+3])/4; 
     } 
     } 

Я определяю blockSizeX и blockSizeY быть 1 и 1 (я пытался сделать их 4,16), но как-то это самый быстрый

dim3 blockSize(blocksizeX,blocksizeY); 
    int new_width=img_width/2; 
    int new_height=img_height/2; 

    int n_blocks_x=new_width/blocksizeX+(new_width/blocksizeY == 0 ?0:1); 
    int n_blocks_y=new_height/blocksizeX+(new_height/blocksizeY == 0 ?0:1); 
    dim3 gridSize(n_blocks_x,n_blocks_y); 

, а затем я называю ядро ​​с gridSize, BLOCKSIZE.

Answer 1

Возможно, ядро ​​не реализовано очень хорошо, или может случиться так, что накладные расходы на перемещение ваших данных на и с карты графического адаптера будут замалчивать любое вычислительное преимущество. Попробуйте сравнить ядро ​​в изоляции (без CPU < -> передача памяти GPU), чтобы узнать, сколько всего вашего времени занимает ядро ​​и сколько происходит с помощью передачи памяти. Затем вы можете решить на основе этих измерений, нужно ли больше работать над ядром.

Answer 2

Хотя я не уверен, что оборудование вы используете эту, вы должны быть в состоянии сделать это ядро ​​выполнять ближе к 1000 кадров в секунду, а не 1000мс/кадр :)

Рекомендация 1: Если это обработка имеет любое взаимодействие с визуализацией, через OpenGL/DirectX или подобное, просто делайте это как шейдер - все детали размера сетки/блока, макета памяти и т. д. обрабатываются для вас. Если вам действительно нужно реализовать это самостоятельно в CUDA, продолжайте читать:

Во-первых, я предполагаю, что вы подбираете изображение 1280x1024 с коэффициентом 2 в каждом направлении, что дает изображение 640x512. Каждый пиксель в результирующем изображении представляет собой среднее значение четырех пикселей исходного изображения. Изображения имеют три канала, RGB.

Вопрос 1: Вы действительно хотите 32 бит на канал или вы хотите RGB888 (8 бит на канал)? RGB888 довольно распространен - ​​я предполагаю, что это то, что вы имели в виду.

Вопрос 2: Являются ли ваши данные на самом деле плоскими или вы извлекаете их из чередующегося формата? RGB888 является чередующимся форматом, где пиксели отображаются в памяти как RGBRGBRGB. Я бы написал ваше ядро ​​для обработки изображения в собственном формате. Я предполагаю, что ваши данные на самом деле плоские, поэтому у вас есть три плоскости: R8, G8 и B8.

Первое, что нужно сделать, это рассмотреть расположение макетов. Вам понадобится один поток для каждого пикселя в целевом изображении. Учитывая, что шаблон доступа к памяти для подвыборки не объединен, вы захотите прочитать данные пикселов в общую память. Рассмотрим размер блока 32х8 потоков. Это позволяет каждому блоку читать 40 * 8 * 4 пикселя или 3072 байта при 3bpp. На самом деле вы будете читать чуть больше, чтобы сохранить нагрузки объединенными, в общей сложности 4096 байт на блок. Это теперь дает вам:

dim3 block(32, 8); 
dim3 grid(1280/2/32, 1024/2/8); // 20x64 blocks of 256 threads 

Теперь наступает интересная часть: выполняется разделяемая память.Ваше ядро ​​может выглядеть так:

__global__ void subsample(uchar* r, uchar* g, uchar* b, // in 
          uchar* ro, uchar* go, uchar* bo) // out 
{ 
    /* Global offset into output pixel arrays */ 
    int gid = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x; 

    /* Global offset into input pixel arrays */ 
    int gidin = gid * 2; 

    __shared__ uchar* rc[1024]; 
    __shared__ uchar* gc[1024]; 
    __shared__ uchar* bc[1024]; 

    /* Read r, g, and b, into shmem cache */ 
    ((int*)rc)[threadIdx.x] = ((int*)r)[gidin + threadIdx.x]; 
    ((int*)gc)[threadIdx.x] = ((int*)g)[gidin + threadIdx.x]; 
    ((int*)bc)[threadIdx.x] = ((int*)b)[gidin + threadIdx.x]; 

    __syncthreads(); 

    /* Shared memory for output */ 
    __shared__ uchar* roc[256]; 
    __shared__ uchar* goc[256]; 
    __shared__ uchar* boc[256]; 

    /* Do the subsampling, one pixel per thread. Store into the output shared memory */ 

    ... 

    __syncthreads(); 

    /* Finally, write the result to global memory with coalesced stores */ 
    if (threadIdx.x < 64) { 
     ((int*)ro)[gid + threadIdx.x] = ((int*)roc)[threadIdx.x]; 
    } else if (threadIdx.x < 128) { 
     ((int*)go)[gid + threadIdx.x-64] = ((int*)goc)[threadIdx.x-64]; 
    } else if (threadIdx.x < 192) { 
     ((int*)bo)[gid + threadIdx.x-128] = ((int*)boc)[threadIdx.x-128]; 
    } 
} 

Whew! Много чего там, извините за свалку кода. Некоторые принципы, которые следует учитывать:

1) Память выполняется быстро, когда вы используете коалесцированные нагрузки/хранилища. Это означает, что для каждого потока в перекосе 32 каждый обращается к 32 байтам. Если индекс 32 байта соответствует индексу потока в warp, тогда все 32 обращения получаются в одну транзакцию 128. Так вы получаете пропускную способность 100 ГБ/с на GPU.

2) Образец доступа к памяти при выполнении подвыборки не объединен, поскольку он основан на пространственной пространственной 2D-памяти, которой необработанной памяти нет. (Для этого также можно использовать текстурную память ...). Сохраняя входные данные в общей памяти, затем обрабатывая, вы минимизируете влияние своей вычислительной производительности.

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