[[CUDA>ノート/CUDA]]~
訪問者数　&counter();　　　　　　最終更新　&lastmod();~

**マルチGPUの扱い方　（2010-08-24） [#n4786eb8]

***参考資料 [#zb0fd0b6]
-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下にあり。学生勉強のため一部引用をお許しください。） [#de225a87]

マルチ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上のブロック・スレッドとして実行
-結果は、デバイスごとに出た結果を、ホストスレッド全体で集める必要がある（当然）。

&multigpu.jpg
&ref(multigpu.jpg);
トップ   編集 差分 バックアップ 添付 複製 名前変更 リロード   新規 一覧 単語検索 最終更新   ヘルプ   最終更新のRSS