Хотя я не уверен, какое оборудование вы используете для этого, вы должны иметь возможность заставить это ядро работать быстрее, чем 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 байта при 3 битах на дюйм.На самом деле вы будете читать чуть больше, чтобы сохранить объединенные нагрузки, в общей сложности 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];
}
}
Вот так!Там много всего, извините за дамп кода.Необходимо помнить несколько принципов:
1) Память работает быстро, когда вы используете объединенные нагрузки / хранилища.Это означает, что для каждого потока в деформации 32 каждый обращается к 32 байтам.Если 32-байтовый индекс совпадает с индексом потока в деформации, то все 32 доступа будут помещены в одну 128-ю транзакцию.Вот как вы получаете пропускную способность 100 ГБ / с для графического процессора.
2) Схема доступа к памяти при выполнении подвыборки не объединяется, поскольку она опирается на пространственную локальность 2D, которой нет в исходной памяти.(Можно также использовать текстурную память для этого ...) Сохраняя входные данные в разделяемой памяти, а затем обрабатывая, вы минимизируете влияние на производительность вычислений.
Надеюсь, это поможет - я могу ответить более подробно о некоторых деталях, если хотите.