素人によるCUDAのお勉強。2

CUDAプログラミングを始める前のお話

CUDAのプログラムはGPU上でスタンドアロンで実行されるのではなく、CPUとGPUの両方を使っていることを念頭に置く必要がある。（それゆえ、CUDAプログラミングでは、GPUのアーキテクチャを深く理解しないといけない。）そのような異なるシステムが混在するのがCUDAプログラミングである。CPUは複雑な命令を実行でき、GPUは簡単な命令を大量のデータに実行することができる。このことを念頭に置くと、少ないデータ量ならば、CPUを使ったほうが遥かに効率的になることも少ない。また、CPUにはキャッシュが存在し、CPUからデータへのアクセスの遅延時間を減らす構造となっている。GPUでは、このキャッシュに相当するのがシェアードメモリである。シェアードメモリを如何に上手く使えるかが、よりよいGPUプログラミングへの近道となるだろう。

CUDAによるプログラミングの基礎

CUDAはCの拡張で、使用するライブラリ名も似ている。前の記事で紹介したように、CUDAプログラミグの流れは以下のようになる。その流れに関数を付け加えると

• GPU上でメモリの確保

cudaError_t cudaMalloc ( void** devPtr, size_t size )

• CPUからGPUへのデータのコピー

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

• CUDA カーネルを実行する。
• GPUからCPUへのデータのコピー

cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

• GPU上のメモリの解放

cudaMemcpy

この関数は、dstとsrcにGPUのメモリ、CPUのメモリを設定し、cudaMemcpyKindに何から何へのコピーかを指定してやる必要が有る。cudaMemcpyKindには以下の値が入る。

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

エラーの処理

cudaの関数は基本的に,cudaError_tを返すように設計されており、以下の関数でエラー内容を取得することができる。

char* cudaGetErrorString(cudaError_t error)

グリッド、ブロック、スレッド

グリッドの中にブロックがあり、ブロックの中にスレッドがある。初めて、CUDAを勉強した時に、私はここで混乱しました。しかも、ブロックやスレッドは、２次元、３次元になっているらしい...。なんのこっちゃ。とりあえず、GPU上での計算範囲を指定しているらしい。

そもそも、どうしてグリッド、ブロック、スレッドの概念が必要なのか

次の記事で書きます...。(参考文献の構成がそうなっているからです...。)

グリッド、ブロック、スレッドのインデックスを調べてみよう。

わからないなりにも先に進みましょう。

//checkDimension.cu
#include <cuda_runtime.h>                                                                                                   
#include <stdio.h>                                                                                                          
                                                                                                                            
__global__ void checkIndex(void){                                                                                           
    printf("threadIdx:(%d, %d, %d) blockIdx:(%d, %d, %d) blockDim:(%d, %d, %d) "                                            
           "gridDim:(%d, %d, %d)\n",threadIdx.x,threadIdx.y,threadIdx.z,                                                    
           blockIdx.x,blockIdx.y,blockIdx.z,blockDim.x,blockDim.y,blockDim.z,                                               
           gridDim.x,gridDim.y,gridDim.z);                                                                                  
}                                                                                                                           
int main(int argc, char **argv){                                                                                            
                                                                                                                            
    int nElem = 6;                                                                                                          
                                                                                                                            
    dim3 block (3);                                                                                                         
    dim3 grid ((nElem+block.x-1)/block.x);                                                                                  
                                                                                                                            
    printf("grid.x %d grid.y %d grid.z %d\n",grid.x,grid.y,grid.z);                                                         
    printf("block.x %d block.y %d block.z %d\n",block.x,block.y,block.z);                                                   
                                                                                                                            
    checkIndex <<<grid,block>>>();                                                                                          
                                                                                                                            
    cudaDeviceReset();                                                                                                      
                                                                                                                            
    return(0);                                                                                                              
}     

実行結果。

nvcc checkDimension.cu
./a.out
grid.x 2 grid.y 1 grid.z 1                                                                                                  
block.x 3 block.y 1 block.z 1                                                                                               
threadIdx:(0, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)                                                 
threadIdx:(1, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)                                                 
threadIdx:(2, 0, 0) blockIdx:(1, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)                                                 
threadIdx:(0, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)                                                 
threadIdx:(1, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)                                                 
threadIdx:(2, 0, 0) blockIdx:(0, 0, 0) blockDim:(3, 1, 1) gridDim:(2, 1, 1)  

前回の記事。oimokihujin.hatenablog.com

参考文献:

Professional CUDA C Programming