來自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獨立顯卡的MBP視網膜。如果需要，我可以從cuda示例代碼運行並粘貼輸出到print_devices。

編輯：請注意，用下面的Makefile編譯過程中沒有錯誤或警告：

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

，並在運行時沒有錯誤。

Answer 1

如果你提供了一個完整的代碼讓別人可以複製，粘貼，編譯和運行，而無需添加任何內容或更改任何內容，那麼在我看來，剝離包含頭文件並沒有什麼幫助，並且使測試代碼如果你需要幫助，依賴於其他人可能沒有的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); 

這是非法的，因爲你正在請求每塊500x500線程（即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)); 

的第五個參數的文件中指出：

寬度 - 矩陣轉移的寬度（以字節爲單位列）

，但你已經通過在元素的寬度（組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; 
} 

注意上面的代碼是一個完整的您可以複製，粘貼，編譯和運行代碼。