CUDAで行列演算：乗算(シェアードメモリ使用版)
をテンプレートにして作成
[
トップ
|
新規
|
一覧
|
検索
|
最終更新
|
ヘルプ
]
開始行:
*CUDAで行列演算 その3 乗算(シェアードメモリ使用版) [#g2d9...
[[CUDAで行列演算：乗算]]で行列の積をやりましたが，メモリ...
CUDAで使えるメモリについて勉強するために，シェアードメモ...
**[[CUDAで行列演算：乗算]]の問題点 [#df62bb9d]
[[CUDAで行列演算：乗算]]でもかなり高速に実行できています...
一番手っ取り早い方法はシェアードメモリを使う方法です．
カーネル関数において，行列Aの列数をMとすると
#code(C){{
for(int k = 0; k < A.width; ++k){
x += A.elements[row*A.width+k]*B.elements[k*B.width+co...
}
}}
の部分だけで，2M回のグローバルメモリ(GPUから見るとグロー...
全体では，スレッド数×2Mで，正方行列を仮定すれば，O(M^3)で...
ブロックごとに部分行列を処理していると考え，ブロックごと...
グローバルメモリへのアクセスをO(M^2)に抑えることができそ...
**シェアードメモリ [#j676529c]
シェアードメモリ(shared memory)はブロック内のスレッドが共...
グローバルメモリに比べて非常に高速です．シェアードメモリ...
スレッド間のバンク・コンフリクトがなければレジスタアクセ...
コンフリクトなしのメモリアクセスについてはCUDA Programmin...
下図は行列の積のイメージです．青の部分が1スレッドがアクセ...
この結果C_ijが計算され結果がグローバルメモリに書き込まれ...
ブロック単位で考えると，黄色で示した部分が1ブロックが処理...
それぞれ元の行列の部分行列と考えることができます．
行列AとBの部分行列(黄色部分)をシェアードメモリに格納すれ...
ただし，1ブロック内のスレッドが1度に処理できるのはBLOCK_S...
行列AとBの部分行列をBLOCK_SIZE×BLOCK_SIZEに分割して処理す...
CENTER:&ref(matrix_mul.jpg);
CENTER:図1 行列の積
**同期処理 [#mb370820]
シェアードメモリを用いるとき，スレッド間の同期が必要にな...
例えば，行列の内容をシェアードメモリに格納して，行列の要...
行列の要素の積を計算する前に，シェアードメモリ内には部分...
各スレッドごとに部分行列の各要素をシェアードメモリへ格納...
ブロック内の全スレッドで格納が完全に終了したことを確かに...
CUDAではブロック内スレッド同期点の設定のために以下の関数...
__syncthreads()
スレッドは，[[CUDAで行列演算：加減算#l7a8f65a]]で述べたワ...
ワープ内のスレッドは暗黙的に同期されます．ワープを意識し...
**実装 [#t6bd24a3]
メモリの確保やデータ転送は前回と同じなので省略します．
行列を格納するためのMatrix構造体ですが，部分行列へのアク...
行列の大きさを示すwidthとheightに加えて，行列要素へのアク...
#code(C){{
struct Matrix
{
int width;
int height;
int stride;
float *elements;
};
}}
部分行列でない場合は，stride = widthとします．
部分行列の場合は，strideは元の行列のwidthとし，elementsに...
先頭アドレスを格納します(図1の行列Cでいうと黄色部分の左上...
元の行列から部分行列を作成する関数は以下です．
#code(C){{
__device__
Matrix GetSubMatrix(Matrix A, int row, int col)
{
Matrix Asub;
Asub.width = BLOCK_SIZE;
Asub.height = BLOCK_SIZE;
Asub.stride = A.stride;
Asub.elements = &A.elements[A.stride*BLOCK_SIZE*row+BLOC...
return Asub;
}
}}
GetSubMatrix関数では第1引数に元の行列，第2,3引数に元の行...
部分行列の位置(0スタート)を渡します．ここで，部分行列は基...
シェアードメモリを用いたカーネルの変更版は以下です．
#code(C){{
__global__
void matrixMulShared(Matrix A, Matrix B, Matrix C)
{
int row = blockIdx.y*blockDim.y+threadIdx.y;
int col = blockIdx.x*blockDim.x+threadIdx.x;
if(row < C.height && col < C.width){
int brow = blockIdx.y;
int bcol = blockIdx.x;
Matrix Csub = GetSubMatrix(C, brow, bcol);
int trow = threadIdx.y;
int tcol = threadIdx.x;
float x = 0.0f;
for(int l = 0; l < (A.width+BLOCK_SIZE-1)/BLOCK_SIZE; +...
Matrix Asub = GetSubMatrix(A, brow, l);
Matrix Bsub = GetSubMatrix(B, l, bcol);
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
// サブ行列の内容をシェアードメモリへ
As[trow][tcol] = Asub.elements[trow*Asub.stride+tcol];
Bs[trow][tcol] = Bsub.elements[trow*Bsub.stride+tcol];
// 他のスレッドがシェアードへの書き込みを終了するのを...
__syncthreads();
for(int k = 0; k < BLOCK_SIZE; ++k){
x += As[trow][k]*Bs[k][tcol];
}
// 次の反復でサブ行列が書き換えられるため，ここで処理...
__syncthreads();
}
Csub.elements[trow*Csub.stride+tcol] = x;
}
}
}}
同期点を2点設定しています(ループ内なので実際に実行される...
シェアードメモリを使用するためには，修飾子 __shared__ を...
**計算時間 [#yab999b2]
[[CUDAで行列演算：乗算]]と同じく，ブロックサイズを16×16と...
計算時間は以下の表です．
| |16|160|1600|
|GPU(シェアードメモリ使用)|0.086|0.15|36|
|GPU|0.085|0.42|398|
|CPU|0.010|12.58|21245|
時間の単位は ms(ミリ秒)です．一回計測しただけです．
シェアードメモリを使うことで高速化ができていることが分か...
終了行:
*CUDAで行列演算 その3 乗算(シェアードメモリ使用版) [#g2d9...
[[CUDAで行列演算：乗算]]で行列の積をやりましたが，メモリ...
CUDAで使えるメモリについて勉強するために，シェアードメモ...
**[[CUDAで行列演算：乗算]]の問題点 [#df62bb9d]
[[CUDAで行列演算：乗算]]でもかなり高速に実行できています...
一番手っ取り早い方法はシェアードメモリを使う方法です．
カーネル関数において，行列Aの列数をMとすると
#code(C){{
for(int k = 0; k < A.width; ++k){
x += A.elements[row*A.width+k]*B.elements[k*B.width+co...
}
}}
の部分だけで，2M回のグローバルメモリ(GPUから見るとグロー...
全体では，スレッド数×2Mで，正方行列を仮定すれば，O(M^3)で...
ブロックごとに部分行列を処理していると考え，ブロックごと...
グローバルメモリへのアクセスをO(M^2)に抑えることができそ...
**シェアードメモリ [#j676529c]
シェアードメモリ(shared memory)はブロック内のスレッドが共...
グローバルメモリに比べて非常に高速です．シェアードメモリ...
スレッド間のバンク・コンフリクトがなければレジスタアクセ...
コンフリクトなしのメモリアクセスについてはCUDA Programmin...
下図は行列の積のイメージです．青の部分が1スレッドがアクセ...
この結果C_ijが計算され結果がグローバルメモリに書き込まれ...
ブロック単位で考えると，黄色で示した部分が1ブロックが処理...
それぞれ元の行列の部分行列と考えることができます．
行列AとBの部分行列(黄色部分)をシェアードメモリに格納すれ...
ただし，1ブロック内のスレッドが1度に処理できるのはBLOCK_S...
行列AとBの部分行列をBLOCK_SIZE×BLOCK_SIZEに分割して処理す...
CENTER:&ref(matrix_mul.jpg);
CENTER:図1 行列の積
**同期処理 [#mb370820]
シェアードメモリを用いるとき，スレッド間の同期が必要にな...
例えば，行列の内容をシェアードメモリに格納して，行列の要...
行列の要素の積を計算する前に，シェアードメモリ内には部分...
各スレッドごとに部分行列の各要素をシェアードメモリへ格納...
ブロック内の全スレッドで格納が完全に終了したことを確かに...
CUDAではブロック内スレッド同期点の設定のために以下の関数...
__syncthreads()
スレッドは，[[CUDAで行列演算：加減算#l7a8f65a]]で述べたワ...
ワープ内のスレッドは暗黙的に同期されます．ワープを意識し...
**実装 [#t6bd24a3]
メモリの確保やデータ転送は前回と同じなので省略します．
行列を格納するためのMatrix構造体ですが，部分行列へのアク...
行列の大きさを示すwidthとheightに加えて，行列要素へのアク...
#code(C){{
struct Matrix
{
int width;
int height;
int stride;
float *elements;
};
}}
部分行列でない場合は，stride = widthとします．
部分行列の場合は，strideは元の行列のwidthとし，elementsに...
先頭アドレスを格納します(図1の行列Cでいうと黄色部分の左上...
元の行列から部分行列を作成する関数は以下です．
#code(C){{
__device__
Matrix GetSubMatrix(Matrix A, int row, int col)
{
Matrix Asub;
Asub.width = BLOCK_SIZE;
Asub.height = BLOCK_SIZE;
Asub.stride = A.stride;
Asub.elements = &A.elements[A.stride*BLOCK_SIZE*row+BLOC...
return Asub;
}
}}
GetSubMatrix関数では第1引数に元の行列，第2,3引数に元の行...
部分行列の位置(0スタート)を渡します．ここで，部分行列は基...
シェアードメモリを用いたカーネルの変更版は以下です．
#code(C){{
__global__
void matrixMulShared(Matrix A, Matrix B, Matrix C)
{
int row = blockIdx.y*blockDim.y+threadIdx.y;
int col = blockIdx.x*blockDim.x+threadIdx.x;
if(row < C.height && col < C.width){
int brow = blockIdx.y;
int bcol = blockIdx.x;
Matrix Csub = GetSubMatrix(C, brow, bcol);
int trow = threadIdx.y;
int tcol = threadIdx.x;
float x = 0.0f;
for(int l = 0; l < (A.width+BLOCK_SIZE-1)/BLOCK_SIZE; +...
Matrix Asub = GetSubMatrix(A, brow, l);
Matrix Bsub = GetSubMatrix(B, l, bcol);
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
// サブ行列の内容をシェアードメモリへ
As[trow][tcol] = Asub.elements[trow*Asub.stride+tcol];
Bs[trow][tcol] = Bsub.elements[trow*Bsub.stride+tcol];
// 他のスレッドがシェアードへの書き込みを終了するのを...
__syncthreads();
for(int k = 0; k < BLOCK_SIZE; ++k){
x += As[trow][k]*Bs[k][tcol];
}
// 次の反復でサブ行列が書き換えられるため，ここで処理...
__syncthreads();
}
Csub.elements[trow*Csub.stride+tcol] = x;
}
}
}}
同期点を2点設定しています(ループ内なので実際に実行される...
シェアードメモリを使用するためには，修飾子 __shared__ を...
**計算時間 [#yab999b2]
[[CUDAで行列演算：乗算]]と同じく，ブロックサイズを16×16と...
計算時間は以下の表です．
| |16|160|1600|
|GPU(シェアードメモリ使用)|0.086|0.15|36|
|GPU|0.085|0.42|398|
|CPU|0.010|12.58|21245|
時間の単位は ms(ミリ秒)です．一回計測しただけです．
シェアードメモリを使うことで高速化ができていることが分か...
ページ名:
![[PukiWiki]](image/pukiwiki.png "[PukiWiki]")
