CUDAその３

研究室にNIVIDA Tesla　C1060が入ったやつが導入されたのでぼちぼちCUDAをいじってみる。

TeslaはGPUのくせに、ビデオ出力の機能がない。ひたすら数値演算に特化している。とりあえず、C1060だとマルチプロセッサが240あってしかも各プロセッサに8つのコアが含まれているらしい。超並列処理とかいっちゃってるわけだ。
GPGPUのウリはこの並列化らしい。各プロセッサの8つのコアはSIMT(single instruction multi thread)という形式で同じ内容の手続きを複数のスレッドに分けて実行する（SIMDと似ているがこっちのほうがややこしい処理を並列化できるんか）。スレッドは8*4=32が基本単位になっていて、32スレッドを1ワープと数えるらしい。変なの。

母体となるコードはCPU上で実行される。コードのなかでも並列化が肝になってくるところをGPU上で行うのが普通。GPU側で処理される手続きをカーネルと呼ぶ。
ふつーカーネルで行われる手続きは並列化される。並列化の基本単位はスレッドなんだが、CUDAではスレッドの塊とさらにその塊の集まったモノを定義する。スレッドの塊はスレッドブロック、スレッドブロックの集まりをグリッドという。ひとつのカーネルに対して、ひとつのグリッドが割り当てられる。スレッドブロックのインデックスは2次元まで、スレッドそのもののインデックスは3次元までとることができる。こんなに入れ子な仕組みにするのはたぶんGPUのハード側の特性（マルチプロセッサが並列＋そのなかのコアが並列）になってるからだろう。それがメモリの階層性にも反映されている。メモリはグローバルメモリとシェアードメモリの2つに大別される。グローバルメモリはグリッド内すべてで共有であり、シェアードメモリはブロック内のみで共有される。シェアードメモリのほうが圧倒的にアクセスが早い。

CUDAのプログラムの大まかな流れは

CPUサイドでデータ処理
↓
GPU上でメモリ確保（cudaMalloc）
↓
GPUへデータ転送　(cudaMemorycpy)
↓
GPUで計算（aaa<<>>(x,y,z)）
↓
GPUからデータ転送(cudaMemorycpy)

こんな感じ。めっちゃ簡単な一次元の配列を足し算するだけのをつくってみた。バージョンはCUDA v2.0。2.0になってちょこちょこ変更があったらしく、Web上の旧バージョンのサンプルとはヘッダファイルの名前が変わってたりというトラップあり。要注意。

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

#define N (32)

__global__ void wa ( int *array1, int *array2, int len)
{
    int i;
    for (i=0; i<len ;i++)
     array1[i]= array1[i]+array2[i];
    return;
}

int main (int argc,char *argv[])
{
    int i;
    int arrayH1[N];
    int arrayH2[N];
    int arrayoutput[N];
    
    int* arrayD1;
    int* arrayD2;
  
    size_t array_size;
        printf("input array \n"); 
    for (i=0; i<N;i++) arrayH1[i]=i;
    for (i=0; i<N;i++) arrayH2[i]=i*2;
    
    printf("input of H1 \n");
    for (i=0;i<N;i++)
    
    printf("%d\n",arrayH1[i]);
    
    printf("input of H2\n");
    for (i=0;i<N;i++)
    
    printf("%d\n",arrayH2[i]);


    array_size = sizeof(int) * N;
    
    cudaMalloc( (void **) &arrayD1, array_size);
    cudaMalloc((void **) &arrayD2, array_size);
    
    cudaMemcpy(arrayD1,arrayH1,array_size,cudaMemcpyHostToDevice);
        cudaMemcpy(arrayD2,arrayH2,array_size,cudaMemcpyHostToDevice);  
    wa <<<dim3(1,1),dim3(1,1,1)>>>(arrayD1,arrayD2,N);
    
    cudaMemcpy(arrayoutput,arrayD1,array_size,cudaMemcpyDeviceToHost);
    
    printf("output\n");
    for (i=0;i<N;i++)
    
    printf("%d\n",arrayoutput[i]);
    
    
    return 0;
    
}

なんか(void**)が？？だったけど、CUDAではこういう型もあるっぽい。void*はvoid型ポインタだからその拡張か？C言語自体あんまりしらんのでなかなか目新しいのが満載だ。今回つかったのは単純のためまったく並列化もなんもせんと1ブロック1スレッド。さて次は並列化を学習しよう。