CUDA
訪問者数 1561      最終更新 2010-08-24 (火) 13:14:46

マルチGPUの扱い方 (2010-08-24)

参考資料

上記の本「CUDA By Example」の例題プログラムを読んでみる(プログラム全体はNVIDIA社のCopyright下にあり。学生勉強のため一部引用をお許しください。)

マルチGPUとなるために、pthreadを使って2つのスレッドを立てる。この例題では

CUTThread thread = start_thread( routine, &(data[0]) );
routine( &(data[1]) );

というところ。1行目で0側のスレッドを生成し、2行目は今走っている自分自身が1側のスレッドとしてプログラムroutineを呼び出す。

start_thread自体はcommon/book.hに置かれていて、基本的にWindowsとLinuxの場合によっての区分けをしていて、Linuxの場合はpthread_createを使って

    CUTThread start_thread(CUT_THREADROUTINE func, void * data){
        pthread_t thread;
        pthread_create(&thread, NULL, func, data);
        return thread;
    }

としている。

つまり、ホスト側で2つのスレッドが走り始めたことになる。どちらも同じ「routine」を実行するスレッドである。それぞれが

cudaMalloc( (void**)&dev_a, size*sizeof(float) )  デバイス側に領域aを確保
cudaMalloc( (void**)&dev_b, size*sizeof(float) ) 同 b
cudaMemcpy( dev_a, a, size*sizeof(float), cudaMemcpyHostToDevice ) aをデバイスへコピー
cudaMemcpy( dev_b, b, size*sizeof(float), cudaMemcpyHostToDevice ) 同 b
dot<<<blocksPerGrid,threadsPerBlock>>>( size, dev_a, dev_b, dev_partial_c );

  カーネル dot を起動。

cudaMemcpy( partial_c, dev_partial_c,
                             blocksPerGrid*sizeof(float),
                             cudaMemcpyDeviceToHost )

としている。

呼び出されるカーネル dot はベクトルの内積(dot product)を計算するもので、本の85ページにリストされているものと同じである。

__global__ void dot( int size, float *a, float *b, float *c ) {
    __shared__ float cache[threadsPerBlock];
    int tid = threadIdx.x + blockIdx.x * blockDim.x;  スレッド番号
    int cacheIndex = threadIdx.x;    キャッシュインデックス=スレッド番号

    float   temp = 0;
    while (tid < size) {
        temp += a[tid] * b[tid];    内積の加算
        tid += blockDim.x * gridDim.x; スレッド番号をblockDim*gridDim分進める
    }

    // set the cache values
    cache[cacheIndex] = temp;   内積の結果をcache[スレッド番号]に格納

    // synchronize threads in this block
    __syncthreads();

    // for reductions, threadsPerBlock must be a power of 2
    // because of the following code
    int i = blockDim.x/2;
    while (i != 0) {
        if (cacheIndex < i)
            cache[cacheIndex] += cache[cacheIndex + i]; cacheを集める。どういう風に?
        __syncthreads();
        i /= 2;
    }

    if (cacheIndex == 0)
        c[blockIdx.x] = cache[0];
}

理解したこと:

multigpu.jpg


添付ファイル: filemultigpu.jpg 210件 [詳細]

トップ   編集 凍結 差分 バックアップ 添付 複製 名前変更 リロード   新規 一覧 単語検索 最終更新   ヘルプ   最終更新のRSS
Last-modified: 2010-08-24 (火) 13:14:46 (2589d)