Плохие данные, полученные от cudaMemcpy2D

Question

Если этот вопрос задан, прошу прощения, напишите мне в теме, пожалуйста!

В любом случае, я новичок в CUDA (я исхожу из OpenCL) и хотел попробовать создать изображение с ним. Соответствующий код CUDA является:

__global__ 
void mandlebrot(uint8_t *pixels, size_t pitch, unsigned long width, unsigned long height) { 
    unsigned block_size = blockDim.x; 
    uint2 location = {blockIdx.x*block_size, blockIdx.y*block_size}; 
    ulong2 pixel_location = {threadIdx.x, threadIdx.y}; 
    ulong2 real_location = {location.x + pixel_location.x, location.y + pixel_location.y}; 
    if (real_location.x >= width || real_location.y >= height) 
    return; 
    uint8_t *row = (uint8_t *)((char *)pixels + real_location.y * pitch); 
    row[real_location.x * 4+0] = 0; 
    row[real_location.x * 4+1] = 255; 
    row[real_location.x * 4+2] = 0; 
    row[real_location.x * 4+3] = 255; 
} 

cudaError_t err = cudaSuccess; 

#define CUDA_ERR(e) \ 
    if ((err = e) != cudaSuccess) { \ 
    fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n", cudaGetErrorString(err)); \ 
    exit(-1); \ 
    } 


int main(void) { 
    ulong2 dims = {1000, 1000}; 
    unsigned long block_size = 500; 
    dim3 threads_per_block(block_size, block_size); 
    dim3 remainders(dims.x % threads_per_block.x, dims.y % threads_per_block.y); 
    dim3 blocks(dims.x/threads_per_block.x + (remainders.x == 0 ? 0 : 1), dims.y/threads_per_block.y + (remainders.y == 0 ? 0 : 1)); 

    size_t pitch; 
    uint8_t *pixels, *h_pixels = NULL; 
    CUDA_ERR(cudaMallocPitch(&pixels, &pitch, dims.x * 4 * sizeof(uint8_t), dims.y)); 
    mandlebrot<<<blocks, threads_per_block>>>(pixels, pitch, dims.x, dims.y); 

    h_pixels = (uint8_t *)malloc(dims.x * 4 * sizeof(uint8_t) * dims.y); 
    memset(h_pixels, 0, dims.x * 4 * sizeof(uint8_t) * dims.y); 
    CUDA_ERR(cudaMemcpy2D(h_pixels, dims.x * 4 * sizeof(uint8_t), pixels, pitch, dims.x, dims.y, cudaMemcpyDeviceToHost)); 

    save_png("out.png", h_pixels, dims.x, dims.y); 

    CUDA_ERR(cudaFree(pixels)); 
    free(h_pixels); 

    CUDA_ERR(cudaDeviceReset()); 
    puts("Success"); 
    return 0; 
} 

save_png функция является обычной функцией полезности я создал для принятия блока данных и сохранения его в формате PNG:

void save_png(const char *filename, uint8_t *buffer, unsigned long width, unsigned long height) { 
    png_structp png_ptr = png_create_write_struct(PNG_LIBPNG_VER_STRING, NULL, NULL, NULL); 
    if (!png_ptr) { 
    std::cerr << "Failed to create png write struct" << std::endl; 
    return; 
    } 
    png_infop info_ptr = png_create_info_struct(png_ptr); 
    if (!info_ptr) { 
    std::cerr << "Failed to create info_ptr" << std::endl; 
    png_destroy_write_struct(&png_ptr, NULL); 
    return; 
    } 
    FILE *fp = fopen(filename, "wb"); 
    if (!fp) { 
    std::cerr << "Failed to open " << filename << " for writing" << std::endl; 
    png_destroy_write_struct(&png_ptr, &info_ptr); 
    return; 
    } 
    if (setjmp(png_jmpbuf(png_ptr))) { 
    png_destroy_write_struct(&png_ptr, &info_ptr); 
    std::cerr << "Error from libpng!" << std::endl; 
    return; 
    } 
    png_init_io(png_ptr, fp); 
    png_set_IHDR(png_ptr, info_ptr, width, height, 8, PNG_COLOR_TYPE_RGBA, PNG_INTERLACE_NONE, PNG_COMPRESSION_TYPE_DEFAULT, PNG_FILTER_TYPE_DEFAULT); 
    png_write_info(png_ptr, info_ptr); 
    png_byte *row_pnts[height]; 
    size_t i; 
    for (i = 0; i < height; i++) { 
    row_pnts[i] = buffer + width * 4 * i; 
    } 
    png_write_image(png_ptr, row_pnts); 
    png_write_end(png_ptr, info_ptr); 
    png_destroy_write_struct(&png_ptr, &info_ptr); 
    fclose(fp); 
} 

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

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

Я использую сетчатую сетчатую сетку с дискретной видеокартой NVIDIA GeForce GT 650M. Я могу запустить и вставить вывод в print_devices из кода образца cuda, если это необходимо.

EDIT: Примечание никаких ошибок или предупреждений во время компиляции с помощью следующей Makefile:

all: 
    nvcc -c mandlebrot.cu -o mandlebrot.cu.o 
    nvcc mandlebrot.cu.o -o mandlebrot -lpng 

и никаких ошибок во время выполнения.

Answer 1

Лучше, если вы предоставите полный код, который кто-то может копировать, вставлять, компилировать и запускать, не добавляя ничего или ничего менять. Отключение заголовков include не помогает, на мой взгляд, и делает ваш тестовый код в зависимости от библиотеки png, которую другие могут не иметь, также неэффективны, если вы хотите получить помощь.

Ошибка в проверке запуска ядра. Вы можете просмотреть proper cuda error checking. Если у вас была правильная проверка ошибок или был запущен ваш код с cuda-memcheck, вы обнаружите ошибку 9 при запуске ядра. Это недопустимая конфигурация. Если вы печатаете ваши blocks и threads_per_block переменные, вы увидите что-то вроде этого:

blocks: 2, 2 
threads: 500, 500 

Вы в настройки темы на блок, чтобы 500,500 здесь факт:

unsigned long block_size = 500; 
dim3 threads_per_block(block_size, block_size); 

Это незаконно, как вы запрашивают 500х500 потоков на блок (т.е. 250000 потоков), который превышает the maximum limit of 1024 threads per block.

Так что ваше ядро ​​не работает вообще, и вы получаете мусор.

Вы можете исправить эту ошибку довольно просто, изменив block_size определение:

unsigned long block_size = 16; 

После того, что есть еще вопрос, как вы неверно истолкованы параметры для cudaMemcpy2D.:

CUDA_ERR(cudaMemcpy2D(h_pixels, dims.x * 4 * sizeof(uint8_t), pixels, pitch, dims.x, dims.y, cudaMemcpyDeviceToHost)); 

Документация состояние для 5-го параметра:

ширина - ширина передачи матрицы (столбцы в байтах)

но вы прошли ширину в элементах (группы из 4 байтов), а не байтов.

Это зафиксирует, что:

CUDA_ERR(cudaMemcpy2D(h_pixels, dims.x * 4 * sizeof(uint8_t), pixels, pitch, dims.x*4, dims.y, cudaMemcpyDeviceToHost)); 

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

#include <stdio.h> 
#include <stdint.h> 

__global__ 
void mandlebrot(uint8_t *pixels, size_t pitch, unsigned long width, unsigned long height) { 
    unsigned block_size = blockDim.x; 
    uint2 location = {blockIdx.x*block_size, blockIdx.y*block_size}; 
    ulong2 pixel_location = {threadIdx.x, threadIdx.y}; 
    ulong2 real_location = {location.x + pixel_location.x, location.y + pixel_location.y}; 
    if (real_location.x >= width || real_location.y >= height) 
    return; 
    uint8_t *row = (uint8_t *)((char *)pixels + real_location.y * pitch); 
    row[real_location.x * 4+0] = 0; 
    row[real_location.x * 4+1] = 255; 
    row[real_location.x * 4+2] = 0; 
    row[real_location.x * 4+3] = 255; 
} 

cudaError_t err = cudaSuccess; 

#define CUDA_ERR(e) \ 
    if ((err = e) != cudaSuccess) { \ 
    fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n", cudaGetErrorString(err)); \ 
    exit(-1); \ 
    } 

int main(void) { 
    ulong2 dims = {1000, 1000}; 
    dim3 threads_per_block(16, 16); 
    dim3 remainders(dims.x % threads_per_block.x, dims.y % threads_per_block.y); 
    dim3 blocks(dims.x/threads_per_block.x + (remainders.x == 0 ? 0 : 1), dims.y/threads_per_block.y + (remainders.y == 0 ? 0 : 1)); 

    size_t pitch; 
    uint8_t *pixels, *h_pixels = NULL; 
    CUDA_ERR(cudaMallocPitch(&pixels, &pitch, dims.x * 4 * sizeof(uint8_t), dims.y)); 

    printf("blocks: %u, %u\n", blocks.x, blocks.y); 
    printf("threads: %u, %u\n", threads_per_block.x, threads_per_block.y); 
    mandlebrot<<<blocks, threads_per_block>>>(pixels, pitch, dims.x, dims.y); 

    h_pixels = (uint8_t *)malloc(dims.x * 4 * sizeof(uint8_t) * dims.y); 
    memset(h_pixels, 0, dims.x * 4 * sizeof(uint8_t) * dims.y); 
    CUDA_ERR(cudaMemcpy2D(h_pixels, dims.x * 4 * sizeof(uint8_t), pixels, pitch, dims.x*4, dims.y, cudaMemcpyDeviceToHost)); 

// save_png("out.png", h_pixels, dims.x, dims.y); 
    for (int row = 0; row < dims.y; row++) 
    for (int col = 0; col < dims.x; col++){ 
     if (h_pixels[(row*dims.x*4) + col*4 ] != 0) {printf("mismatch 0 at %u,%u: was: %u should be: %u\n", row,col, h_pixels[(row*dims.x)+col*4], 0); return 1;} 
     if (h_pixels[(row*dims.x*4) + col*4 +1] != 255) {printf("mismatch 1 at %u,%u: was: %u should be: %u\n", row,col, h_pixels[(row*dims.x)+col*4 +1], 255); return 1;} 
     if (h_pixels[(row*dims.x*4) + col*4 +2] != 0) {printf("mismatch 2: was: %u should be: %u\n", h_pixels[(row*dims.x)+col*4 +2], 0); return 1;} 
     if (h_pixels[(row*dims.x*4) + col*4 +3] != 255) {printf("mismatch 3: was: %u should be: %u\n", h_pixels[(row*dims.x)+col*4 +3 ], 255); return 1;} 
     } 
    CUDA_ERR(cudaFree(pixels)); 
    free(h_pixels); 

    CUDA_ERR(cudaDeviceReset()); 
    puts("Success"); 
    return 0; 
} 

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