Блог

Я новичок в CUDA, и я пытаюсь реализовать плавную свертку на изображении, и до сих пор у меня это есть, но результат неправильный. не уверен, правильно ли я сделал смещение. Какая-нибудь помощь?

 
__global__ void smooth(unsigned char* device_out_image, float kernel_size, unsigned char* device_input_imag, int height, int width)
{
    int pos_x = threadIdx.x   blockIdx.x * blockDim.x;//x coordinate of pixel
    int pos_y = threadIdx.y   blockIdx.y * blockDim.y;//y coordinate of pixel

    if (pos_x < width amp;amp; pos_y < height)
    {
        unsigned char r = device_input_imag[pos_y * width   pos_x];//absolute pixel position
        unsigned char g = device_input_imag[(height   pos_y) * width   pos_x];
        unsigned char b = device_input_imag[(height * 2   pos_y) * width   pos_x];
        //also mix value with the intensity instead of the range x
        float sumR = float(0.0f);
        float sumG = float(0.0f);
        float sumB = float(0.0f);
        for (int i = (-1 * 15); i <= 15; i  )
            for (int j = (-1 *15); j <= 15; j  )
            {
                if (pos_x   j > 0 amp;amp; pos_y   i > 0 amp;amp; pos_x   j <= width amp;amp; pos_y   i <= height)
                {
                    sumR  = (float)device_input_imag[(pos_y   i) * width   (pos_x   j)]/255.0;
                    sumG  = (float)device_input_imag[(height   (pos_y   i)) * width   (pos_x   j)]/255.0;
                    sumB  = (float)device_input_imag[(height * 2   (pos_y   i)) * width   (pos_x   j)]/255.0;
                }   
            }
        sumR = sumR / (15 * 15);
        sumG = sumG / (15 * 15);
        sumB = sumB / (15 * 15);
        device_out_image[pos_y * width   pos_x] = (unsigned char)(sumR * 255.0);
        device_out_image[(height   pos_y) * width   pos_x] = (unsigned char)(sumG * 255.0) ;
        device_out_image[(height * 2   pos_y) * width   pos_x] = (unsigned char)(sumB *255.0 );
        if (device_out_image[pos_y * width   pos_x] > 255)
            device_out_image[pos_y * width   pos_x] = 255;
        if (device_out_image[(height   pos_y) * width   pos_x] > 255)
            device_out_image[(height   pos_y) * width   pos_x] = 255;
        if (device_out_image[(height * 2   pos_y) * width   pos_x] > 255)
            device_out_image[(height * 2   pos_y) * width   pos_x] = 255;   
    }
}
 

Попробуйте выполнить следующие шаги отладки:

Можете ли вы настроить выходное изображение на серую рампу, где все 3 канала одинаковы, а значение равно координате x? Как насчет тебя? Это проверяет правильность интерпретации выходного изображения, memcpy, компоновки канала, размеров и т. Д. Если это неверно, продолжайте копать и используйте внешний вид результата, чтобы помочь диагностировать любые проблемы. Например, если вы видите эффект рампы, но он смещен, у вас, вероятно, неправильный интервал между строками. Если результат не серый, вы неправильно понимаете, где расположены r, g и b. Сделайте шаг назад и установите только один канал за раз.

Можете ли вы скопировать входное изображение в выходное изображение, используя ту же структуру ядра, которая у вас есть? То есть закомментируйте цикл и просто установите выходное значение на входное значение. Это проверяет правильность интерпретации входных данных, memcpy и т.д.

Перейдите к ядру размером 3×3. Затем вернемся ко всему этому.

Несколько других советов:

Вам не должен понадобиться поплавковый аккумулятор. Поскольку изображение состоит из 8-разрядных значений, вы никогда не переполните 32-разрядный накопитель целых чисел, даже с гораздо большими ядрами (32 * 32 * 256 составляет 18 бит).

Внимательно просмотрите свой > vs >>=. 0-это нормально, поэтому вы хотите протестировать idx>=0 , но ширина не в порядке, поэтому вам нужно протестировать idx<=width-1 или, более идиоматично, idx<width .

Выполните зажим выходного диапазона перед назначением в ячейку памяти. Оптимизатор почти наверняка исправит это, но ваш код также будет выглядеть проще и его будет легче проверить, если вы зафиксируете временные значения sumR/SUM/SUM.