LoginSignup
11
7

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?

More than 3 years have passed since last update.

@HnniTns(Tns Hnni)

一番速いfor文はどれか？【中編】

Last updated at Posted at 2021-06-23

最初に

　この記事は前編の続きになっています。ですので、事前コードなど詳細は前編を参照して下さい。

環境

  • エディター：VisualStudio2019 C++17
  • 実行環境：Release x64
  • OS：Windows 10 Home
  • CPU：Intel(R) Core(TM) i7-9750H CPU ＠ 2.60GHz
  • GPU：NVIDIA GeForce RTX 2070

検証方法

　配列に入っている数値を適当な変数に四則演算したうえで時間を計り、検証します。今回はGPUでのマルチスレッド性能を検証します。前回はCPUでのシングル・マルチスレッドを検証しました。

使用技術

　今回をきっかけに初めて触った技術もあるので、その点はご了承ください。それと、参考程度に難易度を書いていますが、実際に導入して使ってみて自分なりの感想＆結果なので、あくまで「参考」程度にとどめておいてください。

GPU-マルチスレッド(GPGPU)

1. C++ AMP(Accelerated Massive Parallelism)

  • Microsoftが開発している標準で付属しているGPU並列処理ライブラリです。CUDAでいう「ブロック・グリッド」といった概念を考えなくても実装が可能なライブラリになっているので、非常に簡単にGPGPUを可能にします。逆にいうと、複雑な事が出来ないという事でもあります。
  • 難易度
    • 導入コスト：★☆☆（インクルードするだけです）
    • 学習コスト：★☆☆（変数を書く際に書き方を覚えないといけませんが、GPU処理部はC++で書くことができます）
    • 資料数　　：★★★（前回の「OpenMP(CPU)」と検索数は同じぐらいです。あくまで、GPUの中での星判定なので注意してください。1）
  • 参考サイト
  • コードへ

2. CUDA

3. OpenCL

4. OpenMP

追記：↓
現時点では、OpenMPオフロードターゲットはWindowsではサポートされていません。.png
う～ん？
現時点では、OpenMPオフロードターゲットはWindowsではサポートされていません。- 集中線.png
以上！現場からでした！（私の数日間返して...）
※README.md参照元
※集中線ダウンロードサイト

5. OpenACC

6. OpenGL(Compute Shader)

  • 皆さんも聞いた・触った事があるはずのOpenGLです7。OpenCLとは違って名前の通りグラフィック機能がメインの機能です。Ver4.3からCompute Shaderが実装されました。歴史は古く、OpenCLの初版が2008年なのに対し、OpenGLの初版は1992年となっています。（Wiki参照）
  • 難易度
    • 導入コスト：★★☆（OpenGL単体は標準でサポートされており、ですが動かそうとすると他のライブラリを導入する事になると思うのでこの評価です。）
    • 学習コスト：★★★（直接関係無いところも含めて、覚える事がOpenCLより多いです。仕方ないですが、本来は描画目的なので仕方ありません。（そう考えれば、DirectX11よ(ry ））
    • 資料数　　：★★☆（OpenMP(GPU)と同じぐらいの量になりますね。やはり、比較的最近出た機能なので検索数が少ないのでしょう。）
  • 参考サイト
  • コードへ

7. SYCL

  • このライブラリは紹介だけします。というのもかなり多くなっているので...8。このライブラリはOpenCLとCUDAをサポートしているフロントエンドライブラリになっています。実は、C++標準規格に選ばれるというのを最終目標にしているようです。
  • 参考サイト

↑ 様々なGPGPUライブラリを紹介・解説しています。

結論

　「合計時間（一処理毎の平均時間）」という形で計算結果をのせます。基本的に単位はミリ秒（ms）としていますが、一部分変更しています。ですので、変更部分に関してはキチンと明記するようにしていますのでご了承ください。それと、当然のことながら、結果に「OpenMP」と「OpenACC」は記入できないので、その点は注意して下さい。

メモリ管理も含めた場合（確保・転送）

　この欄はメモリ管理を含めて時間を計測しています。GPGPUに置ける大体のオーバーヘッドの部分を含むやり方となっております。

メモリ管理（確保・転送）.png

要素数 / 種類 C++ AMP CUDA OpenCL OpenGL
100万 591.17 (5.91) 190.99 (1.91) 79.37 (0.79) 256.88 (2.57)
500万 770.42 (7.70) 466.74 (4.67) 304.60 (3.05) 1357.49 (13.57)
1000万 992.72 (9.93) 707.95 (7.08) 567.68 (5.68) 2712.21 (27.12)
5000万 2871.96 (28.72) 2870.70 (28.71) 2740.09 (27.40) 14217.95 (142.18)
1億 6972.03 (69.72) 5602.62 (56.03) 5454.64 (54.55) 26895.08 (268.95)

　上記の結果になり、とてつもなく遅い結果になってしまいました。しかも、OpenGLに関してはかなり遅いことになっています。ですが、実際は、メモリ確保に関しては初期化で行うことは多いと思うので、上記の計測結果よりは早くなるはずです。しかも、GPU側で処理させているコードは実質無いようなものなので、その点も注意すべきでしょう。

メモリ管理も含めた場合（転送のみ）

　この欄はメモリ管理を含めて時間を計測していますが、GPU側への転送のみとなっておりますので、現実的な使い方になるとは思います。

メモリ管理（転送）.png

要素数 / 種類 C++ AMP CUDA OpenCL OpenGL
100万 9.61 (0.10) 44.95 (0.45) 40.59 (0.41) 91.98 (0.92)
500万 19.84 (0.20) 209.38 (2.09) 214.78 (2.15) 260.47 (2.60)
1000万 22.47 (0.22) 448.35 (4.48) 437.63 (4.38) 489.05 (4.89)
5000万 45.68 (0.46) 2258.27 (22.58) 2225.22 (22.25) 2606.78 (26.07)
1億 77.40 (0.77) 4485.25 (44.85) 4236.12 (42.36) 5347.58 (53.48)

　大体予想した感じになりましたが、C++AMPが他と比べて圧倒的に早くなっています。他に、OpenGLに関してはメモリ確保でかなり時間を取られていたことが分かりますね。

処理部分のみ（確保・転送は除く）

　この欄はメモリ管理を含みません、なので処理部分（実行部分）のみループさせています。ですので、かなり高速に動くため、ここだけ単位を「ns（ナノ秒）」に変更します。というのも、グラフにする時にとてつもなく見にくくなるのでここだけ別の単位を用います。

処理部分のみ.png

要素数 / 種類 C++ AMP CUDA OpenCL OpenGL
100万 1235.50 (12.36) 3959.50 (39.60) 1361.90 (13.62) 8.00 (0.08)
500万 1255.80 (12.56) 114.50 (1.15) 235.30 (2.35) 6.40 (0.06)
1000万 1464.30 (14.64) 115.50 (1.16) 240.30 (2.40) 8.00 (0.08)
5000万 1290.90 (12.91) 113.00 (1.13) 240.30 (2.40) 6.60 (0.07)
1億 1356.20 (13.56) 106.10 (1.06) 1490.10 (14.90) 6.80 (0.07)

　もう一度、確認の為に書いておきます。「ここだけ単位を「ns（ナノ秒）」に変更します」
　上記の結果になりましたが、自分としてはかなり驚く結果になっています。C++ AMPとOpenGLはまだ分かる。なぜ、CUDAは一番負荷が少ない100万がとびぬけて一番時間がかかっている？なぜ、OpenCLはサンドイッチ状態でややとびぬけて時間がかかっている？なぜだか訳が分かりません。誰か詳しい人、教えて下さい。m(_ _)m
　更に、OpenGLに関しては急激に速度が速くなっています。グラフでは完全に床と化しています（笑）
　といっても、結果を見る限り、要素数が増えれば増える程に経過時間がかかるわけではなく、要素数と処理時間はあまり一致していないと思われます。ですが、GPUに投げている計算（現状は皆無）にもっと重たい負荷を書ければ結果は大幅に変わってくると思います。

CPU・GPUの平均値

CPU・GPU.png
　このグラフは各GPU・CPU処理方法をまとめた上で平均値にしてグラフ化したものになっています。こう見ると、「処理部分のみ」の軽さが際立っています。やはり、GPGPUにおけるメモリ確保・転送などの「オーバーヘッド」はかなり大きのが分かります。
　そして、CPUとGPUを比べてみると、CPUとGPUにおいて中身が無いと、一部を除いてほとんど経過時間は変わらないということも分かりました。

各項目の順位

速度.png
導入コスト.png
学習コスト.png
資料数.png
　速度に関しては、「メモリ管理も含めた場合（転送のみ）」を参照しています。そして、他の項目に関しては、私の体感・経験を基にしているデータなので、皆さんは参考程度にして下さい。

コード

main.cpp
// 予めブロック数を手動で指定しておく
static const std::unordered_map<uint32_t, std::array<uint32_t, 3>> block_nums
	{ { 1000000, { { 10,10,10 } } }, { 5000000, { { 25,20,10 } } }, { 10000000, { { 25,20,20 } } },
	{ 50000000, { { 40,25,25 } } }, { 100000000, { { 50,50,40 } } } };

const auto& block_num{ block_nums.at(ArrSize) };

C++ AMP

main.cpp
for (size_t i = 0; i < LoopNum; i++)
{
	timer.Start(); // 開始
	{
		// デバイス側のバッファを確保し、データを転送
		Concurrency::array<int, 1> arr(ArrSize, num_array.begin(), num_array.end());

		parallel_for_each(arr.extent, [&arr](index<1> idx) restrict(amp)
			{
				if (idx[0] > ArrSize)	return;

				arr[idx] += ConstNumber;
				arr[idx] -= ConstNumber;
				arr[idx] *= ConstNumber;
				arr[idx] /= ConstNumber;
				arr[idx] += 10;
			});

		// ホスト側にデータを転送
		//num_array = arr;
	}
	timer.End(); // 終了

	times[i] = GetTime(timer);
}


cout << "C++ AMP(GPU : Multithread) " << endl;

OutPutResult(times);

CUDA

長くなったので折りたたんでいます
main.cpp
int* d_mem{};

const size_t ArrayMemSize{ ArrSize * sizeof(int) };

// グリッド数はXYZともに10で固定
const dim3 grid{ 10, 10, 10 }, block{ block_num[0], block_num[1], block_num[2] };

// ブロック数とグリッド数の表示
cout << endl << "block.x: " << block.x << ", y: " << block.y << ", z: " << block.z << endl;
cout << "grid.x : " << grid.x << ", y: " << grid.y << ", z: " << grid.z << endl << endl;

for (size_t i = 0; i < LoopNum; i++)
{
	timer.Start(); // 開始

	// デバイス側のメモリを確保
	CUDA_CHECK(cudaMalloc(&d_mem, ArrayMemSize));

	// 初期化
	CUDA_CHECK(cudaMemset(d_mem, 0, ArrayMemSize));

	// ホスト側からデバイス側にコピー
	CUDA_CHECK(cudaMemcpy(d_mem, num_array.data(), ArrayMemSize, cudaMemcpyHostToDevice));

	// 処理開始
	CudaCalc(grid, block, d_mem, ConstNumber, ArrSize);

	// 処理を完了するまでCPUを待機させる
	CUDA_CHECK(cudaDeviceSynchronize());

	// デバイス側からホスト側にコピー
	//CUDA_CHECK(cudaMemcpy(d_mem, num_array.data(), ArrayMemSize, cudaMemcpyDeviceToHost));

	// デバイス側のメモリを解放
	CUDA_CHECK(cudaFree(d_mem));

	timer.End(); // 終了

	times[i] = timer.GetMicroTimer();
}

// CUDAで使用したリソースのリセット
CUDA_CHECK(cudaDeviceReset());

cout << "CUDA(GPU: Multithread) " << endl;

OutPutResult(times, "CUDA");
cudatest.cu
#include <cuda_runtime.h>

#define CUDA_DEBUG false

#if CUDA_DEBUG
#include <cstdio>
#endif

#include "cuda_test.cuh"

__global__ void CudaCalcTest(int* num_arr, const int const_number, const int max_count)
{
	int index_x = blockIdx.x * blockDim.x + threadIdx.x;
	int index_y = blockIdx.y * blockDim.y + threadIdx.y;
	int index_z = blockIdx.z * blockDim.z + threadIdx.z;

	int global_size_x = blockDim.x * gridDim.x;
	int global_size_y = blockDim.y * gridDim.y;

	int index = (index_z * global_size_x * global_size_y) + (index_y * global_size_x) + index_x;

#if CUDA_DEBUG
	printf("threadIdx:(%d, %d, %d) blockIdx:(%d, %d, %d) blockDim:(%d, %d, %d) gridDim:(%d, %d, %d)\n",
		threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y, blockIdx.z, blockDim.x, blockDim.y,
		blockDim.z, gridDim.x, gridDim.y, gridDim.z);
	printf("index: %d\n", index);
#endif

	if (index >= max_count)	return;

	num_arr[index] += const_number;
	num_arr[index] -= const_number;
	num_arr[index] *= const_number;
	num_arr[index] /= const_number;
	num_arr[index] += 10;
}

// 処理
void CudaCalc(const dim3& grid, const dim3& block, int* d_mem, const int const_number, const int max_count)
{
	CudaCalcTest << <grid, block >> > (d_mem, const_number, max_count);
}
cudatest.cuh
void CudaCalc(const dim3& grid, const dim3& block, int* d_mem, const int const_number, const int max_count);
cudacheck.c
#define USE_CUDA_CHECK false

#if USE_CUDA_CHECK
#define CUDA_CHECK(call)\
{\
	const cudaError_t error = call;\
	if (error != cudaSuccess)\
	{\
		printf("Error: %s:%d, ", __FILE__, __LINE__);\
		printf("code:%d, reason: %s\n", error, cudaGetErrorString(error));\
		system("pause");\
		exit(1);\
	}\
}
#else
#define CUDA_CHECK(call) (call)
#endif

OpenCL

長くなったので折りたたんでいます
main.cpp
cl_device_id device_id{};
cl_context context{};
cl_command_queue command_queue{};
cl_mem memobj{};
cl_program program{};
cl_kernel kernel{};
cl_int ret{};

// デバイスを使う準備
{
	cl_platform_id platform_id{};
	cl_uint ret_num_devices{}, ret_num_platforms{};

	// プラットフォーム・デバイスの情報の取得
	CL_CHECK(clGetPlatformIDs(1, &platform_id, &ret_num_platforms));
	cout << "platforms: " << ret_num_platforms << endl;

	CL_CHECK(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices));
	cout << "devices: " << ret_num_devices << endl;

	// OpenCLコンテキストの作成
	context = clCreateContext(nullptr, 1, &device_id, nullptr, nullptr, &ret);
	CL_CHECK(ret);

	// コマンドキューの作成
	command_queue = clCreateCommandQueue(context, device_id, 0, &ret);
	CL_CHECK(ret);
}

// プログラムの準備
{
	size_t source_size{};
	const char* source_str{};
	std::string source_buffer;

	// ファイルからソースコードを読み込む
	{
		FILE* fp{};
		const std::string fileName{ "./opencl_test.cl" };

		// カーネルを含むソースコードをロード
		//（fstreamで読み込まないと「source file is not valid UTF-8」でビルドエラーになる模様）
		std::fstream kernelFile(fileName);

		if (!kernelFile.is_open())
		{
			cout << "Failed to load kernel." << endl;
			system("pause");
			exit(1);
		}

		source_buffer = std::string
		{ std::istreambuf_iterator<char>(kernelFile), std::istreambuf_iterator<char>() };

		source_str = source_buffer.c_str();
	}

	// 読み込んだソースからカーネルプログラムを作成
	program = clCreateProgramWithSource(
		context, 1, (const char**) &source_str, (const size_t*) &source_size, &ret);
	CL_CHECK(ret);

	// カーネルプログラムをビルド
	ret = clBuildProgram(program, 1, &device_id, nullptr, nullptr, nullptr);
	if (ret != CL_SUCCESS)
	{
		auto before_ret = ret;
		std::array<char, 2048> buffer{};
		size_t len{};

		CL_CHECK(clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG,
			sizeof(buffer), buffer.data(), &len));

		cout << "error source: " << source_buffer << endl << endl;

		cout << "ProgramBuildInfo: " << buffer.data() << endl;
		system("pause");
		exit(1);
	}

	// OpenCLカーネルの作成
	kernel = clCreateKernel(program, "vecAdd", &ret);
}

for (size_t i = 0; i < LoopNum; i++)
{
	timer.Start(); // 開始

	// デバイス側のメモリバッファの作成
	memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, ArrSize * sizeof(float), nullptr, &ret);
	CL_CHECK(ret);

	// デバイスの使用
	{
		// メモリバッファにデータを転送
		CL_CHECK(clEnqueueWriteBuffer(command_queue, memobj, CL_TRUE, 0, ArrSize * sizeof(float),
			num_array.data(), 0, nullptr, nullptr));

		// OpenCLカーネル引数の設定
		CL_CHECK(clSetKernelArg(kernel, 0, sizeof(cl_mem), &memobj));
		// 普通の変数を渡す時だけバッファ転送は不要
		CL_CHECK(clSetKernelArg(kernel, 1, sizeof(int), &(ArrSize)));
		CL_CHECK(clSetKernelArg(kernel, 2, sizeof(int), &(ConstNumber)));

		constexpr int dim{ 3 };
		const std::array<size_t, dim>
			global_work_size{ block_num[0], block_num[1], block_num[2] },
			local_work_size{ 10, 10, 10 };

		// OpenCLカーネルを実行
		CL_CHECK(clEnqueueNDRangeKernel(command_queue, kernel, dim, nullptr,
			global_work_size.data(), local_work_size.data(), 0, nullptr, nullptr));

		// 処理を完了するまでCPUを待機させる
		CL_CHECK(clFinish(command_queue));

		// デバイス側からホスト側にデータを転送
		//CL_CHECK(clEnqueueReadBuffer(command_queue, memobj, CL_TRUE, 0, ArrSize * sizeof(float),
		//	num_array.data(), 0, nullptr, nullptr));
	}

	// デバイス側のメモリバッファを削除
	CL_CHECK(clReleaseMemObject(memobj));

	timer.End(); // 終了

	times[i] = GetMicroTimer(timer);
}

// 終了処理
{
	CL_CHECK(clReleaseKernel(kernel));
	CL_CHECK(clReleaseProgram(program));
	CL_CHECK(clReleaseCommandQueue(command_queue));
	CL_CHECK(clReleaseContext(context));
}

cout << "OpenCL(GPU: Multithread) " << endl;

OutPutResult(times, "OpenCL");
opencl_test.cl
#define OPENCL_DEBUG false

__kernel void vecAdd(__global float* num_arr, const int const_number, const int max_size)
{
	size_t global_id_0 = get_global_id(0);
	size_t global_id_1 = get_global_id(1);
	size_t global_id_2 = get_global_id(2);

	size_t global_size_0 = get_global_size(0);
	size_t global_size_1 = get_global_size(1);

	int index =
		(global_id_2 * global_size_0 * global_size_1) + (global_id_1 * global_size_0) + global_id_0;

#if OPENCL_DEBUG
	size_t local_id_0 = get_local_id(0);
	size_t local_id_1 = get_local_id(1);
	size_t local_id_2 = get_local_id(2);

	size_t offset_0 = get_global_offset(0);
	size_t offset_1 = get_global_offset(1);
	size_t offset_2 = get_global_offset(2);

	size_t global_size_2 = get_global_size(2);

	printf("global_id_0: %d, global_id_1: %d, global_id_2: %d\n", index_0, index_1, index_2);
	printf("global_size_0: %d, global_size_1: %d, global_size_2: %d\n",
		global_size_0, global_size_1, global_size_2);
	printf("local_id_0: %d, local_id_1: %d, local_id_2: %d\n", local_id_0, local_id_1, local_id_2);
	printf("index: %d\n", index);
#endif

	if (index >= max_size)	return;

	num_arr[index] += const_number;
	num_arr[index] -= const_number;
	num_arr[index] *= const_number;
	num_arr[index] /= const_number;
	num_arr[index] += 10;
}
cLErrorCheck.cl
// CUDAみたいに標準関数でエラーメッセージを取得して欲しい（願望）
char* CLGetErrorString(const cl_int error)
{
	static char* err_msg[]
	{
		"CL_SUCCESS",
		"CL_DEVICE_NOT_FOUND",
		"CL_DEVICE_NOT_AVAILABLE",
		"CL_COMPILER_NOT_AVAILABLE",
		"CL_MEM_OBJECT_ALLOCATION_FAILURE",
		"CL_OUT_OF_RESOURCES",
		"CL_OUT_OF_HOST_MEMORY",
		"CL_PROFILING_INFO_NOT_AVAILABLE",
		"CL_MEM_COPY_OVERLAP",
		"CL_IMAGE_FORMAT_MISMATCH",
		"CL_IMAGE_FORMAT_NOT_SUPPORTED",
		"CL_BUILD_PROGRAM_FAILURE",
		"CL_MAP_FAILURE",
		"CL_MISALIGNED_SUB_BUFFER_OFFSET",
		"CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST",
		"CL_COMPILE_PROGRAM_FAILURE     ",
		"CL_LINKER_NOT_AVAILABLE",
		"CL_LINK_PROGRAM_FAILURE",
		"CL_DEVICE_PARTITION_FAILED",
		"CL_KERNEL_ARG_INFO_NOT_AVAILABLE",
		"",
		"",
		"",
		"",
		"",
		"",
		"",
		"",
		"",
		"",
		"CL_INVALID_VALUE",
		"CL_INVALID_DEVICE_TYPE",
		"CL_INVALID_PLATFORM",
		"CL_INVALID_DEVICE",
		"CL_INVALID_CONTEXT",
		"CL_INVALID_QUEUE_PROPERTIES",
		"CL_INVALID_COMMAND_QUEUE",
		"CL_INVALID_HOST_PTR",
		"CL_INVALID_MEM_OBJECT",
		"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
		"CL_INVALID_IMAGE_SIZE",
		"CL_INVALID_SAMPLER",
		"CL_INVALID_BINARY",
		"CL_INVALID_BUILD_OPTIONS",
		"CL_INVALID_PROGRAM",
		"CL_INVALID_PROGRAM_EXECUTABLE",
		"CL_INVALID_KERNEL_NAME",
		"CL_INVALID_KERNEL_DEFINITION",
		"CL_INVALID_KERNEL",
		"CL_INVALID_ARG_INDEX",
		"CL_INVALID_ARG_VALUE",
		"CL_INVALID_ARG_SIZE",
		"CL_INVALID_KERNEL_ARGS",
		"CL_INVALID_WORK_DIMENSION",
		"CL_INVALID_WORK_GROUP_SIZE",
		"CL_INVALID_WORK_ITEM_SIZE",
		"CL_INVALID_GLOBAL_OFFSET",
		"CL_INVALID_EVENT_WAIT_LIST",
		"CL_INVALID_EVENT",
		"CL_INVALID_OPERATION",
		"CL_INVALID_GL_OBJECT",
		"CL_INVALID_BUFFER_SIZE",
		"CL_INVALID_MIP_LEVEL",
		"CL_INVALID_GLOBAL_WORK_SIZE",
		"CL_INVALID_PROPERTY",
		"CL_INVALID_IMAGE_DESCRIPTOR",
		"CL_INVALID_COMPILER_OPTIONS",
		"CL_INVALID_LINKER_OPTIONS",
		"CL_INVALID_DEVICE_PARTITION_COUNT", };

	return err_msg[error];
}

#define USE_CL_CHECK true

#if USE_CL_CHECK
#define CL_CHECK(call)\
{\
	const cl_int error = call * -1;\
	if (error != CL_SUCCESS)\
	{\
		printf("Error: %s:%d, ", __FILE__, __LINE__);\
		printf("code:%d, reason:%s\n", error, CLGetErrorString(error));\
		system("pause");\
		exit(1);\
	}\
}
#else
#define CL_CHECK(call) (call)
#endif

OpenMP(GPU)

少し長くなったので折りたたんでいます
恐らく、下記のようなコードになるとは思います。ですが、**コンパイルと実行が出来ていないので正しいコードとは言えませんし、実際にGPUで動いているのを確認も出来ていない**のでその点だけ注意して下さい。
main.cpp
int* arr{ num_array.data() };

for (size_t i = 0; i < LoopNum; i++)
{
	timer.Start();

#pragma omp target data map(tofrom: arr[0:ArrSize])
#pragma omp target teams distribute parallel for
	for (int j = 0; j < ArrSize; j++)
	{
		if (j >= ArrSize)	continue;

		arr[j] += ConstNumber;
		arr[j] -= ConstNumber;
		arr[j] *= ConstNumber;
		arr[j] /= ConstNumber;
		arr[j] += 10;
	}

	timer.End();

	times[i] = timer.GetMicroTimer();
}

cout << "OpenMP(GPU: Multithread) " << endl;

OutPutResult(times, "OpenMP");

流石に動くであろうコードをのせておいた方がいいので、下記に記載しておきます。
#pragma omp distribute parallel for（参照サイト）

sample.cpp
const int N = 8;
int A[N], B[N], C[N];
int k = 4;
int nteams = 16;
int block_threads = N/nteams;
for(int i=0; i<N; ++i) 
{
   A[i] = 0; 
   B[i] = i;
   C[i] = 3*i;
}
#pragma omp target map(tofrom: A) map(to: B, C)
#pragma omp teams num_teams(nteams)
#pragma omp distribute parallel for dist_schedule(static, block_threads)
for(int i=0; i<N; ++i) 
{
  A[i] = B[i] + k*C[i];
}

OpenACC

長くなったので折りたたんでいます
上記のOpenMPと同じく「恐らく」コードになってしまいますが、例を記述します。ですので、**コンパイルと実行が出来ていないので正しいコードとは言えませんし、実際にGPUで動いているのを確認も出来ていない**のでその点だけ注意して下さい。
main.cpp

int* arr{ num_array.data() };

for (size_t i = 0; i < LoopNum; i++)
{
	timer.Start();

#pragma acc kernels loop independent
	for (int j = 0; j < ArrSize; j++)
	{
		if (j >= ArrSize)	continue;

		num_array[j] += ConstNumber;
		num_array[j] -= ConstNumber;
		num_array[j] *= ConstNumber;
		num_array[j] /= ConstNumber;
		num_array[j] += 10;
	}

	timer.End();

	times[i] = timer.GetMicroTimer();
}

cout << "OpenMP(GPU: Multithread) " << endl;

OutPutResult(times, "OpenMP");

こちらも上記のOpenMPと同じく、動くであろうコードをのせておいた方がいいので、下記に記載しておきます。
4章　OpenACC を使って、まず始めてみよう（参照サイト）

sample.cpp
/*
 *  Copyright 2012 NVIDIA Corporation
 *
 *  Licensed under the Apache License, Version 2.0 (the "License");
 *  you may not use this file except in compliance with the License.
 *  You may obtain a copy of the License at
 *
 *      http://www.apache.org/licenses/LICENSE-2.0
 *
 *  Unless required by applicable law or agreed to in writing, software
 *  distributed under the License is distributed on an "AS IS" BASIS,
 *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
 *  See the License for the specific language governing permissions and
 *  limitations under the License.
 */
#define NN 4096
#define NM 4096

double A[NN][NM];
double Anew[NN][NM];

const int n = NN;
const int m = NM;
const int iter_max = 1000;

const double tol = 1.0e-6;
double error = 1.0;

memset(A, 0, n* m * sizeof(double));
memset(Anew, 0, n* m * sizeof(double));

for (int j = 0; j < n; j++)
{
	A[j][0] = 1.0;
	Anew[j][0] = 1.0;
}

printf("Jacobi relaxation Calculation: %d x %d mesh\n", n, m);

int iter = 0;

#pragma acc data copy(A), create(Anew)
// while ループの前に A[] とAnew[] 配列のアロケート＆コピーを行う
while (error > tol && iter < iter_max)
{
	error = 0.0;

#pragma acc kernels
	for (int j = 1; j < n - 1; j++)
	{
		for (int i = 1; i < m - 1; i++)
		{
			Anew[j][i] = 0.25 * (A[j][i + 1] + A[j][i - 1] + A[j - 1][i] + A[j + 1][i]);
			error = fmax(error, fabs(Anew[j][i] - A[j][i]));
		}
	}
#pragma acc kernels
	for (int j = 1; j < n - 1; j++)
	{
		for (int i = 1; i < m - 1; i++)
		{
			A[j][i] = Anew[j][i];
		}
	}

	if (iter % 100 == 0) printf("%5d, %0.6f\n", iter, error);

	iter++;
}

exit(0);

OpenGL(Compute Shader)

長くなったので折りたたんでいます
main.cpp
auto InitOpenGL{ []()
{
	auto inits_glfw = glfwInit();
	if (inits_glfw != GLFW_TRUE)
	{
		throw std::runtime_error("error occurred: glfwInit!");
	}
	glfwWindowHint(GLFW_CONTEXT_VERSION_MAJOR, 4);
	glfwWindowHint(GLFW_CONTEXT_VERSION_MINOR, 3);
	glfwWindowHint(GLFW_OPENGL_FORWARD_COMPAT, GLFW_TRUE);
	glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE);
	glfwWindowHint(GLFW_VISIBLE, GLFW_FALSE);

	GLFWwindow* window = glfwCreateWindow(1, 1, "invisible window", nullptr, nullptr);
	if (window == nullptr)
	{
		throw std::runtime_error("error occurred: glfwCreateWindow!");
	}
	glfwMakeContextCurrent(window);

	auto inits_glew = glewInit();
	if (inits_glew != GLEW_OK)
	{
		throw std::runtime_error("error occurred: glewInit!");
	}
} };

auto CreateComputeShaderProgram{ [](const char* shader_src)
{
	GLuint shader = glCreateShader(GL_COMPUTE_SHADER);
	glShaderSource(shader, 1, &shader_src, nullptr);
	glCompileShader(shader);

	GLint compiles = 0;
	glGetShaderiv(shader, GL_COMPILE_STATUS, &compiles);
	if (compiles == GL_FALSE)
	{
		GLint log_length = 0;
		glGetShaderiv(shader, GL_INFO_LOG_LENGTH, &log_length);
		std::vector<GLchar> info_log(log_length);
		glGetShaderInfoLog(shader, log_length, &log_length, info_log.data());

		glDeleteShader(shader);

		std::string error_msg = "error occurred in compiling shader: ";
		throw std::runtime_error(error_msg + info_log.data());
	}

	GLuint program = glCreateProgram();

	glAttachShader(program, shader);
	glLinkProgram(program);

	GLint links = 0;
	glGetProgramiv(program, GL_LINK_STATUS, &links);
	if (links == GL_FALSE)
	{
		GLint log_length = 0;
		glGetProgramiv(program, GL_INFO_LOG_LENGTH, &log_length);
		std::vector<GLchar> info_log(log_length);
		glGetProgramInfoLog(program, log_length, &log_length, info_log.data());

		glDeleteProgram(program);
		glDeleteShader(shader);

		std::string error_msg = "error occurred in linking shader: ";
		throw std::runtime_error(error_msg + info_log.data());
	}

	glDetachShader(program, shader);
	glDeleteShader(shader);

	return program;
} };

GLuint shader_program{};

try
{
	InitOpenGL();

	{
		GLint x{}, y{}, z{};

		glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 0, &x);
		glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 1, &y);
		glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_COUNT, 2, &z);

		cout << "num_groups_x: " << x << ", y: " << y << ", z: " << z << endl;

		glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_SIZE, 0, &x);
		glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_SIZE, 1, &y);
		glGetIntegeri_v(GL_MAX_COMPUTE_WORK_GROUP_SIZE, 2, &z);

		cout << "size_groups_x: " << x << ", y: " << y << ", z: " << z << endl;

		cout << "MAX_COMPUTE_WORK_GROUP_INVOCATIONS: " << GL_MAX_COMPUTE_WORK_GROUP_INVOCATIONS << endl << endl;
	}

	{
		const std::string fileName{ "./OpenglTest.glsl" };

		// カーネルを含むソースコードをロード
		std::fstream kernelFile(fileName);

		if (!kernelFile.is_open())
		{
			cout << "Failed to load" << "glsl: " << fileName << endl;
			system("pause");
			exit(1);
		}

		std::string source_buffer = std::string
		{ std::istreambuf_iterator<char>(kernelFile), std::istreambuf_iterator<char>() };

		shader_program = CreateComputeShaderProgram(source_buffer.data());
	}
}
catch (const std::exception& e)
{
	std::cerr << e.what() << std::endl;
	system("pause");
	exit(1);
}

// メモリーバッファの作成
GLuint ssbo{};
glGenBuffers(1, &ssbo);

const GLint uniform_max_count{ glGetUniformLocation(shader_program, "max_count") },
	uniform_const_number{ glGetUniformLocation(shader_program, "const_number") };

for (size_t i = 0; i < LoopNum; i++)
{
	timer.Start(); // 開始

	// デバイス側へデータを転送（配列の更新）
	glBindBuffer(GL_SHADER_STORAGE_BUFFER, ssbo);
	glBufferData(GL_SHADER_STORAGE_BUFFER, num_array.size() * sizeof(int), num_array.data(),
		GL_DYNAMIC_COPY);
	glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, ssbo);
	glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0); // 解放する

	// シェーダプログラムの使用開始
	glUseProgram(shader_program);

	// デバイス側へデータを転送（一部を更新）
	glUniform1ui(uniform_max_count, ArrSize);
	glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, ssbo);
	glUniform1ui(uniform_const_number, ConstNumber);
	glBindBufferBase(GL_SHADER_STORAGE_BUFFER, 3, ssbo);

	// OpenCLカーネルを実行
	glDispatchCompute(block_num[0], block_num[1], block_num[2]);

	// 処理を完了するまでCPUを待機させる
	glMemoryBarrier(GL_SHADER_STORAGE_BARRIER_BIT);

	// ホスト側にデータを転送（配列情報取得）
	//glBindBuffer(GL_SHADER_STORAGE_BUFFER, ssbo);
	//glGetBufferSubData(GL_SHADER_STORAGE_BUFFER, 0, num * sizeof(int), num_array.data());
	//glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0);

	// 終了
	glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0);

	// シェーダープログラムの使用終了
	glUseProgram(0);

	timer.End(); // 終了

	times[i] = GetTime(timer);
}

glBindBuffer(GL_SHADER_STORAGE_BUFFER, 0);
glDeleteBuffers(1, &ssbo);

glDeleteProgram(shader_program);

glfwTerminate();

cout << endl << "＜＜＜ OpenGL(GPU: Multithread)  ＞＞＞" << endl << endl;

OutPutResult(times, "OpenGL");
opengltest.glsl
#version 430

#define GL_DEBUG true

uniform uint max_count;
uniform int const_number;

layout(std430, binding = 3) buffer layout_dest
{
	int num_arr[];
};

layout(local_size_x = 10, local_size_y = 10, local_size_z = 10) in;

void main()
{
	uint index = gl_LocalInvocationIndex;

	if (index >= max_count)	return;

	num_arr[index] += const_number;
	num_arr[index] -= const_number;
	num_arr[index] *= const_number;
	num_arr[index] /= const_number;
	num_arr[index] += 10;
}

まとめ

　今回はGPUでの処理方法を試しました。記事に書くにあたって、ほとんど全て初めて使う技術ばかりでしたので、時間と労力がかかってしまいました。ですが、それ以上にいい経験を得ることが出来たとは思っています。
　今回の速度トップは実用的な使い方ではC++AMPが優勝、最速はOpenGLという事になりました。そう考えると、OpenGLのオーバーヘッドがいかに高いかが分かりますね。そして、あくまで自分の環境での結果という事は理解するようにお願いします。

　実際にゲーム制作などで使う際はCUDAかC++AMPを使う事がほとんどになるとは思いますが、環境的に無理な場合などはOpenCL辺りが無難と思われます。事実、CUDAの場合コンパイルはVSコンパイラーでコンパイル時に行うので、汎用エンジンなどは難しい事になるはずです。ですが、OpenCLは実行中にコンパイルするのでそういった点は楽になるはずです。なので、選ぶ基準は「速度＜環境」になり、使い分けてGPGPU生活をしていきましょう！

次回へ

　今回はGPU側の処理はほとんど書いておらず、試験的な意味合いが強い検証でした。なので、次回は実践的な場合の時間検証を試します。

　おまけなんですが、CUDAやOpenCLなどはカーネル関数内でのprintf()に対応しており、GPU側の処理後に直接出力して見ることが可能になっています。そういう機能は何気にあったほうが便利なのでありがたい機能ですね。
　それと、クロノスグループさん、早くOpenMP(GPU)・OpenACCをWindowsに正式対応して下さい（願望）。

  1. 世界中はCPUよりもGPUの方が中心なのでしょうか？ ↩

  2. 開発はクロノス・グループが担当しています。もし、Appleが「開発」もしていたら、絶対ロイヤルフリーではないでしょうね（笑） ↩

  3. 「Intel(R) oneAPI DPC++ Compiler 2021.2.0」を導入して、プロパティーから「Generate device binary for Intel(R) GPUs (/Qopenmp-targets:spir64_gen)」の設定をONにすると「Device name missing.」か「region cannot be nested inside 'target' region」とエラーを吐きます。教えてください詳しい人...。 ↩

  4. というのも、C/C++コード・OpenMPコード共にVSCode上でビルド・デバッグも出来て動かせたのですが、どうみてもGPUで動いていないので訳わかめ状態になっています。 ↩

  5. 現状VS2019はOpenMP 2.0が標準サポートされています。つまりSIMDを除いてOpenMP 4.0はサポートされていません。因みに、SIMDを使うには、C/C++コマンドラインに「/openmp:experimental」を記入してください。ですが、元から対応しているコンパイラの場合は★１判定ですね。 ↩

  6. というのも、NVIDIAのサイトからダウンロード出来るのですが、対応がLinuxのみとなっています。一番下にあるのですが、「今後対応するね！」と書かれていて現状無理ということなのでしょう。 ↩

  7. これもOpenCL同様にクロノス・グループが開発を担当しています。 ↩

  8. あと、導入がとてつもなくめんどく(ry ↩

11
7
0

Register as a new user and use Qiita more conveniently

  1. You get articles that match your needs
  2. You can efficiently read back useful information
  3. You can use dark theme
What you can do with signing up
Sign upLogin
11
7

Delete article

Deleted articles cannot be recovered.

Draft of this article would be also deleted.

Are you sure you want to delete this article?