ビットマップをOpenGLで描画する

#include <GL/gl.h>
#include <GL/glut.h>
#include <cuda_runtime.h>
#include <stdio.h>

#define WIDTH 512
#define HEIGHT 512
#define IMAGE_SIZE_IN_BYTE (4 * WIDTH * HEIGHT)

// エラー用マクロ
#define EXIT_IF_FAIL(call)                                                     \
  do {                                                                         \
    cudaError_t retval = call;                                                 \
    if (retval != cudaSuccess) {                                               \
      printf("error in file %s line at %d: %s\n", __FILE__, __LINE__,          \
             cudaGetErrorString(retval));                                      \
      exit(1);                                                                 \
    }                                                                          \
  } while (0)

// グローバルのデータ(へのポインタ)をまとめたもの
struct DataBlock {
  unsigned char *bitmap;
  unsigned char *dev_bitmap;
  static DataBlock *get_data() {
    static DataBlock g_data;
    return &g_data;
  }
};

// カーネル関数: bitmapに適当に色を塗る
__global__ void kernel(unsigned char *bitmap) {
  int x = threadIdx.x + blockIdx.x * blockDim.x;
  int y = threadIdx.y + blockIdx.y * blockDim.y;
  int offset = x + y * blockDim.x * gridDim.x;

  bitmap[offset * 4 + 0] = (x / (WIDTH / 255));
  bitmap[offset * 4 + 1] = (y / (HEIGHT / 255));
  bitmap[offset * 4 + 2] = ((x + y) / (WIDTH + HEIGHT) / 255);
  bitmap[offset * 4 + 3] = 255;
}

// 描画用コールバック
static void draw() {
  DataBlock *data = DataBlock::get_data();
  glClearColor(0.0, 0.0, 0.0, 1.0);
  glClear(GL_COLOR_BUFFER_BIT);
  glDrawPixels(WIDTH, HEIGHT, GL_RGBA, GL_UNSIGNED_BYTE, data->bitmap);
  glFlush();
}

int main(int argc, char *argv[]) {
  DataBlock *data = DataBlock::get_data();

  // ビットマップをデバイス/ホストに確保
  data->bitmap = new unsigned char[IMAGE_SIZE_IN_BYTE];
  EXIT_IF_FAIL(cudaMalloc(&data->dev_bitmap, IMAGE_SIZE_IN_BYTE));

  dim3 threads(16, 16);                // 16x16 スレッド per グリッド
  dim3 grids(WIDTH / 16, HEIGHT / 16); // グリッドの数はthreadsとDIMから求まる

  // カーネル呼び出し
  kernel<<<grids, threads>>>(data->dev_bitmap);
  // 出来上がったビットマップをホストへコピー
  EXIT_IF_FAIL(cudaMemcpy(data->bitmap, data->dev_bitmap, IMAGE_SIZE_IN_BYTE,
                          cudaMemcpyDeviceToHost));
  // メモリの開放
  EXIT_IF_FAIL(cudaFree(data->dev_bitmap));

  // 画像の描画
  glutInit(&argc, argv);
  glutInitDisplayMode(GLUT_SINGLE | GLUT_RGBA);
  glutInitWindowSize(WIDTH, HEIGHT);
  glutCreateWindow("bitmap");
  glutDisplayFunc(draw);
  glutMainLoop();

  // メモリの開放
  free(data->bitmap);

  return 0;
}