Плохие данные, поступающие от 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);
}

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

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

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

РЕДАКТИРОВАТЬ: Обратите внимание, нет ошибок или предупреждений во время компиляции со следующим make-файлом:

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

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

Answer 1

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

Ваша проверка ошибок при запуске ядра не работает. Вы можете рассмотреть правильную проверку ошибок cuda. Если у вас была правильная проверка ошибок, или вы запустили свой код с 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 потоков), что превышает максимальный предел в 1024 потока на блок.

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

Вы можете исправить эту ошибку, просто изменив 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;
}

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