CUDA | Basics

##Architecture
https://qiita.com/Keech/items/381495ed90e012d69f1a#cuda%E3%81%AE%E3%82%A2%E3%83%BC%E3%82%AD%E3%83%86%E3%82%AF%E3%83%81%E3%83%A3

##GPUメモリのアロケーション/領域解放
cudaMalloc(void ** pointer, size_t nbytes) //gpu(デバイス)メモリ確保
cudaMemset(void * pointer, int value, size_t count)
cudaFree(void* pointer)

memorySet

int n = 1024;
int nbytes = 1024*sizeof(int);
int *d_a = 0;
cudaMalloc( (void**)&d_a, nbytes );
cudaMemset( d_a, 0, nbytes);
cudaFree(d_a)

##データコピー

cudaMemcpy(void *dst, void *src, size_t nbytes,enum cudaMemcpyKind direction);
directionでsrcとdstの場所(ホスト、デバイス)を指定
CPUスレッドをブロック：コピー完了後に戻る
以前のCUDAコールが完了するまでコピーを開始しない

enum cudaMemcpyKind{
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice
}

copyData

int main(void)
{
  float *a_h, *b_h; // host data
  float *a_d, *b_d; // device data
  int N = 14, nBytes, i ;
  nBytes = N*sizeof(float); // needed memory space
  a_h = (float *)malloc(nBytes);
  b_h = (float *)malloc(nBytes); // preserve data on host
  cudaMalloc((void **) &a_d, nBytes);
  cudaMalloc((void **) &b_d, nBytes); // preserve data on device
  for (i=0, i<N; i++) a_h[i] = 100.f + i; // initialize host data with 100.f~113.f
  cudaMemcpy(a_d, a_h, nBytes, cudaMemcpyHostToDevice); // copy from host to device
  cudaMemcpy(b_d, a_d, nBytes, cudaMemcpyDeviceToDevice); // copy from device to device
  cudaMemcpy(b_h, b_d, nBytes, cudaMemcpyDeviceToHost); // copy from device to host
  for (i=0; i< N; i++) assert( a_h[i] == b_h[i] ); // check if b_h is successfully initialized with above copying
  free(a_h); free(b_h); cudaFree(a_d); cudaFree(b_d); // free the host and device memory
  return 0;
}

##カーネル関数
GPUで行う処理は、カーネル関数に記述する。
カーネル関数の呼び出しは以下のように記述する。
https://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Basics_PartII_jp.pdf
###カーネル関数名
<<<グリッド数(dim3) , スレッド数(dim3)>>>()
kernel<<< dim3 dG, dim3 dB >>>(…)
####global、__device__の関数からアクセスできる自動定義の変数
dim3 gridDim; ブロックによるグリッドの次元(最大2D)
dim3 blockDim; スレッドによるブロックの次元
dim3 blockIdx; グリッド内のブロックのインデックス
dim3 threadIdx; ブロック内のスレッドのインデックス

##Concurrent computation 並列計算
###データの並列
https://www.gsic.titech.ac.jp/supercon/main/attwiki/index.php?plugin=attach&pcmd=open&file=gpu-prog16-1ow.pdf&refer=SupercomputingContest2016

GPU は１つの中に Ｎ 個もの演算装置（計算コア）がいる

xxx.cu

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

カーネル関数(GPU上で実行される関数)

__global__ void mm_gpu(double *A, double *B, double *C, int n)
{
    int i, j, k;
    i = blockIdx.y * blockDim.y + threadIdx.y; // 自分のY方向の背番号を計算. i列
    j = blockIdx.x * blockDim.x + threadIdx.x; // 自分のX方向の背番号を計算. j行
    // 自分の背番号から担当する (i,j) を決める

    if (i >= n || j >= n) return; // 行列からはみ出す部分は計算しない

    for (k = 0; k < n; k++) { // 総和を計算する部分. 背番号にしたがって、Cijを一個だけ計算する
        C[i*n+j] += A[i*n+k] * B[k*n+j]; // C[i][k] += A[i][k] * B[k][j]に相当
    }
}

ホスト関数(CPU上で実行される関数)

// GPUに対してデータ転送やGPUカーネル関数呼び出しを実行
int main(int argc, char *argv[])
{
    int i, j, n;
    double *A, *B, *C; // ホストメモリ用のポインタ
    double *DA, *DA, *DC; // デバイスメモリ用のポインタ

    n = atoi(argv[1]); // 行列の大きさ (atoi:整数値の文字列型データをint型に変換する)

    // A, B, Cのためにホストメモリを確保
    A = (double *)malloc(sizeof(double)*n*n);
    B = (double *)malloc(sizeof(double)*n*n);
    C = (double *)malloc(sizeof(double)*n*n);

    // A, Bの内容を設定し、Cをゼロクリア
    // このプログラムでは、二次元配列でなく「サイズn*nの一次元配列」にした
    for (i = 0; i < n; i++) {
        for (j = 0; j < n; j++) {
            A[i*n+j] = 3.0; // ここでのA, Bの内容はいい加減
            B[i*n+j] = 0.1; //に決めた．（あくまで例なので）
            C[i*n+j] = 0.0;
    } }

    // A, B, Cのためにデバイスメモリを確保
    cudaMalloc((void**)&DA, sizeof(double)*n*n);
    cudaMalloc((void**)&DB, sizeof(double)*n*n);
    cudaMalloc((void**)&DC, sizeof(double)*n*n);

    // A, B, Cの内容を、ホストメモリからデバイスメモリへコピー
    // 注意! ホスト関数(CPU側)からは、DA, DB, DCの中身には触れない。
    // DA[i*n+j] = … などしてはダメ！
    cudaMemcpy(DA, A, sizeof(double)*n*n, cudaMemcpyHostToDevice);
    cudaMemcpy(DB, B, sizeof(double)*n*n, cudaMemcpyHostToDevice);
    cudaMemcpy(DC, C, sizeof(double)*n*n, cudaMemcpyHostToDevice);

    // GPUカーネル関数を呼び出す!! 約n*n個のスレッドを使う.それぞれのスレッドには、二次元の背番号をつけます
    mm_gpu<<<dim3((n+BS‐1)/BS, ((n+BS‐1)/BS)), dim3(BS, BS)>>>(DA, DB, DC, n);

    // 結果のCを、デバイスメモリからホストメモリへコピー
    cudaMemcpy(C, DC, sizeof(double)*n*n, cudaMemcpyDeviceToHost);
    // Cを出力などに利用(略)
    return 0;
}

1.データをスレードによって区切る