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) {

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

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

  //xposは0～になっているので、offsetを足す
  int x = xpos + min(sx, ex);

  if (x < 0 || x >= width)return;
  if (x < min(sx,ex) || x >= max(sx, ex))return;

  int y = a * x + b;//xの時のy座標を求める

  if (y < 0 || y >= height)return;
  if (y < min(sy, ey) || y >= max(sy, ey))return;


  int pos = y * width + x;

  pimage[pos] = 1;//画素へ書き込み

}

/////
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 = 40; // 線分の始点・終点を設定
  int sy = 5;
  int ex = 10;
  int ey = 20;

  //書き込む画素数を算出
  int xlen = max(sx, ex) - min(sx, ex);

  // xlenを grids×blocksで表現
  int blocks = 512;
  int grids = (xlen+511) / blocks;


  dim3 block(blocks,1);
  dim3 grid(grids,1);

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

  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("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] );
      k++;
    }
    fprintf(fp, "\n");
  }

  fclose(fp);
}