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

前提

つい先日から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

行動規範の内容に同意します

回答2件

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

投稿2020/03/14 16:49

TaroToyotomi

総合スコア1449

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

行動規範の内容に同意します

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

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

投稿2020/03/15 02:58

episteme

総合スコア16612

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に置き換えるとよりはっきりと効果が表れたのでよろしければそちらもご覧ください。

行動規範の内容に同意します

あなたの回答