[
トップ
] [
新規
|
一覧
|
単語検索
|
最終更新
|
ヘルプ
]
開始行:
[[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上のブロック・スレッドとして実行
-結果は、デバイスごとに出た結果を、ホストスレッド全体で集める必要がある（当然）。
&ref(multigpu.jpg);
終了行:
[[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上のブロック・スレッドとして実行
-結果は、デバイスごとに出た結果を、ホストスレッド全体で集める必要がある（当然）。
&ref(multigpu.jpg);
ページ名:
![[Toho University]](image/CLogo_Kihon_process.jpg "[Toho University]")
![[YYLAB]](image/yamanouc2.png "[YYLAB]")
