Cuda Bayer / CFA пример демозаики

Question

Я написал подпрограмму демодуляции CUDA4 Bayer, но она медленнее однопоточного кода ЦП, работающего на 16-ядерном GTS250.
Размер блока равен (16,16), а яркость изображения кратна 16 - но изменив этоне улучшает его.

Я делаю что-то явно глупое?

--------------- calling routine ------------------
uchar4 *d_output;
size_t num_bytes; 

cudaGraphicsMapResources(1, &cuda_pbo_resource, 0);    
cudaGraphicsResourceGetMappedPointer((void **)&d_output, &num_bytes, cuda_pbo_resource);

// Do the conversion, leave the result in the PBO fordisplay
kernel_wrapper( imageWidth, imageHeight, blockSize, gridSize, d_output );

cudaGraphicsUnmapResources(1, &cuda_pbo_resource, 0);

--------------- cuda -------------------------------
texture<uchar, 2, cudaReadModeElementType> tex;
cudaArray *d_imageArray = 0;

__global__ void convertGRBG(uchar4 *d_output, uint width, uint height)
{
    uint x = __umul24(blockIdx.x, blockDim.x) + threadIdx.x;
    uint y = __umul24(blockIdx.y, blockDim.y) + threadIdx.y;
    uint i = __umul24(y, width) + x;

    // input is GR/BG output is BGRA
    if ((x < width) && (y < height)) {

        if ( y & 0x01 ) {
            if ( x & 0x01 ) {  
                d_output[i].x =  (tex2D(tex,x+1,y)+tex2D(tex,x-1,y))/2;  // B                
                d_output[i].y = (tex2D(tex,x,y));     // G in B
                d_output[i].z = (tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/2;  // R                    
            } else {
                d_output[i].x = (tex2D(tex,x,y));        //B
                d_output[i].y = (tex2D(tex,x+1,y) + tex2D(tex,x-1,y)+tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/4;  // G
                d_output[i].z = (tex2D(tex,x+1,y+1) + tex2D(tex,x+1,y-1)+tex2D(tex,x-1,y+1)+tex2D(tex,x-1,y-1))/4;   // R
            }
        } else {
            if ( x & 0x01 ) {
                 // odd col = R
                d_output[i].y = (tex2D(tex,x+1,y+1) + tex2D(tex,x+1,y-1)+tex2D(tex,x-1,y+1)+tex2D(tex,x-1,y-1))/4;  // B
                d_output[i].z = (tex2D(tex,x,y));        //R
                d_output[i].y = (tex2D(tex,x+1,y) + tex2D(tex,x-1,y)+tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/4;  // G    
            } else {    
                d_output[i].x = (tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/2;  // B
                d_output[i].y = (tex2D(tex,x,y));               // G  in R               
                d_output[i].z = (tex2D(tex,x+1,y)+tex2D(tex,x-1,y))/2;  // R                    
            }
        }                                
    }
}



void initTexture(int imageWidth, int imageHeight, uchar *imagedata)
{

    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 0, 0, 0, cudaChannelFormatKindUnsigned);
    cutilSafeCall( cudaMallocArray(&d_imageArray, &channelDesc, imageWidth, imageHeight) ); 
    uint size = imageWidth * imageHeight * sizeof(uchar);
    cutilSafeCall( cudaMemcpyToArray(d_imageArray, 0, 0, imagedata, size, cudaMemcpyHostToDevice) );
    cutFree(imagedata);

    // bind array to texture reference with point sampling
    tex.addressMode[0] = cudaAddressModeClamp;
    tex.addressMode[1] = cudaAddressModeClamp;
    tex.filterMode = cudaFilterModePoint;
    tex.normalized = false; 

    cutilSafeCall( cudaBindTextureToArray(tex, d_imageArray) );
}

Answer 1

В вашем коде нет очевидных ошибок, но есть несколько очевидных возможностей для повышения производительности:

1) для лучшей производительности вы должны использовать текстуру для перехода в общую память - см. SDK 'SobelFilter'sample.

2) Как написано, код записывает байты в глобальную память, что гарантированно приведет к значительному снижению производительности.Вы можете использовать разделяемую память для поэтапного представления результатов перед их фиксацией в глобальной памяти.

3) Удивительно большое преимущество в производительности заключается в том, что размеры блоков соответствуют параметрам аппаратного кэша текстур.На оборудовании класса Tesla оптимальный размер блока для ядер, использующих ту же схему адресации, что и у вашего ядра, составляет 16x4.(64 потока на блок)

Для таких рабочих нагрузок может быть сложно конкурировать с оптимизированным кодом ЦП.SSE2 может выполнять 16 операций размером в байт в одной инструкции, а процессоры работают примерно в 5 раз быстрее.

Answer 2

Основываясь на ответах на форумах Nvidia, здесь (для поисковых систем) приведена чуть более оптимизированная версия, которая записывает блок пикселей 2x2 в каждой теме.Хотя разница в скорости не может быть измерена на моей установке.

Обратите внимание, что ее следует вызывать с размером сетки половина размера изображения;

dim3 blockSize(16, 16); // for example
dim3 gridSize((width/2) / blockSize.x, (height/2) / blockSize.y);


__global__ void d_convertGRBG(uchar4 *d_output, uint width, uint height)
{
    uint x = 2 * (__umul24(blockIdx.x, blockDim.x) + threadIdx.x);
    uint y = 2 * (__umul24(blockIdx.y, blockDim.y) + threadIdx.y);
    uint i = __umul24(y, width) + x;

    // input is GR/BG output is BGRA
    if ((x < width-1) && (y < height-1)) {
        // x+1, y+1:

        d_output[i+width+1] = make_uchar4( (tex2D(tex,x+2,y+1)+tex2D(tex,x,y+1))/2,  // B                
                                             (tex2D(tex,x+1,y+1)),     // G in B
                                             (tex2D(tex,x+1,y+2)+tex2D(tex,x+1,y))/2,  // R                    
                                             0xff);

        // x, y+1:
        d_output[i+width] =   make_uchar4( (tex2D(tex,x,y+1)),        //B
                                             (tex2D(tex,x+1,y+1) + tex2D(tex,x-1,y+1)+tex2D(tex,x,y+2)+tex2D(tex,x,y))/4,  // G
                                             (tex2D(tex,x+1,y+2) + tex2D(tex,x+1,y)+tex2D(tex,x-1,y+2)+tex2D(tex,x-1,y))/4,   // R
                                             0xff);


        // x+1, y:
        d_output[i+1] =       make_uchar4( (tex2D(tex,x,y-1) + tex2D(tex,x+2,y-1)+tex2D(tex,x,y+1)+tex2D(tex,x+2,y-1))/4,  // B
                                            (tex2D(tex,x+2,y) + tex2D(tex,x,y)+tex2D(tex,x+1,y+1)+tex2D(tex,x+1,y-1))/4,  // G
                                            (tex2D(tex,x+1,y)),        //R
                                            0xff);


        // x, y:
        d_output[i] =         make_uchar4( (tex2D(tex,x,y+1)+tex2D(tex,x,y-1))/2,  // B
                                             (tex2D(tex,x,y)),               // G  in R           
                                             (tex2D(tex,x+1,y)+tex2D(tex,x-1,y))/2,  // R                    
                                             0xff);

    }
}

Answer 3

В коде много символов if и else.Если вы структурируете код так, чтобы исключить все условные операторы, то вы получите огромный прирост производительности, поскольку ветвление снижает производительность.Это действительно возможно, чтобы удалить ветви.Есть ровно 30 случаев, которые вам придется явно кодировать.Я реализовал его на CPU, и он не содержит никаких условных операторов.Я думаю сделать блог, объясняющий это.Опубликуем его, как только сделаем.