ノート/CUDA/マルチGPU

サイトトップ

ノート
データ処理
├マイニング
├バスケット解析
├バスケット解析をＲで
├図書貸出をＲで
├モデルクラスタリング
└
テキストマイニング
├テキストマイニングTM
├TMとシソーラス
├PubMedをTM
├TMとMeSH
├TMとNLTK
├テキストマイニングとtagger
├医薬品添付文書ＤＢ
├論文の処理1
├剽窃１
├Bing検索
└
分子進化学
├分子進化学
└
Pythonと論文アクセス
├Pythonを使ってみる
├PythonでPubMed
├Pythonで...続き
└
Macメモ
├Script

&ref(): File not found: "favicon_story.txt" at page "ノート/ノート";

自信ないので
├Winのバッチ
│├スクリプト入門
│├バッチファイル
│├SETコマンド
│├自分のメモ
│└
└

メニュー編集

最新の5件

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

マルチGPUの扱い方　（2010-08-24） †

参考資料 †

• Jason Sanders, Edward Kandrot: "CUDA By Example" (Addison-Wesley) Chapter 11 Section 3 下記の具体例はこの本の例から引用

• NvidiaのForumから
  • http://forums.nvidia.com/index.php?showtopic=70604　役に立つかも
  • http://forums.nvidia.com/index.php?showtopic=162645　あまり関係ない
  • https://www.wiki.ed.ac.uk/display/ecdfwiki/Use+multiple+GPU+devices+with+OpenMP+and+CUDA　CUDA + OpenMP というアプローチ。
    うちらでは未だOpenMPをCUDA下で使う環境を整えていない。
    ちなみに、このEDCF (Edinburgh Compute and Data Facility) ではGPGPU環境を整えているらしい　⇒　https://www.wiki.ed.ac.uk/display/ecdfwiki/Documentation　

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

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

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

というところ。１行目で0側のスレッドを生成し、２行目は今走っている自分自身が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;
    }

としている。

つまり、ホスト側で２つのスレッドが走り始めたことになる。どちらも同じ「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];
}

理解したこと：

• ホストCPU側でpthreadによってスレッドを起動する。この点でOpenMPを使ってもと同じと考えていいだろう。
  この時点で、どのように問題（やデータ）を分割するかを考慮済み。
• ホスト側スレッドのそれぞれにおいて、GPUデバイスを起動
  GPUデバイス側では、既にデバイス数分に分割されている問題を、GPU上のブロック・スレッドとして実行
• 結果は、デバイスごとに出た結果を、ホストスレッド全体で集める必要がある（当然）。