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

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

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

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

CUDA

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

C++

C++はC言語をもとにしてつくられた最もよく使われるマルチパラダイムプログラミング言語の1つです。オブジェクト指向、ジェネリック、命令型など広く対応しており、多目的に使用されています。

Q&A

2回答

8333閲覧

【CUDA】cudaMemcpyに時間がかかり速度が出ないときの対処法

shukrin

総合スコア14

C

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

CUDA

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

C++

C++はC言語をもとにしてつくられた最もよく使われるマルチパラダイムプログラミング言語の1つです。オブジェクト指向、ジェネリック、命令型など広く対応しており、多目的に使用されています。

1グッド

0クリップ

投稿2020/03/14 11:47

編集2020/03/15 09:11

前提

つい先日からCUDAの勉強を始め、こちらのサイトを参考に配列同士の足し算を行うプログラムを作ってみました。並列計算による高速化を期待していましたが、あまり思ったようにいかず困っております。

発生している問題・エラーメッセージ

512*2048個の要素を持つfloat型配列2つの各要素を足し合わせて別の配列に値を格納するようなプログラムをGPU,CPUでそれぞれ実行し、その計算時間を比較してみたのですがCPUの方が速いという結果になってしまいました。
どこがボトルネックになっているのかを調べるために、GPUによる計算時間をホストからデバイスへのメモリコピー、GPUによる計算、デバイスからホストへのメモリコピーという3つの段階に分けてそれぞれにかかった時間を計測しました。以下がその結果です。

memcpy(HtD): 3314 microsec
gpu_calc: 451 microsec
memcpy(DtH): 783 microsec
cpu_calc: 1999 microsec

計算自体は一応高速化されているようですが、ホストとデバイス間のメモリコピーに時間がかかっていたのが原因でGPUの計算時間が増大していたようです。

該当のソースコード

上記サイトのソースコードを時間計測用に一部改変したものとなっています。

C++

1#include <iostream> 2#include <stdio.h> 3#include <algorithm> 4#include <cuda_runtime.h> 5#include <device_launch_parameters.h> 6#include <chrono> 7 8//N = 512×2048 9int N = 1 << 20; 10 11void vec_sum_c(float k, float c[], float a[], float b[]) { 12 13 for (int i = 0; i < N; i++) { 14 c[i] = k * a[i] + b[i]; 15 } 16} 17 18__global__ 19void vec_sum_k(float k, float* a, float* b, float* c) 20{ 21 int i = blockIdx.x * blockDim.x + threadIdx.x; 22 c[i] = k * a[i] + b[i]; 23} 24 25int main() { 26 27 28 //a,b,cはホスト用、d_a,d_b,d_cはデバイス用のポインタ 29 float* a, * b, * c, * d_a, * d_b, * d_c; 30 31 //ホスト側の配列を用意 32 //a = (float*)malloc(N * sizeof(float)); 33 //b = (float*)malloc(N * sizeof(float)); 34 //c = (float*)malloc(N * sizeof(float)); 35 cudaMallocHost(&a, N * sizeof(float)); 36 cudaMallocHost(&b, N * sizeof(float)); 37 cudaMallocHost(&c, N * sizeof(float)); 38 39 //デバイス側の配列を用意 40 cudaMalloc(&d_a, N * sizeof(float)); 41 cudaMalloc(&d_b, N * sizeof(float)); 42 cudaMalloc(&d_c, N * sizeof(float)); 43 44 //a,bの配列にそれぞれ1,2を代入し、cを初期化 45 for (int i = 0; i < N; i++) { 46 a[i] = 1.0f; 47 b[i] = 2.0f; 48 c[i] = 0.0f; 49 } 50 51 //時間計測開始 52 std::chrono::system_clock::time_point start, memcpyh2d, memcpyd2h, gpus, gpue, cpu; 53 std::cout << "-----TimeStart-----" << std::endl; 54 55 for (int i = 0; i < 10; i++) { 56 57 start = std::chrono::system_clock::now(); 58 59 //ホスト側の配列の内容をデバイス側にコピー 60 cudaMemcpy(d_a, a, N * sizeof(float), cudaMemcpyHostToDevice); 61 cudaMemcpy(d_b, b, N * sizeof(float), cudaMemcpyHostToDevice); 62 cudaMemcpy(d_c, c, N * sizeof(float), cudaMemcpyHostToDevice); 63 64 //ホストの配列をデバイス側にコピーするのにかかった時間 65 memcpyh2d = std::chrono::system_clock::now(); 66 67 //スレッドの設定 68 int blocksize = 256; 69 70 //ブロックあたりのスレッド数(blocksize)を512、 71 //ブロックの総数(gridsize)をN/512用意する 72 //したがって総スレッド数は blocksize × gridsize = N 個 73 dim3 block(blocksize, 1, 1); 74 dim3 grid(N / block.x, 1, 1); 75 76 //GPUの計算時間計測開始 77 gpus = std::chrono::system_clock::now(); 78 79 // カーネル関数の呼び出し 80 vec_sum_k << <grid, block >> > (2.0f, d_a, d_b, d_c); 81 cudaDeviceSynchronize(); 82 83 //GPUの計算時間計測終了 84 gpue = std::chrono::system_clock::now(); 85 86 //計算結果をホストへコピー 87 cudaMemcpy(c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost); 88 89 //計算結果をホストにコピー。GPUの計算終わり。 90 memcpyd2h = std::chrono::system_clock::now(); 91 92 // 計測時間の表示 93 std::cout << "-------GPULoop: " << i << " -------" << std::endl; 94 double HtDtime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(memcpyh2d - start).count()); 95 std::cout << "memcpy(HtD): " << HtDtime << " microsec" << std::endl; 96 double gputime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(gpue - gpus).count()); 97 std::cout << "gpu: " << gputime << " microsec" << std::endl; 98 double DtHtime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(memcpyd2h - gpue).count()); 99 std::cout << "memcpy(DtH): " << DtHtime << " microsec" << std::endl; 100 101 } 102 103 //計算結果をホストにコピー。GPUの計算終わり。CPUの計算開始 104 memcpyd2h = std::chrono::system_clock::now(); 105 106 vec_sum_c(2.0f, c, a, b); 107 108 //CPUの計算時間計測終了 109 cpu = std::chrono::system_clock::now(); 110 111 float maxError = 0.0f; 112 113 //計算結果の確認 114 for (int i = 0; i < N; i++) maxError = std::max(maxError, abs(c[i] - 4.0f)); 115 std::cout << "MaxError: : " << maxError << std::endl; 116 117 //計測時間の表示 118 std::cout << "-------CPU-------" << std::endl; 119 double cputime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(cpu - memcpyd2h).count()); 120 std::cout << "cpu: " << cputime << " microsec" << std::endl; 121 122 //メモリの開放 123 free(a); 124 free(b); 125 free(c); 126 127 cudaFree(d_a); 128 cudaFree(d_b); 129 cudaFree(d_c); 130 131 return 0; 132}

実現したいこと

メモリコピーの時間を含めてCPUより高速で動作するようなCUDAプログラムを作成するためにはどのような方法があるでしょうか?
また計算自体は高速化されているとはいえ、期待したほど(十倍~数百倍以上)ではないのも気になります。上記サイトの結果では少なくとも3000倍の高速化に成功していたようです。私のプログラムがそれほど高速化されていない理由は何でしょうか? また計算の高速化のためにはどのような工夫ができるでしょうか?

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

OS:Windows 10
CUDA version:10.2
GPU:nVidia GeForce GTX1060 6Gb

以下がdevice_queryを実行した結果です。

CUDA Device Query (Runtime API) version (CUDART static linking) Detected 1 CUDA Capable device(s) Device 0: "GeForce GTX 1060 6GB" CUDA Driver Version / Runtime Version 10.2 / 10.2 CUDA Capability Major/Minor version number: 6.1 Total amount of global memory: 6144 MBytes (6442450944 bytes) (10) Multiprocessors, (128) CUDA Cores/MP: 1280 CUDA Cores GPU Max Clock rate: 1709 MHz (1.71 GHz) Memory Clock rate: 4004 Mhz Memory Bus Width: 192-bit L2 Cache Size: 1572864 bytes Maximum Texture Dimension Size (x,y,z) 1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384) Maximum Layered 1D Texture Size, (num) layers 1D=(32768), 2048 layers Maximum Layered 2D Texture Size, (num) layers 2D=(32768, 32768), 2048 layers Total amount of constant memory: 65536 bytes Total amount of shared memory per block: 49152 bytes Total number of registers available per block: 65536 Warp size: 32 Maximum number of threads per multiprocessor: 2048 Maximum number of threads per block: 1024 Max dimension size of a thread block (x,y,z): (1024, 1024, 64) Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535) Maximum memory pitch: 2147483647 bytes Texture alignment: 512 bytes Concurrent copy and kernel execution: Yes with 2 copy engine(s) Run time limit on kernels: Yes Integrated GPU sharing Host Memory: No Support host page-locked memory mapping: Yes Alignment requirement for Surfaces: Yes Device has ECC support: Disabled CUDA Device Driver Mode (TCC or WDDM): WDDM (Windows Display Driver Model) Device supports Unified Addressing (UVA): Yes Device supports Compute Preemption: Yes Supports Cooperative Kernel Launch: No Supports MultiDevice Co-op Kernel Launch: No Device PCI Domain ID / Bus ID / location ID: 0 / 1 / 0 Compute Mode: < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) > deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 10.2, CUDA Runtime Version = 10.2, NumDevs = 1 Result = PASS
s.k👍を押しています

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

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

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

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

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

guest

回答2

0

CPUの計算に使っているvec_sum_c関数ですがループ内のコードが間違ってる気がします。
c[0] = k * a[i] + b[i];
ではなく、
c[i] = k * a[i] + b[i];
では?
この間違い方だと最適化によって全然ちがう処理時間になってしまいますよ?

投稿2020/03/14 16:49

TaroToyotomi

総合スコア1430

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

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

shukrin

2020/03/14 17:47

ご回答ありがとうございます。おっしゃる通りでコードを書き間違えておりました。書き直してから改めて時間を計測してみました。 memcpy(HtD): 2709 microsec gpu: 457 microsec memcpy(DtH): 810 microsec cpu: 2197 microsec これでCPUとGPUで同じ処理を行ったことになると思いますが、結果はこれまでとあまり変わらないようです。
TaroToyotomi

2020/03/14 23:42

じゃあ、最適化がかかっていないのですね。 cudaでのホストとデバイスのメモリ転送ですが、ホスト側のメモリ確保にmallocではなくcudaMallocHost関数を使うといいかもしれません。 これを使うとアロケートと同時にメモリをページロックしてくれるので多少早くなるかもしれません。 あと、HOSTからDEVICEへの入力データの転送での cudaMemcpy(d_c, c, N * sizeof(float), cudaMemcpyHostToDevice); は必要ないと思います。
shukrin

2020/03/15 09:27

ご回答ありがとうございます。mallocをcudaMallocHostに置き換えて時間を測りなおしてみました。epistemeさんに最初のメモリ転送/カーネル実行は遅くなる場合があると教えていただいたので、cudaMallocHostに書き換えたうえで一連のGPU処理を10回繰り返して時間を計測しました。書き換える前の実行時間はepistemeさんへの返信に記載いたしましたので、もしよろしければそちらもご覧ください。 cudaMemcpy(d_c, c, N * sizeof(float), cudaMemcpyHostToDevice)はd_cをcと同じ内容で初期化するために必要なのではないかと考え、残してあります。 -----TimeStart----- -------GPULoop: 0 ------- memcpy(HtD): 1344 microsec gpu: 300 microsec memcpy(DtH): 446 microsec -------GPULoop: 1 ------- memcpy(HtD): 1184 microsec gpu: 221 microsec memcpy(DtH): 395 microsec -------GPULoop: 2 ------- memcpy(HtD): 1219 microsec gpu: 220 microsec memcpy(DtH): 384 microsec -------GPULoop: 3 ------- memcpy(HtD): 1204 microsec gpu: 233 microsec memcpy(DtH): 403 microsec -------GPULoop: 4 ------- memcpy(HtD): 1240 microsec gpu: 274 microsec memcpy(DtH): 439 microsec -------GPULoop: 5 ------- memcpy(HtD): 1182 microsec gpu: 236 microsec memcpy(DtH): 384 microsec -------GPULoop: 6 ------- memcpy(HtD): 1173 microsec gpu: 220 microsec memcpy(DtH): 385 microsec -------GPULoop: 7 ------- memcpy(HtD): 1173 microsec gpu: 220 microsec memcpy(DtH): 385 microsec -------GPULoop: 8 ------- memcpy(HtD): 1206 microsec gpu: 221 microsec memcpy(DtH): 385 microsec -------GPULoop: 9 ------- memcpy(HtD): 1325 microsec gpu: 445 microsec memcpy(DtH): 467 microsec MaxError: : 0 -------CPU------- cpu: 2206 microsec
guest

0

最初のメモリ転送/カーネル実行がかなり遅くなることはよくあります。

ホスト→デバイス / カーネル実行 / デバイス→ホスト
を数回繰り返したとき、2回目以降のメモリ転送にも長い時間かかっていますか?

投稿2020/03/15 02:58

episteme

総合スコア16614

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

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

shukrin

2020/03/15 09:19

ご回答ありがとうございます。GPUによるホスト→デバイス / カーネル実行 / デバイス→ホスト を10回繰り返してそれぞれの時間を計測してみました。 -----TimeStart----- -------GPULoop: 0 ------- memcpy(HtD): 2438 microsec gpu: 450 microsec memcpy(DtH): 902 microsec -------GPULoop: 1 ------- memcpy(HtD): 2643 microsec gpu: 434 microsec memcpy(DtH): 886 microsec -------GPULoop: 2 ------- memcpy(HtD): 2342 microsec gpu: 425 microsec memcpy(DtH): 1176 microsec -------GPULoop: 3 ------- memcpy(HtD): 2184 microsec gpu: 386 microsec memcpy(DtH): 777 microsec -------GPULoop: 4 ------- memcpy(HtD): 2136 microsec gpu: 383 microsec memcpy(DtH): 777 microsec -------GPULoop: 5 ------- memcpy(HtD): 2094 microsec gpu: 382 microsec memcpy(DtH): 775 microsec -------GPULoop: 6 ------- memcpy(HtD): 2428 microsec gpu: 427 microsec memcpy(DtH): 881 microsec -------GPULoop: 7 ------- memcpy(HtD): 2414 microsec gpu: 414 microsec memcpy(DtH): 783 microsec -------GPULoop: 8 ------- memcpy(HtD): 2190 microsec gpu: 381 microsec memcpy(DtH): 785 microsec -------GPULoop: 9 ------- memcpy(HtD): 2386 microsec gpu: 471 microsec memcpy(DtH): 785 microsec MaxError: : 0 -------CPU------- cpu: 2220 microsec 確かに数回繰り返した後の方が最初よりも少し速くなっているようです。TaroToyotomiさんにいただいたアドバイスをもとにmallocをcudaMallocHostに置き換えるとよりはっきりと効果が表れたのでよろしければそちらもご覧ください。
guest

あなたの回答

tips

太字

斜体

打ち消し線

見出し

引用テキストの挿入

コードの挿入

リンクの挿入

リストの挿入

番号リストの挿入

表の挿入

水平線の挿入

プレビュー

まだベストアンサーが選ばれていません

会員登録して回答してみよう

アカウントをお持ちの方は

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

ただいまの回答率
85.45%

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

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

質問する

関連した質問