Ответ 1

Это классический случай неловко параллельной проблемы с обработкой изображений, которая может быть очень легко сопоставлена ​​с каркасом CUDA. Фильтр усреднения известен как Box Filter в доменах обработки изображений.

Самый простой подход - использовать текстуры CUDA для процесса фильтрации, так как граничные условия могут обрабатываться очень легко текстурами.

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

• Выделите достаточно большую память для хранения исходных и целевых изображений на устройстве.
• Скопировать исходное изображение с хоста на устройство.
• Привязать указатель устройства исходного изображения к текстуре.
• Укажите подходящий размер блока и сетку, достаточную для покрытия каждого пикселя изображения.
• Запустите ядро ​​фильтрации с использованием заданного размера сетки и блока.
• Скопировать результаты на хост.
• Отвяжите текстуру
• Свободные указатели на устройства.

Пример реализации фильтра ящиков

Kernel

texture<unsigned char, cudaTextureType2D> tex8u;

//Box Filter Kernel For Gray scale image with 8bit depth
__global__ void box_filter_kernel_8u_c1(unsigned char* output,const int width, const int height, const size_t pitch, const int fWidth, const int fHeight)
{
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y;

    const int filter_offset_x = fWidth/2;
    const int filter_offset_y = fHeight/2;

    float output_value = 0.0f;

    //Make sure the current thread is inside the image bounds
    if(xIndex<width && yIndex<height)
    {
        //Sum the window pixels
        for(int i= -filter_offset_x; i<=filter_offset_x; i++)
        {
            for(int j=-filter_offset_y; j<=filter_offset_y; j++)
            {
                //No need to worry about Out-Of-Range access. tex2D automatically handles it.
                output_value += tex2D(tex8u,xIndex + i,yIndex + j);
            }
        }

        //Average the output value
        output_value /= (fWidth * fHeight);

        //Write the averaged value to the output.
        //Transform 2D index to 1D index, because image is actually in linear memory
        int index = yIndex * pitch + xIndex;

        output[index] = static_cast<unsigned char>(output_value);
    }
}

Функция Wrapper:

void box_filter_8u_c1(unsigned char* CPUinput, unsigned char* CPUoutput, const int width, const int height, const int widthStep, const int filterWidth, const int filterHeight)
{

    /*
     * 2D memory is allocated as strided linear memory on GPU.
     * The terminologies "Pitch", "WidthStep", and "Stride" are exactly the same thing.
     * It is the size of a row in bytes.
     * It is not necessary that width = widthStep.
     * Total bytes occupied by the image = widthStep x height.
     */

    //Declare GPU pointer
    unsigned char *GPU_input, *GPU_output;

    //Allocate 2D memory on GPU. Also known as Pitch Linear Memory
    size_t gpu_image_pitch = 0;
    cudaMallocPitch<unsigned char>(&GPU_input,&gpu_image_pitch,width,height);
    cudaMallocPitch<unsigned char>(&GPU_output,&gpu_image_pitch,width,height);

    //Copy data from host to device.
    cudaMemcpy2D(GPU_input,gpu_image_pitch,CPUinput,widthStep,width,height,cudaMemcpyHostToDevice);

    //Bind the image to the texture. Now the kernel will read the input image through the texture cache.
    //Use tex2D function to read the image
    cudaBindTexture2D(NULL,tex8u,GPU_input,width,height,gpu_image_pitch);

    /*
     * Set the behavior of tex2D for out-of-range image reads.
     * cudaAddressModeBorder = Read Zero
     * cudaAddressModeClamp  = Read the nearest border pixel
     * We can skip this step. The default mode is Clamp.
     */
    tex8u.addressMode[0] = tex8u.addressMode[1] = cudaAddressModeBorder;

    /*
     * Specify a block size. 256 threads per block are sufficient.
     * It can be increased, but keep in mind the limitations of the GPU.
     * Older GPUs allow maximum 512 threads per block.
     * Current GPUs allow maximum 1024 threads per block
     */

    dim3 block_size(16,16);

    /*
     * Specify the grid size for the GPU.
     * Make it generalized, so that the size of grid changes according to the input image size
     */

    dim3 grid_size;
    grid_size.x = (width + block_size.x - 1)/block_size.x;  /*< Greater than or equal to image width */
    grid_size.y = (height + block_size.y - 1)/block_size.y; /*< Greater than or equal to image height */

    //Launch the kernel
    box_filter_kernel_8u_c1<<<grid_size,block_size>>>(GPU_output,width,height,gpu_image_pitch,filterWidth,filterHeight);

    //Copy the results back to CPU
    cudaMemcpy2D(CPUoutput,widthStep,GPU_output,gpu_image_pitch,width,height,cudaMemcpyDeviceToHost);

    //Release the texture
    cudaUnbindTexture(tex8u);

    //Free GPU memory
    cudaFree(GPU_input);
    cudaFree(GPU_output);
}

Хорошей новостью является то, что вам не нужно самостоятельно внедрять фильтр. CUDA Toolkit поставляется со свободной библиотекой обработки сигналов и обработки изображений, названной NVIDIA Performance Primitives, ака NPP, сделанной NVIDIA. АЭС использует графические процессоры с поддержкой CUDA для ускорения обработки. Фильтр усреднения уже реализован на АЭС. Текущая версия NPP (5.0) поддерживает 8-разрядные, 1-канальные и 4-канальные изображения. Функции:

• nppiFilterBox_8u_C1R для 1-канального изображения.
• nppiFilterBox_8u_C4R для 4-канального изображения.

Ответ 2

Некоторые основные мысли/этапы:

• Скопировать данные изображения из ЦП на графический процессор
• Вызвать ядро ​​для построения среднего для каждой строки (по горизонтали) и сохранить его в общей памяти.
• Вызвать ядро ​​для создания среднего значения для каждого столбца (по вертикали) и сохранить его в глобальной памяти.
• Скопируйте данные обратно в CPU-Memory.

Вы должны иметь возможность масштабировать это довольно легко с помощью 2D-памяти и многомерных вызовов ядра.

Ответ 3

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

Если вы храните изображение в глобальной памяти, его можно запрограммировать легко, но у вас будет много банковских конфликтов. Одна из возможных оптимизаций заключается в том, что вы загружаете блоки изображения в блок Shared Memory. Используя элементы phantom (чтобы вы не превышали размеры разделяемого блока при поиске соседних пикселей), вы можете рассчитать среднее значение пикселей внутри блока.

Единственное, что вы должны быть осторожны, это то, как "сшивание" будет сделано в конце, потому что блоки разделяемой памяти будут перекрываться (из-за дополнительных "отступов" пикселей), и вы не хотите вычислить их значения дважды.