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

Question

前提

つい先日から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の計算時間が増大していたようです。

該当のソースコード

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

#include <iostream>
#include <stdio.h>
#include <algorithm>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <chrono>

//N = 512×2048
int N = 1 << 20;

void vec_sum_c(float k, float c[], float a[], float b[]) {

	for (int i = 0; i < N; i++) {
		c[i] = k * a[i] + b[i];
	}
}

__global__
void vec_sum_k(float k, float* a, float* b, float* c)
{
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	c[i] = k * a[i] + b[i];
}

int main() {


	//a,b,cはホスト用、d_a,d_b,d_cはデバイス用のポインタ
	float* a, * b, * c, * d_a, * d_b, * d_c;

	//ホスト側の配列を用意
	//a = (float*)malloc(N * sizeof(float));
	//b = (float*)malloc(N * sizeof(float));
	//c = (float*)malloc(N * sizeof(float));
	cudaMallocHost(&a, N * sizeof(float));
	cudaMallocHost(&b, N * sizeof(float));
	cudaMallocHost(&c, N * sizeof(float));

	//デバイス側の配列を用意
	cudaMalloc(&d_a, N * sizeof(float));
	cudaMalloc(&d_b, N * sizeof(float));
	cudaMalloc(&d_c, N * sizeof(float));

	//a,bの配列にそれぞれ1，2を代入し、cを初期化
	for (int i = 0; i < N; i++) {
		a[i] = 1.0f;
		b[i] = 2.0f;
		c[i] = 0.0f;
	}

	//時間計測開始
	std::chrono::system_clock::time_point start, memcpyh2d, memcpyd2h, gpus, gpue, cpu;
	std::cout << "-----TimeStart-----" << std::endl;

	for (int i = 0; i < 10; i++) {

		start = std::chrono::system_clock::now();

		//ホスト側の配列の内容をデバイス側にコピー
		cudaMemcpy(d_a, a, N * sizeof(float), cudaMemcpyHostToDevice);
		cudaMemcpy(d_b, b, N * sizeof(float), cudaMemcpyHostToDevice);
		cudaMemcpy(d_c, c, N * sizeof(float), cudaMemcpyHostToDevice);

		//ホストの配列をデバイス側にコピーするのにかかった時間
		memcpyh2d = std::chrono::system_clock::now();

		//スレッドの設定
		int blocksize = 256;

		//ブロックあたりのスレッド数（blocksize)を512、
		//ブロックの総数（gridsize）をN/512用意する
		//したがって総スレッド数は blocksize × gridsize = N 個
		dim3 block(blocksize, 1, 1);
		dim3 grid(N / block.x, 1, 1);

		//GPUの計算時間計測開始
		gpus = std::chrono::system_clock::now();

		// カーネル関数の呼び出し
		vec_sum_k << <grid, block >> > (2.0f, d_a, d_b, d_c);
		cudaDeviceSynchronize();

		//GPUの計算時間計測終了
		gpue = std::chrono::system_clock::now();

		//計算結果をホストへコピー
		cudaMemcpy(c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);

		//計算結果をホストにコピー。GPUの計算終わり。
		memcpyd2h = std::chrono::system_clock::now();

		// 計測時間の表示
		std::cout << "-------GPULoop: " << i << " -------" << std::endl;
		double HtDtime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(memcpyh2d - start).count());
		std::cout << "memcpy(HtD): " << HtDtime << " microsec" << std::endl;
		double gputime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(gpue - gpus).count());
		std::cout << "gpu: " << gputime << " microsec" << std::endl;
		double DtHtime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(memcpyd2h - gpue).count());
		std::cout << "memcpy(DtH): " << DtHtime << " microsec" << std::endl;

	}

	//計算結果をホストにコピー。GPUの計算終わり。CPUの計算開始
	memcpyd2h = std::chrono::system_clock::now();

	vec_sum_c(2.0f, c, a, b);

	//CPUの計算時間計測終了
	cpu = std::chrono::system_clock::now();

	float maxError = 0.0f;	

	//計算結果の確認
	for (int i = 0; i < N; i++) maxError = std::max(maxError, abs(c[i] - 4.0f));
	std::cout << "MaxError: : " << maxError << std::endl;

	//計測時間の表示
	std::cout << "-------CPU-------" << std::endl;
	double cputime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(cpu - memcpyd2h).count());
	std::cout << "cpu: " << cputime << " microsec" << std::endl;

	//メモリの開放
	free(a);
	free(b);
	free(c);

	cudaFree(d_a);
	cudaFree(d_b);
	cudaFree(d_c);

	return 0;
}

実現したいこと

メモリコピーの時間を含めて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

Answer 1

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

Comments

[shukrin]

ご回答ありがとうございます。おっしゃる通りでコードを書き間違えておりました。書き直してから改めて時間を計測してみました。

memcpy(HtD): 2709 microsec
gpu: 457 microsec
memcpy(DtH): 810 microsec
cpu: 2197 microsec

これでCPUとGPUで同じ処理を行ったことになると思いますが、結果はこれまでとあまり変わらないようです。

[TaroToyotomi]

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

[shukrin]

ご回答ありがとうございます。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

Answer 2

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

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

Comments

[shukrin]

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