質問をすることでしか得られない、回答やアドバイスがある。

15分調べてもわからないことは、質問しよう!

新規登録して質問してみよう
ただいま回答率
85.50%
C

C言語は、1972年にAT&Tベル研究所の、デニス・リッチーが主体となって作成したプログラミング言語です。 B言語の後継言語として開発されたことからC言語と命名。そのため、表記法などはB言語やALGOLに近いとされています。 Cの拡張版であるC++言語とともに、現在世界中でもっとも普及されているプログラミング言語です。

CUDA

CUDAは並列計算プラットフォームであり、Nvidia GPU(Graphics Processing Units)向けのプログラミングモデルです。CUDAは様々なプログラミング言語、ライブラリ、APIを通してNvidiaにインターフェイスを提供します。

Q&A

解決済

2回答

6083閲覧

CUDAでのcudaMemcpyでかかる時間について

退会済みユーザー

退会済みユーザー

総合スコア0

C

C言語は、1972年にAT&Tベル研究所の、デニス・リッチーが主体となって作成したプログラミング言語です。 B言語の後継言語として開発されたことからC言語と命名。そのため、表記法などはB言語やALGOLに近いとされています。 Cの拡張版であるC++言語とともに、現在世界中でもっとも普及されているプログラミング言語です。

CUDA

CUDAは並列計算プラットフォームであり、Nvidia GPU(Graphics Processing Units)向けのプログラミングモデルです。CUDAは様々なプログラミング言語、ライブラリ、APIを通してNvidiaにインターフェイスを提供します。

0グッド

1クリップ

投稿2019/03/07 16:42

編集2019/03/07 17:51

質問内容

cuda初心者です。
現在cudaを使って高速に大型の二次元行列同士の内積を行えるプログラムを作ろうとしております。
そこで下記のようなコードでCUDAではどの部分でどれだけ時間がかかっているのか
確認することにしました。
計算する行列の形状は
| x dot w = out
| ↑ ↑ ↑
| (2000,3000)(3000,5000)(2000,5000) です。

結果は下記の通りです。

output

1cudaMalloc:0.124760秒 2cudaMemcpy(Host_to_Device):0.021225秒 3内積計算:0.000032秒 4cudaMemcpy(Devise_to_Host):14.315247秒

内積計算の速度自体はとても早いです、しかし
デバイス側のメモリからホスト側のメモリへコピーする処理で14秒もかかっています...
この遅さは私の書き方が原因なのでしょうか。
それともCUDAとはこのようなものなのでしょうか?
どうかご教授よろしくお願いします。

あともしよろしければ下記のコードに対するご指摘がございましたら
教えて頂きたいです。

実行したコード

c

1///////////////////////////////////////////////////////////////////////// 2// 3// 備考:このコードでは二次元行列同士の内積を行い実行時間を観測する。 4// 5///////////////////////////////////////////////////////////////////////// 6 7#include <stdio.h> 8#include <stdlib.h> 9#include <time.h> 10 11//////////////////////// 12// ↓行列構造体定義 13//////////////////////// 14typedef struct matrix{ 15 int col; 16 int row; 17 double* element; 18}Matrix; 19 20////////////////////////////////// 21// ↓内積を行うカーネル関数 22////////////////////////////////// 23__global__ 24void gpu_dot(const double* x,const double* w,double* out,const int* x_row,const int* w_row, const int* w_col){ 25 26 int i,j,k; 27 double sum = 0; 28 29 //変数iには現在実行中のスレッドの固有の番号(y軸)が入る(スレッド数は結果保存用行列構造体の行数個ある) 30 //変数jには現在実行中のスレッドの固有の番号(x軸)が入る(スレッド数は結果保存用行列構造体の列数個ある) 31 i = blockIdx.y * blockDim.y + threadIdx.y; 32 j = blockIdx.x * blockDim.x + threadIdx.x; 33 34 for(k=0; k<(*w_col); k++){ 35 sum += x[i*(*x_row)+k] * w[k*(*w_row)+j]; 36 } 37 out[i*(*w_row)+j] = sum; 38} 39 40//////////////////////////////////////////////////// 41// ↓行列構造体を操作する関数のプロトタイプ宣言 42//////////////////////////////////////////////////// 43//↓コンストラクタ 44void Matrix_constructor(Matrix* self,const int col,const int row); 45 46//↓行列の要素に要素番号代入 47void Matrix_init(Matrix* self); 48 49//↓行列をゼロクリア 50void Matrix_zeros(Matrix* self); 51 52//↓行列の中身をすべて表示 53void Matrix_print(Matrix* self); 54 55//行列構造体内の要素開放 56void Matrix_free(Matrix* self); 57 58/////////////////////// 59// ↓メイン関数 60/////////////////////// 61 62int main(){ 63 64 //↓タイマー用変数宣言 65 time_t start,stop; 66 67 //↓行列構造体宣言 68 Matrix x; 69 Matrix w; 70 Matrix out;//←計算結果保存用 71 72 //↓行列構造体のコンストラクタ 73 Matrix_constructor(&x,2000,3000); 74 Matrix_constructor(&w,3000,5000); 75 Matrix_constructor(&out,2000,5000); 76 77 //↓入力用の行列に数値代入 78 Matrix_init(&x); 79 Matrix_init(&w); 80 81 //↓出力保存用行列を0クリア 82 Matrix_zeros(&out); 83 84 //↓カーネル関数の引数に使う変数宣言(配列として使う) 85 double* gpu_x; 86 double* gpu_w; 87 double* gpu_out; 88 89 //↓カーネル関数の引数に使う変数宣言(定数として使う) 90 int* x_row; 91 int* w_row; 92 int* w_col; 93 94 //↓cudaMallocでかかる時間測定開始 95 start = clock(); 96 97 //↓カーネルの引数に使う配列の動的確保 98 cudaMalloc(&gpu_x,sizeof(double)*x.col*x.row); 99 cudaMalloc(&gpu_w,sizeof(double)*w.col*w.row); 100 cudaMalloc(&gpu_out,sizeof(double)*out.col*out.row); 101 102 //↓カーネルの引数に使う定数の動的確保 103 cudaMalloc(&x_row,sizeof(int)); 104 cudaMalloc(&w_row,sizeof(int)); 105 cudaMalloc(&w_col,sizeof(int)); 106 107 //↓cudaMallocでかかる時間測定終了 108 stop = clock(); 109 110 //↓cudaMallocでかかる時間表示 111 printf("cudaMalloc:%lf秒\n",(double)(stop-start)/CLOCKS_PER_SEC); 112 113 //↓cudaMemcpyでかかる時間測定開始 114 start = clock(); 115 116 //↓計算で使う行列の中身をカーネルの引数で使う変数へコピー 117 cudaMemcpy(gpu_x,x.element,sizeof(double)*x.col*x.row,cudaMemcpyHostToDevice); 118 cudaMemcpy(gpu_w,w.element,sizeof(double)*w.col*w.row,cudaMemcpyHostToDevice); 119 120 //↓計算で使う定数の中身をカーネルの引数で使う変数へコピー 121 cudaMemcpy(x_row,&(x.row),sizeof(int),cudaMemcpyHostToDevice); 122 cudaMemcpy(w_row,&(w.row),sizeof(int),cudaMemcpyHostToDevice); 123 cudaMemcpy(w_col,&(w.col),sizeof(int),cudaMemcpyHostToDevice); 124 125 //↓cudaMemcpyでかかる時間測定終了 126 stop = clock(); 127 128 //↓cudaMemcpyでかかる時間表示 129 printf("cudaMemcpy(Host_to_Device):%lf秒\n",(double)(stop-start)/CLOCKS_PER_SEC); 130 131 //↓内積計算でかかる時間測定開始 132 start = clock(); 133 134 //↓内積計算実行 135 gpu_dot<<<dim3(out.row,out.col,1),dim3(1,1,1)>>>(gpu_x,gpu_w,gpu_out,x_row,w_row,w_col); 136 137 //↓内積計算でかかる時間測定終了 138 stop = clock(); 139 140 //↓内積計算でかかる時間表示 141 printf("内積計算:%lf秒\n",(double)(stop-start)/CLOCKS_PER_SEC); 142 143 //↓カーネル用変数からホスト用変数に内容をコピーするのにかかる時間測定開始 144 start = clock(); 145 146 //↓カーネル用変数からホスト用変数に内容をコピー 147 cudaMemcpy(out.element,gpu_out,sizeof(double)*out.col*out.row,cudaMemcpyDeviceToHost); 148 149 //↓カーネル用変数からホスト用変数に内容をコピーするのにかかる時間測定終了 150 stop = clock(); 151 152 //↓カーネル用変数からホスト用変数に内容をコピーするのにかかる時間表示 153 printf("cudaMemcpy(Devise_to_Host):%lf秒\n",(double)(stop-start)/CLOCKS_PER_SEC); 154 155 //↓ホスト側デバイス側共に動的確保した領域開放 156 cudaFree(gpu_x); 157 cudaFree(gpu_w); 158 cudaFree(gpu_out); 159 160 cudaFree(x_row); 161 cudaFree(w_row); 162 cudaFree(w_col); 163 164 Matrix_free(&x); 165 Matrix_free(&w); 166 Matrix_free(&out); 167 168 return 0; 169} 170 171/////////////////////////////////////////////////////////////////////////////////////////////// 172// ↓ここから行列構造体を操作する関数の実装(関数の解説は上記のプロトタイプ宣言部分に記載) 173/////////////////////////////////////////////////////////////////////////////////////////////// 174 175void Matrix_constructor(Matrix* self,const int col,const int row){ 176 self->col = col; 177 self->row = row; 178 179 self->element = (double*)malloc(sizeof(double)*col*row); 180} 181 182void Matrix_init(Matrix* self){ 183 for(int i=0;i<self->col;i++){ 184 for(int j=0;j<self->row;j++){ 185 self->element[i*self->row+j] = i*self->row+j; 186 } 187 } 188} 189 190void Matrix_zeros(Matrix* self){ 191 for(int i=0;i<self->col;i++){ 192 for(int j=0;j<self->row;j++){ 193 self->element[i*self->row+j] = 0; 194 } 195 } 196} 197 198void Matrix_print(Matrix* self){ 199 for(int i=0;i<self->col;i++){ 200 for(int j=0;j<self->row;j++){ 201 printf("[%lf]",self->element[i*self->row+j]); 202 } 203 printf("\n"); 204 } 205} 206 207void Matrix_free(Matrix* self){ 208 free(self->element); 209 self->element = NULL; 210}

補足情報(FW/ツールのバージョンなど)

os:linux mint18.3 cinnamon
cuda version:9.0
GPU:nVidia GeForce GTX1060 3Gb

気になる質問をクリップする

クリップした質問は、後からいつでもMYページで確認できます。

またクリップした質問に回答があった際、通知やメールを受け取ることができます。

バッドをするには、ログインかつ

こちらの条件を満たす必要があります。

takabosoft

2019/03/08 04:13 編集

1年以上前に本を読んだのでほとんど忘れましたが、カーネル関数の呼び出しはたしか非同期だったはずなので、そのコードだと正確に測れないです。mmmisakiさんがおっしゃるようにすべてのカーネル関数が終わるのを待つか、公式のプロファイルツールを素直に使ったほうが良いかと思います。
guest

回答2

0

ベストアンサー

cudaDeviceSynchronize();
をカーネル呼び出しの後につけて実行してみてると正常に計測できるハズ
詳しい説明はcudaDeviceSynchronizeで調べればいい記事がいっぱいでてくるのでそちらの方をご参考にしてみて下さい(^。^)

投稿2019/03/08 01:24

mmmisaki

総合スコア34

バッドをするには、ログインかつ

こちらの条件を満たす必要があります。

退会済みユーザー

退会済みユーザー

2019/03/08 08:36

おっしゃる通りcudaDeviceSynchronize(); を計算直後に挿入して全てのスレッドの計算が終了 したあとに結果をホストに転送したところ 転送に14秒かかっていたのではなく計算に14秒かかっていることが 分かりました。 そして転送にかかった時間はhost→deviceとだいたい同じになったので スレッドの同期が私のコードには抜けていることが分かりました。 今後もっと計算速度を上げれれるようにシェアードメモリにも手を出していこうと 思います。 ご回答ありがとうございました。
退会済みユーザー

退会済みユーザー

2019/03/08 08:44

ちなみに同期処理を挿入後の結果です。 cudaMalloc:0.120117秒 cudaMemcpy(Host_to_Device):0.019329秒 内積計算:14.243599秒 cudaMemcpy(Devise_to_Host):0.011023秒
退会済みユーザー

退会済みユーザー

2019/03/08 09:30 編集

度々失礼します。 GPU計算でこの規模の内積計算が14秒は遅いなと感じ TensorflowとNumpyとCudaで計算速度を競わせてみたところ結果は tensorflow(GPU):0.7秒 cuda(GPU):14秒 Numpy(CPU):1.196869134902954秒 となってしまい、おかしいと感じたので カーネル呼び出し部分を gpu_dot<<<dim3(50,200,1),dim3(100,10,1)>>>(gpu_x,gpu_w,gpu_out,x_row,w_row,w_col); としたところ計算速度が cuda:0.5秒と一瞬で終わるようになりました。 しかし使用するblock/gridとthread/blockを手動で指定するのではなく 入力された行列の形状から自動で指定できるようにしたいのですが何か いい知恵がございましたらご教授願いたいです。 ちなみになぜthread/blockを全て1にしたのかと申しますと thread/blockのx軸数*y軸数*z軸数が1024を超えるとカーネル関数自体が呼び出されなくなってしまうので 1024より上限の広いblock/gridに希望のスレッド数を記載しました。
guest

0

device→host が14秒は異常値です、通常host→deviceと同程度。
最初の転送には初期化モロモロで時間がかかることがままありますから、
deviceに転送/計算/hostに転送 を何度か繰り返して計測してみません?

投稿2019/03/07 19:35

episteme

総合スコア16614

バッドをするには、ログインかつ

こちらの条件を満たす必要があります。

退会済みユーザー

退会済みユーザー

2019/03/08 08:40

毎度ご回答ありがとうございます。 どうやらmmmisakiさんのおっしゃる通り 私の測定方法に問題があったようです。 計算直後にcudaDeviceSynchronize(); を挿入してみたところ転送に14秒かかっていたのではなく 計算に14秒かかっていたことが分かりました。
退会済みユーザー

退会済みユーザー

2019/03/08 09:39 編集

上記のmmmisakiさんへの追加の質問と同じですが失礼します。 GPU計算でこの規模の内積計算が14秒は遅いなと感じ TensorflowとNumpyとCudaで計算速度を競わせてみたところ結果は tensorflow(GPU):0.7秒 cuda(GPU):14秒 Numpy(CPU):1.196869134902954秒 となってしまい、おかしいと感じたので カーネル呼び出し部分を gpu_dot<<<dim3(50,200,1),dim3(100,10,1)>>>(gpu_x,gpu_w,gpu_out,x_row,w_row,w_col); としたところ計算速度が cuda:0.5秒と一瞬で終わるようになりました。 しかし使用するblock/gridとthread/blockを手動で指定するのではなく 入力された行列の形状から自動で指定できるようにしたいのですが何か いい知恵がございましたらご教授願いたいです。 ちなみになぜthread/blockを全て1にしたのかと申しますと thread/blockのx軸数*y軸数*z軸数が1024を超えるとカーネル関数自体が呼び出されなくなってしまうので 1024より上限の広いblock/gridに希望のスレッド数を記載しました。
episteme

2019/03/09 00:24 編集

threads/blockのx*y*zが32の倍数となるようblocks/gridを調整するのがオススメ。32個のCUDA-coreが1-setなので。
退会済みユーザー

退会済みユーザー

2019/03/09 08:30

ご返信ありがとうございます。 貴方の返信を見ながらいろいろ考えてみたところ 素因数分解のアルゴリズムで解決できそうだなと思いました。 まさか10代の頃いやいや学校でやらされていたことが 今更役に立つなんて皮肉なもんです...(汗)
guest

あなたの回答

tips

太字

斜体

打ち消し線

見出し

引用テキストの挿入

コードの挿入

リンクの挿入

リストの挿入

番号リストの挿入

表の挿入

水平線の挿入

プレビュー

15分調べてもわからないことは
teratailで質問しよう!

ただいまの回答率
85.50%

質問をまとめることで
思考を整理して素早く解決

テンプレート機能で
簡単に質問をまとめる

質問する

関連した質問