CUDAを使ってGPUプログラミングに挑戦してみた。 - のんびりしているエンジニアの日記

皆さんこんにちは
お元気ですか。私は家の近くに一風堂ができて感動しています。
本日はCUDAを使ったGPUプログラミングに挑戦します。

最近、ふと思うことがあります。
GPUをよく使いますが、GPUの心みたいなものがわからない状態です。
そこで、GPUプログラミングに挑戦し、コードを書きながらGPUのことを
知ることができたらと思って書いてみました。

GPUわからないとTwitter呟いていたら、
@dandelion1124さんにおすすめの書籍を紹介してもらいました！
GPUがどんな仕組みで動作しているのか非常にわかりやすいです。

CUDA C プロフェッショナルプログラミング impress top gearシリーズ

作者: John Cheng,Max Grossman,Ty McKercher
出版社/メーカー: インプレス
発売日: 2015/09/24
メディア: Kindle版
この商品を含むブログを見る

そもそもGPUって何？

第１回 GPUコンピューティングおよびCUDAについて | G-DEPより引用します。

GPUとはGraphics Processing Unit の略称で、その名の通りPCやワークステーションにおいて画像処理を担当する主要な部品の一つです。高速のVRAM（ビデオメモリ；グラフィックボード上のGPU専用メモリ）と接続され、グラフィクスシェーディングに特化したプロセッサが多く集まった構造を持っています。一つ一つのプロセッサの構造は単純なためその機能はCPUに比べて限定されたものですが、大量のデータを複数のプロセッサで同時かつ並列処理することができます。

近年、DeepLearingの演算で利用されることの多いのは最後にある大量のデータを
複数のプロセッサで同時かつ並列処理できる点でしょう。
これにより特に行列演算において、CPUより高速に演算を行えます。

ただし、GPUは条件分岐処理が含まれるとCPUより処理が遅くなる傾向があるとのこと。
条件分岐があると逐次処理になり、処理自身を並列にできないようです。

CUDAとは

第１回 GPUコンピューティングおよびCUDAについて | G-DEPより引用します。

CUDA（クーダ）とはCompute Unified Device Archtectureの略称で、半導体メーカーNVIDIA社が提供するGPUコンピューティング向けの統合開発環境です。プログラム記述、コンパイラ、ライブラリ、デバッガなどから構成されており、プログラム言語はC言語ベースに拡張を加えたものであるため、C言語によるプログラミングの経験があれば扱いやすくなっています。

GPUコンピューティング向けの環境であるようです。（そのまま）
CUDAはC言語ベースに拡張を加えたものであるため、C言語を扱ったことがある人であれば
簡単に実装することができるそう。

さっそく、CUDAを使ったGPUプログラミングに挑戦しましょう。

準備

GPUとCUDA、GPUドライバーを準備してください。

nonbiri-tereka.hatenablog.com

GPUでHello Worldを書く

まずは、GPUでHello Worldです。
次のソースコードを「hello_world.cu」のファイル名で作成しましょう。

#include <stdio.h>

int main( void ) {
  printf("Hello, World!\n");
  return 0;
}

そして、コンパイルします。コンパイラはgccではなくnvccを使いましょう。

nvcc hello_world.cu

nvccは「ソースの中のCPUで動く部分とGPUで動く部分を分け、CPU用コードをCコンパイラに渡し、GPU用コードをコンパイル（正確にはPTXコードに変換）する」コンパイラとのことです。
そのため、Cのみでも、実行可能となります。

さて、実行しましてみましょう。

$ ./a.out
Hello, World!

これでGPU計算の第一歩を踏み出しました（~~ぶっちゃけ、GPU使ってない~~）

GPU上で計算させてみる。

先程実行したのは出力をする簡単なものです。
GPUを使って計算させてみましょう。
次のコードは２つの配列を合計する計算をする処理を行います。

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

#define N 2000000000

__global__
void sum_of_array(float *arr1, float *arr2, float *arr3, int size){
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    arr3[i] = arr1[i] + arr2[i];
}

void initialize_array(float *arr, int size){
    for (int i = 0; i < size; i++){
        arr[i] = (float)rand();
    }
}

int main(void){
    float *arr1, *arr2, *arr3, *d_arr1, *d_arr2, *d_arr3;
    size_t n_byte = N * sizeof(float);

    arr1 = (float *)malloc(n_byte);
    arr2 = (float *)malloc(n_byte);
    arr3 = (float *)malloc(n_byte);

    initialize_array(arr1, n_byte);
    initialize_array(arr2, n_byte);
    initialize_array(arr3, n_byte);

    printf("start cudaMalloc\n");
    cudaMalloc((void**)&d_arr1, N);
    cudaMalloc((void**)&d_arr2, N);
    cudaMalloc((void**)&d_arr3, N);
    printf("finish cudaMalloc\n");

    printf("start cudaMemcpy\n");
    cudaMemcpy(d_arr1, arr1, n_byte, cudaMemcpyHostToDevice);
    cudaMemcpy(d_arr2, arr2, n_byte, cudaMemcpyHostToDevice);
    cudaMemcpy(d_arr3, arr3, n_byte, cudaMemcpyHostToDevice);
    printf("finish cudaMemcpy\n");

    printf("start kernel function\n");
    sum_of_array<<<(N+255)/256, 256>>>(d_arr1, d_arr2, d_arr3, n_byte);
    printf("finish kernel function\n");
    cudaMemcpy(arr3, d_arr3, n_byte, cudaMemcpyDeviceToHost);
}

GPU上で演算するには、次の4つの手順が必要になります。

GPUのメモリを確保する・・cudaMalloc
GPUのメモリにコピーする・・cudaMemcpy（cudaMemcpyHostToDevice）
計算を実行する・・カーネル関数
CPUのメモリに結果をコピーする・・cudaMemcpy（cudaMemcpyDeviceToHost）

まずは、GPUメモリを確保します。これはcudaMallocを使って行います。
次にGPUのメモリにCPUで確保したデータをコピーします。そして、カーネル関数と呼ばれる
関数で計算し、最後結果をCPUに書き戻します。

これにより、GPUを使った演算ができます。
CUDAにはまだまだライブラリがあるので、どんどん使ってみましょう！

参考文献

第１回 GPUコンピューティングおよびCUDAについて | G-DEP
GPGPUによるアプリケーションの高速化|PearLab, Utsunomiya Univ.