• トップ
  • G-DEPについて
  • ご購入ガイド
  • サポート
  • お問い合わせ

G-DEPトップ  >  第8回 プログラミング演習 ~基礎編~

第8回 プログラミング演習 ~基礎編~

第8回 「プログラミング演習 ~基礎編~」

<< 第7回   |   目次に戻る   |   第9回>>

 第8回はプログラミング演習の基礎編ということで、今まで学んできた知識を実際の例題を通して確認しながら、CUDAプログラミングの基礎を身に付けていきたいと思います。

プログラミングスキルを身に付けるには、やはり自分で手を動かしてミスを繰り返しながら覚えていくのが一番の近道です。

今回の目的はズバリ、「CUDAの文法やプログラミングスタイルに慣れること」です。高速化(並列化)などの高度な計算は一切行っていません!とりあえず、CUDAがどんなものか見て慣れることから始めましょう。

 

それでは早速例題から。

 

Ex.1 "Hello, World!"


 第4回のおさらいになりますが、プログラミング言語の最初の問題といえば”Hello, World!”です。

まずは簡易的な表現から見ていきましょう。

#include <stdio.h>

__global__ void
    kernel( void ) {
}


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

上は第4回でも紹介した例題ですが、通常のC言語では見られないCUDAで拡張された部分を赤文字で示しています。

__global__ void kernel( void ) { }
 ここではkernel()という関数が__global__で修飾されており、この関数がホストではなくデバイスで実行されるためのものとしてコンパイルされます。
 
kernel<<<1,1>>>();
 これはmain()関数におけるkernel()関数の呼び出しで、<<<1,1>>>は以前のコラムで説明した様にブロックとスレッドのサイズを指定しています。(ブロック:1、スレッド:1)

 実際にはここで書いたCUDA拡張部分は「kernel関数の指定とホストからの呼び出し」という動作はしますが、Hello, World!の記述には関与していません。コンパイルして実行してみれば分かるかと思いますが、この2つの部分は無くても実行できます。その場合通常のC言語で作られた"Hello, World!"となるわけです。

そこで今度はGPUで文字列を形成し、それをホスト側に移して表示する方法を見てみましょう。

#include <stdio.h>

__device__ void
device_strcpy(char *dst, const char *src)
{
    while (*dst++ = *src++);
}


__global__ void kernel(char *A)
{
    device_strcpy(A, "Hello, World!");

}

int main() {
   char *d_hello;
   char hello[32];

   cudaMalloc((void**)&d_hello, 32);

   kernel<<<1,1>>>(d_hello);

   cudaMemcpy(hello, d_hello, 32, cudaMemcpyDeviceToHost);

   cudaFree(d_hello);


   puts(hello);
}

先ほど見られた拡張部分を茶色、新たな拡張部分を赤色で表示しています。

それではひとつずつ見ていきましょう。

① #include <stdio.h>
  
② __device__ void
  device_strcpy(char *dst, const char *src)
  {
      while (*dst++ = *src++);
  }
  
② __global__ void kernel(char *A)
  {
      device_strcpy(A, "Hello, World!");
  }

①通常のC言語と同じく標準I/Oをインクルードします。今回は使用していませんが、GPU用の特別なライブラリを用いるときは、同様にインクルードしてあげます。

②ここでデバイスメモリ上に"Hello, World!"という文字列を生成します。device_strcpyという関数を__device__識別子によってデバイス側から呼び出しデバイス側で実行される関数として、kernelを__global__識別子によってホスト側から呼び出しデバイス側で実行される関数として指定しています。GPUには標準ライブラリが存在しないため、strcpyを用意して__device__によりデバイス側で実行するように指定してあげることでC言語と同様に扱うことができます。

③ int main() {
④  char *d_hello;
⑤  char hello[32];

    cudaMalloc((void**)&d_hello, 32);

⑦    kernel<<<1,1>>>(d_hello);

⑧  cudaMemcpy(hello, d_hello, 32, cudaMemcpyDeviceToHost);

⑨  cudaFree(d_hello);

⑩   puts(hello);
  }

③ int main() {

 メインとなる関数です。

⑤char hello[32];

 ホストメモリを確保しています。結果を格納するのに十分な配列サイズを指定します。

cudaMalloc((void**)&d_hello, 32);

 デバイスメモリ(グローバルメモリ)を確保しています。引数は、cudaMalloc(確保したメモリのアドレスを格納するポインタ, 確保するメモリのサイズ); となります。

kernel<<<1,1>>>(d_hello);

 kernel()関数の呼び出し。

cudaMemcpy(hello, d_hello, 32, cudaMemcpyDeviceToHost);

 デバイスメモリに生成された文字列データをホスト側にメモリコピーしています。引数は、cudaMemcpy(書き込み先アドレス, 読み込み元アドレス, データサイズ, 転送方向); となります。

cudaFree(d_hello);

 デバイスメモリを解放してあげます。これを忘れてしまうと、C言語malloc同様にメモリリークが発生して、プログラム起動中徐々にメモリを圧迫していくことがありますので必ず記述してください。

⑩puts(hello);

 結果の出力です。このとき結果の文字列データはメインメモリにあるので、通常の出力方法で記述できます。

 

それでは、ここまでを踏まえて簡単な演習です。

演習1


「こんにちは、世界!私は、GPUコンピューティングを始めました。」という文章を、CUDAを使って表示するプログラムを作ってみてください。

ヒント:上のプログラムで"Hello, World!"の文章を置き換えただけでは上手く表示されません。どこかを変える必要があります。
 
答えはコチラ(別ウィンドウで開きます)

 

 

 

Ex.2 "数値データの受け渡し"


 Ex.1では文字列データの受け渡しを行いました。もちろん、数値計算および数字データの受け渡しも可能です。

2つ目の例では、簡単な数値計算を通して数値データの受け渡しに関するプログラムを見てみましょう。

#include <stdio.h>

__global__ void add (int a, int b, int *c) {
    *c = a+ b;
}


int main( void) {
    int c;
    int *dev_c;

    cudaMalloc( (void**)&dev_c, sizeof(int) );

    add<<<1,1>>>( 12, 2000, dev_c);

    cudaMemcpy( &c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree( dev_c);

    printf( "12 + 2000 = %d\n", c);

    return 0;
}

Ex.1と共通しているのは、

  • cudaMallocによりデバイス側のメモリ確保
  • カーネル関数addをホスト側で呼び出し、デバイス側で実行
  • 計算結果をcudaMemcpyによりホスト側に返す
  • cudaFreeでデバイスメモリの解放

add関数の( )内は引数です。呼び出す際に数値代入することでデバイス側にも反映されます。

cudaMemcpy(○, ○, ○, HostToDevice)を使って数値を一つずつ送ることもできます。また、DeviceToDeviceでデバイス内での数値の移動を行うこともできます。

Ex.1、Ex.2いずれも計算部分とメモリの確保・解放、数値の移動に関する拡張が増えただけで、C言語をGPUコード化することができています。

 

では早速、通常のC言語による簡単な計算問題を、上の例題を参考にGPUコード化してみましょう。

演習2


自分の今年(2012年)の満年齢を入力し、何年生まれかを答えるプログラムです。

以下がC言語による(CPU処理のみ)のコードになります。

#include <stdio.h>

int main()
{
    int age, birth;

    printf("あなたは今年で何歳になりますか?");
    scanf("%d", &age);

    birth = 2012 - age;

    printf("すると、%d生まれですね。\n", birth);

    return 0;
}
 
ヒント:GPUコード化させるのはる計算部分とそれに関するデータ転送部分だけで構いません。
 
答えはコチラ(別ウィンドウで開きます)

 

 

 

今回は基礎編ということで、とりあえずここまでとしましょう。C言語をすでにある程度習得されてる方にとっては非常に簡単でつまらなかったかもしれません。

ということで、次回の演習ではGPU計算の速さを感じられるプログラムについて見ていきます。

 

<< 第7回   |   目次に戻る   |   第9回>>


 

(執筆 G-DEP Associate Research Engineer 東京大学大学院工学系研究科 岡安優)