素人によるCUDAのお勉強。

どうして並列計算をするのか

　現在のCPUのクロック周波数は、ほぼ頭打ちです。クロック周波数を増やそうとすると、発熱が起こり、その熱を処理できないためだそうです。では、これからは計算速度をどのように向上させるのかというと、その答えの一つが並列化です。大きな問題(データ数が大きい)に対しては、データを分割することによって異なるプロセスにおいて計算を実行することができます。単純に、系のサイズNに対して、プロセッサーを４つ使うと、time(N)/4になるはずです。（もちろんそんなことはありません。）以上より、実行速度が早いプログラムを書こうとすると並列計算は避けて通れないものだと思います。

どうしてCUDAなのか

GPUがスパコンに比べて安い、これに尽きます。その代わり、大規模な計算をしようとすると、１つのGPUに系が収まり切らず、実行速度が落ちます。この場合、TUBAMEみたいにGPUのクラスター組み、大規模な計算をチャレンジするような方法がとられます。

CUDAによる並列化

• Task parallelism
  • 複数の関数をコアに割り当て、同時実行させる。
• Data parallelism
  • データを分解し、複数のデータを同時に処理する。

CUDAによる並列化は、主にData parallelismにあります。よって、どのようにデータを分解するかが重要になります。

データの分解

4x4の行列で例を示します。

• Block partition:
  • 2x2の正方行列を４つ生成。
• Cyclic partition:
  • 1x4の行列を４つ生成する。

CUDAのAPI

• CUDA Driver API
  • 低レベルのAPI
• CUDA Runtime API
  • CUDA Driver APIに比べ、高レベルのAPI

Hello, World

とりあえず、Hello World

#include <stdio.h>  
// hello.cu                                                                                                                            
                                                                                                                                                
__global__ void helloFromGPU(void){                                                                                                             
    printf("Hello World from GPU\n");                                                                                                           
}                                                                                                                                               
//__global__はCPUで呼び出し、GPUで実行する指示句                                                                                                                                           
int main(void){                                                                                                                                 
    printf("Hello World from CPU\n");                                                                                                           
    helloFromGPU<<<1,10>>>();  
    //10スレッドを生成                                                                                                                 
    cudaDeviceReset();
    //このプログラムで使用したGPUのリソースを解放。解放する前に同期を取っている。この一文を消して、実行してみてください。                                                                                                                          
    // cudaDeviceSynchronize();
}   

コンパイル

nvcc hello.cu

実行結果

./a.out
Hello World from CPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU                                                                                                                            
Hello World from GPU    

CUDAプログラムの流れ

• GPU上でメモリの確保
• CPUからGPUへのデータのコピー
• CUDA カーネルを実行する。
• GPUからCPUへのデータのコピー
• GPU上のメモリの解放

Professional CUDA C Programming