CUDAで直線を書くことを考える（２）（書いてみる）

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

///////////////////////////////////////////////
// GPU側 ////////////////////////////////////// 
__global__ void line(
  unsigned char* pimage,
  const int width,
  const int height,
  const int sx,
  const int sy,
  const int ex,
  const int ey) {

  //アクセス法
  //このスレッドが担当する画素の位置を二次元座標で求める
  int xpos = blockIdx.x * blockDim.x + threadIdx.x;
  int ypos = blockIdx.y * blockDim.y + threadIdx.y;


  int pos = ypos * width + xpos;
  if (xpos >= width || ypos >= height)return;
  if (pos < 0 || pos >= width * height)return;

  double a = (ey - sy) / double(ex - sx);//傾き
  double b = sy - a * sx;//切片
  double y = int(a * xpos + b);

  if (ypos < min(sy, ey))return;
  if (xpos < min(sx, ex))return;
  if (ypos > max(sy, ey))return;
  if (xpos > max(sx, ex))return;
  
  if (y == ypos) {
    pimage[pos] = 1;
  }
  else {
    pimage[pos] = 0;
  }

}
/////
void pbmP1_Write(const char* const fname, const int width, const int height, const unsigned char* const p);

int main()
{
  int width = 100;
  int height = 50;

  int sx = 98; // 線分の始点・終点を設定
  int sy = 48;
  int ex = 1;
  int ey = 3;

  // 最小単位 合計で512未満
  //10×10の領域に分けて計算する
  dim3 block(10, 10);

  //グリッド数
  // スレッド数はblock数×grid数なので(10x100)x(10x100)=1000x1000
  // 多すぎるが今回はそこまで考えない
  dim3 grid(100, 100);

  unsigned char* p_cpu = new unsigned char[width*height];
  for (int i = 0; i < width*height; i++) {
    p_cpu[i] = 1;
  }

  unsigned char* p_gpu;//GPU側メモリ確保
  cudaError_t ce = cudaMalloc((void**)&p_gpu, width*height);
  printf("error:%d\n", ce);
  line << <grid, block >> > (p_gpu, width, height,sx,sy,ex,ey);

  //GPU→CPU側へメモリコピー
  cudaMemcpy(p_cpu, p_gpu, width*height, cudaMemcpyDeviceToHost);
  cudaFree(p_gpu);


  ///////////////////////////////////////////////
  // 結果出力

  pbmP1_Write("C:\\dev\\cudaline.pbm", width, height, p_cpu);

  delete[] p_cpu;

  return 0;
}

//! @brief PBM(1byte,テキスト)を書き込む
//! @param [in] fname ファイル名
//! @param [in] width 画像の幅
//! @param [in] height 画像の高さ
//! @param [in] p 画像のメモリへのアドレス
//! @details 1画素1Byteのメモリを渡すと、0,1テキストでファイル名fnameで書き込む
void pbmP1_Write(const char* const fname, const int width, const int height, const unsigned char* const p) { // PPM ASCII

  FILE* fp = fopen(fname, "wb");
  fprintf(fp, "P1\n%d\n%d\n", width, height);

  size_t k = 0;
  for (size_t i = 0; i < (size_t)height; i++) {
    for (size_t j = 0; j < (size_t)width; j++) {
      fprintf(fp, "%d ", p[k] ? 1 : 0);
      k++;
    }
    fprintf(fp, "\n");
  }

  fclose(fp);
}