CUDA копировать только копию изображения изображения

Question

Сегодня я пытался создать программу, которая копировала изображение с помощью графического процессора. Я создал простую программу, которая делает это. Чтобы загрузить изображение, я использую lodepng. Проблема заключается не в копировании через cudaMemcpy, потому что, когда я копирую изображение на GPU и обратно, он остается нетронутым, но когда я пытаюсь скопировать его в ядро, это не так. Не стесняйтесь задавать любые вопросы о моей проблеме.

Код:

#include <cuda_runtime.h> 
#include <device_launch_parameters.h> 
#include <iostream> 
#include <Windows.h> 
#include <math.h> 
#include <LodePNG\lodepng.h> 

const int BLOCK_WIDTH = 32; 


using namespace std; 

__global__ void expousure(unsigned char *in, unsigned char *out) 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int pitch = blockDim.x * gridDim.x; 
    int absIdx = x + y * pitch; 

    out[absIdx] = in[absIdx]; 

} 

void decode(std::vector<unsigned char>& image, const char* filename, int& width, int& height) 
{ 
    unsigned widthU, heightU; 
     //decode 
    unsigned error = lodepng::decode(image, widthU, heightU, filename); 

    width = int(widthU); 
    height = int(heightU); 

    //if there's an error, display it 
    if (error) std::cout << "decoder error " << error << ": " << lodepng_error_text(error) << std::endl; 

    //the pixels are now in the vector "image", 4 bytes per pixel, ordered RGBARGBA..., use it as texture, draw it, ... 
} 

void encodeAndSave(const std::vector<unsigned char>& inPixels, const char* filename, int width, int height) 
{ 
    std::vector<unsigned char> outEncoded; 

    unsigned error = lodepng::encode(outEncoded, inPixels, unsigned(width), unsigned(height)); 

    if (error){ 

     std::cout << "encoder error" << error << ": " << lodepng_error_text(error) << std::endl; 

     return; 
    } 

    lodepng::save_file(outEncoded, filename); 
} 

void encodeAndSave(unsigned char* inPixels, const char* filename, int width, int height) 
{ 
    std::vector<unsigned char> outEncoded; 

    unsigned error = lodepng::encode(outEncoded, inPixels, unsigned(width), unsigned(height)); 

    if (error){ 

     std::cout << "encoder error" << error << ": " << lodepng_error_text(error) << std::endl; 

     return; 
    } 

    lodepng::save_file(outEncoded, filename); 
} 


int main(int argc, char *argv[]) 
{ 

    // decode the image to image from filename 
    int width, height; 
    const char* filename = argc > 1 ? argv[1] : "C:/Users/Russell/Documents/Visual Studio 2013/Projects/Hello CUDA/Release/test.png"; 
    vector <unsigned char> h_image; 
    decode(h_image, filename, width, height); 

    unsigned char *d_in; 
    unsigned char *d_out; 

    cudaMalloc(&d_in, sizeof(unsigned char) * width * height * 4); 
    cudaMalloc(&d_out, sizeof(unsigned char) * width * height * 4); 

    cudaMemcpy(d_in, &h_image[0], sizeof(unsigned char) * width * height * 4, cudaMemcpyHostToDevice); 

    expousure<<<dim3(width/BLOCK_WIDTH, height/BLOCK_WIDTH, 1), dim3(BLOCK_WIDTH, BLOCK_WIDTH, 1) >>>(d_in, d_out); 

    unsigned char h_out[256 * 256 * 4]; 

    cudaMemcpy(h_out, d_out, sizeof(unsigned char) * width * height * 4, cudaMemcpyDeviceToHost); 

    // encode and save image from image to filename 
    vector <unsigned char> imageOUT; 
    const char* outname = "C:/Users/Russell/Documents/Visual Studio 2013/Projects/Hello CUDA/Release/testOUT.png"; 
    encodeAndSave(h_out, outname, width, height); 

} 

Исходное изображение: http://i.stack.imgur.com/Rx0mF.png

Выход изображения: http://i.stack.imgur.com/HLmPQ.png

Answer 1

Я хотел бы отметить несколько вещей:

1. Основной вопрос у вас есть то, что ваш массив потоков имеет размер, чтобы обеспечить 1 поток на пиксель, но поскольку каждый пиксель состоит из 4 байтов, а ваше ядро ​​копирует только один байт на поток, вы получаете только 1/4 копируемого изображения. Исправить это, в двух словах, можно было бы запустить в 4 раза больше потоков в x-размерности, чтобы учесть 4 байта на пиксель.

2. В любое время, когда у вас возникли проблемы с кодом CUDA, рекомендуется делать proper cuda error checking, хотя я не думаю, что это могло бы что-то придумать. В качестве быстрой проверки вы также можете запустить код CUDA с cuda-memcheck.

3. Лучше, если вы предоставите надлежащее MCVE. Такой код не зависит от внешних вещей, таких как lodepng.

4. Ваш код, как написано (с исправлением выше или выше), будет зависеть от того, что размеры изображения равномерно делятся на BLOCK_WIDTH. Это хорошая идея написать код, который не имеет этих зависимостей, и изменения включают в себя обеспечение того, что вы запускаете достаточно или больше, чем достаточно потоков в обоих измерениях во время запуска ядра, а затем включаете «проверку потока» в ядре убедитесь, что только действующие потоки выполняют какую-либо работу (в этом случае копирование).

Вот полный пример, который не зависит от lodepng, но демонстрирует подходящие исправления для пунктов 1,3 и 4 выше.

#include <iostream> 
#include <vector> 

const int BLOCK_WIDTH = 32; 

#define DUMMY_SIZE 256 

unsigned create_dummy_image(std::vector<unsigned char>& image, unsigned & widthU, unsigned &heightU, const char* filename){ 

    for (int i = 0; i < 4*DUMMY_SIZE; i++) 
    for (int j = 0; j < DUMMY_SIZE; j++) 
     image.push_back(j%8); 
    widthU = DUMMY_SIZE; 
    heightU = DUMMY_SIZE; 
    return 0; 
} 

unsigned dummy_encode(std::vector<unsigned char> &outEncoded, unsigned char *inPixels, unsigned width, unsigned height){ 

    for (int j = 0; j < height; j++) 
    for (int i = 0 ; i < 4*width; i++) 
     outEncoded.push_back(inPixels[(j*4*width)+i]); 
    return 0; 
} 

void dummy_save(std::vector<unsigned char> &outEncoded, const char * filename){ 

    for (int i = 0; i < outEncoded.size(); i++) 
    if (outEncoded[i] != (i%8)) {printf("mismatch at %d, was %d, should be %d\n", i, outEncoded[i], i%8); exit(1);} 
} 

using namespace std; 

__global__ void expousure(unsigned char *in, unsigned char *out, const int width, const int height) 
{ 
    int x = threadIdx.x + blockIdx.x * blockDim.x; 
    int y = threadIdx.y + blockIdx.y * blockDim.y; 
    int pitch = width*4; 
    int absIdx = x + y * pitch; 
    if ((x<(width*4)) && (y<height)) 
     out[absIdx] = in[absIdx]; 

} 


void decode(std::vector<unsigned char>& image, const char* filename, int& width, int& height) 
{ 
    unsigned widthU, heightU; 
     //decode 
    unsigned error = create_dummy_image(image, widthU, heightU, filename); 

    width = int(widthU); 
    height = int(heightU); 

    //if there's an error, display it 
    if (error) std::cout << "decoder error " << error << ": " << error << std::endl; 

    //the pixels are now in the vector "image", 4 bytes per pixel, ordered RGBARGBA..., use it as texture, draw it, ... 
} 

void encodeAndSave(unsigned char* inPixels, const char* filename, int width, int height) 
{ 
    std::vector<unsigned char> outEncoded; 

    unsigned error = dummy_encode(outEncoded, inPixels, unsigned(width), unsigned(height)); 

    if (error){ 

     std::cout << "encoder error" << error << ": " << error << std::endl; 

     return; 
    } 

    dummy_save(outEncoded, filename); 
} 


int main(int argc, char *argv[]) 
{ 

    // decode the image to image from filename 
    int width, height; 
    const char* filename = argc > 1 ? argv[1] : "C:/Users/Russell/Documents/Visual Studio 2013/Projects/Hello CUDA/Release/test.png"; 
    std::vector<unsigned char> h_image; 
    decode(h_image, filename, width, height); 

    unsigned char *d_in; 
    unsigned char *d_out; 

    cudaMalloc(&d_in, sizeof(unsigned char) * width * height * 4); 
    cudaMalloc(&d_out, sizeof(unsigned char) * width * height * 4); 

    cudaMemcpy(d_in, &h_image[0], sizeof(unsigned char) * width * height * 4, cudaMemcpyHostToDevice); 

    expousure<<<dim3((4*width/BLOCK_WIDTH)+1, (height/BLOCK_WIDTH)+1, 1), dim3(BLOCK_WIDTH, BLOCK_WIDTH, 1) >>>(d_in, d_out, width, height); 

    unsigned char h_out[DUMMY_SIZE * DUMMY_SIZE * 4]; 

    cudaMemcpy(h_out, d_out, sizeof(unsigned char) * width * height * 4, cudaMemcpyDeviceToHost); 

    // encode and save image from image to filename 
    vector <unsigned char> imageOUT; 
    const char* outname = "C:/Users/Russell/Documents/Visual Studio 2013/Projects/Hello CUDA/Release/testOUT.png"; 
    encodeAndSave(h_out, outname, width, height); 
    std::cout << "Success!" << std::endl; 
}