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

# 最初に
　この記事は[前編](https://qiita.com/HnniTns/items/70ae1e6868412a9698f3)の続きになっています。ですので、事前コードなど詳細は[前編](https://qiita.com/HnniTns/items/70ae1e6868412a9698f3)を参照して下さい。

## 環境
+ エディター：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の中での星判定なので注意してください。[^0]）
+ 参考サイト
    + [C++ AMPによるGPGPU入門](https://qiita.com/Riyaaaa_a/items/40054b893e70b54f5a26)
    + [C++ AMP の概要](https://docs.microsoft.com/ja-jp/cpp/parallel/amp/cpp-amp-overview?view=msvc-160)（公式サイト）
    + [C++ AMPを用いたGPUプログラミングwiki](https://aeroastro.sd.tmu.ac.jp/hydrodynamics/main/colums/AMPindex.html)
+ [コードへ](#c-amp)

[^0]: 世界中はCPUよりもGPUの方が中心なのでしょうか？

#### 2. CUDA

+ NVIDIAが開発しているGPGPU処理方法で、C言語の書き方で出来るのが使いやすいポイントです。
+ 難易度
    + 導入コスト：★★☆（ライブラリをダウンロードして、パスを通さなければなりませんが、慣れている人なら簡単ですよね？）
    + 学習コスト：★★☆（変数のメモリー管理や関数の書き方など覚えることがややありますが、慣れれば簡単にかけるはず。）
    + 資料数　　：★★★★（CUDAだけ検索ヒット数が一桁違いますね。これだけあれば、参考資料には困らないと思います。）
+ 参考サイト
    + [CUDAを使ったGPUプログラミング超入門](https://qiita.com/Keech/items/381495ed90e012d69f1a) 
    + [CUDA C++でGPUプログラミング【CUDA 第1回】](http://tecsingularity.com/cuda/cudaprogram1/)
    + [CUDAの導入，CUDAの基礎](https://co-crea.jp/wp-content/uploads/2016/07/File_2.pdf)
    + [CUDAプログラミングの基本](https://http.download.nvidia.com/developer/cuda/jp/CUDA_Programming_Basics_PartII_jp.pdf)
    + [繰り返し処理をCUDAで書く（配列同士の足し算）](https://qiita.com/wazakkyd/items/8a5694e7a001465b6025)
    + [GPUコンピューティング（CUDA)講習会](http://gpu-computing.gsic.titech.ac.jp/Japanese/Lecture/2010-06-28/multi-gpu.pdf)
    + [CUDAコードの実行時間を測定する方法のまとめ](https://qiita.com/syo0901/items/7ea3b8dfc01fd5cc2cf4)
+ [コードへ](#cuda)

#### 3. OpenCL

+ 実はAppleが提案した処理方法で、幅広い分野で使われるものです[^2]。即ち、GPGPUに特化していません。
+ 難易度
    + 導入コスト：★★★（ライブラリをダウンロードする為にIntelに登録しなければならないのがメンドクサイですね...。）
    + 学習コスト：★★★（コンテキスト・コマンドキューやカーネル関数の読み込み・ビルド・引数の設定など手動で書くことが多く、結構時間がかかりました。恐らく一番コード量が多いと思います。）
    + 資料数　　：★★★（様々な用途で使われるのでかなり多いです。流石です。）
+ 参考サイト
    + [C++でOpenCL(環境構築編)](https://qiita.com/hkiribayashi/items/5ad6253661f1c43a8381)
    + [C++でOpenCL(使ってみよう編)](https://qiita.com/hkiribayashi/items/acf04e5fbbc6dabde1bc)
    + [9. OpenCL](http://www.e-em.co.jp/tutorial/chap9.htm)
    + [改訂新版 OpenCL入門向けサンプルコード](https://www.fixstars.com/ja/company/books/sample/)
    + [NVIDIAのドライバでOpenCLをやる](https://qiita.com/GRGSIBERIA/items/d8290c1a3ded59f4075a)
    + [10.4. 組み込み関数](https://mkguytone.github.io/java-opencl/ch10s04.html)
    + [OpenCLでいくぜ！　(8)](https://ameblo.jp/xcc/entry-10799610964.html)
    + [OpenCLと共通コード（１） 基礎編](https://research.kek.jp/people/matufuru/Research/Docs/Programming/OpenCL/Slide/lattice-code_opencl-1.pdf)
    + [ワークアイテム・ワークグループ・次元数について](http://neareal.net/index.php?Programming%2FOpenCL%2FDimentionWorkGroupWorkItem)
+ [コードへ](#opencl)

[^2]: 開発は[クロノス・グループ](https://ja.wikipedia.org/wiki/%E3%82%AF%E3%83%AD%E3%83%8E%E3%82%B9%E3%83%BB%E3%82%B0%E3%83%AB%E3%83%BC%E3%83%97)が担当しています。もし、Appleが「開発」もしていたら、絶対ロイヤルフリーではないでしょうね（笑）

#### 4. OpenMP

+ 前回登場した``OpenMP``君です。記事を見て知ったのですが、``OpenMP``を使ってGPU処理が出来るんですね。ただし、OpenMP4.0以上から対応なので、「OpenMP gpu」「OpenMP offloading」などで検索しないとあまり情報が出てきません。<font color="red">**導入する為にインテルコンパイラやGCC(G++)を使って色々と手を尽くしたのですが、残念ながら前者はビルドが通らず[^4]、後者はGPUで動かすことが叶わず[^5]、時間をこれ以上割けないので諦める事にしました。悔しいです。**</font>
+ 難易度
    + 導入コスト：★★★★（設定を変更してインクルードするだけと思いきや、VS2019では``OpenMP SIMD(4.0の一部機能)``までしか対応していません。[^3]）
    + 学習コスト：★★☆（サンプルを見て判断した評価となっています。）
    + 資料数　　：★★☆（CPUでの``OpenMP``とGPUでの``OpenMP``では検索数が大きく異なります。やはり「OpenMP＝CPU」のイメージが強いのでしょうか？）
+ 参考サイト
    + [OpenMP* 4.x における拡張](http://jp.xlsoft.com/documents/intel/seminar/openmp_vector_offload/02_What%20is%20OpenMP%2040%20offload_Jan18.pdf)
    + [第３回 OpenMPの基礎](https://www.r-ccs.riken.jp/wp/wp-content/uploads/2020/09/katagiri190425.pdf)
    + [OpenMP 並列プログラミング入門](https://www2.ccs.tsukuba.ac.jp/workshop/HPCseminar/2013/material/2013-03-openmp.pdf)
    + [OpenACC vs OpenMP GPU Offloading (環境構築編)](https://proc-cpuinfo.fixstars.com/2020/06/openacc-vs-openmp-gpu-offloading/)
    + [Programming Your GPU with OpenMP](https://openmpusers.org/wp-content/uploads/uk-openmp-users-2018-OpenMP45Tutorial_new.pdf)
    + [IBM ドキュメンテーション（#pragma omp target data）](https://www.ibm.com/docs/en/xl-c-and-cpp-linux/13.1.5?topic=parallelization-pragma-omp-target-data)
    + [OpenMP 5.0 API シンタックス・クイック・リファレンス・カード](https://www.openmp.org/wp-content/uploads/OpenMPRef-5.0-0519-JA.pdf)
+ [コードへ](#openmpgpu)（といっても、私の環境ではコンパイル・コード実行が出来ないので、サンプルだと思ってください。）
+ 備考
    + 下記の[OpenMP](#openmpgpu)のコードをVS2019で実行するには、``Intel C++ Compiler``をダウンロード・インストールしなければなりません。ですので、とてつもなくメンドクサイ事になります。しかも、これで動けば良かったのですが...。
+ インテルコンパイラ導入などの参考サイト
    + [インテル® C++ コンパイラーの使用](https://www.xlsoft.com/jp/products/intel/compilers/manual/cpp_all_os/GUID-0DFD85A7-AAC8-44E2-9CF0-5B015F4F8288.htm)
    + [Free Intel® Software Development Tools](https://software.intel.com/content/www/us/en/develop/articles/free-intel-software-developer-tools.html)
    + [oneAPI-samples](https://github.com/oneapi-src/oneAPI-samples)（Intel® oneAPI DPC++/C++ Compilerのサンプル）
    + [Get Started with the Intel® oneAPI DPC++/C++ Compiler](https://software.intel.com/content/www/us/en/develop/documentation/get-started-with-dpcpp-compiler/top.html)

追記：↓
![現時点では、OpenMPオフロードターゲットはWindowsではサポートされていません。.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/40bc01d0-e60d-3aa0-3b2a-303e481c8a13.png)
**う～ん？**
![現時点では、OpenMPオフロードターゲットはWindowsではサポートされていません。- 集中線.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/19602da2-0c64-af5b-b6cb-a4abe500e369.png)
以上！現場からでした！（私の数日間返して...）
※[README.md参照元](https://github.com/oneapi-src/oneAPI-samples/tree/master/DirectProgramming/DPC%2B%2B/DenseLinearAlgebra/matrix_mul)
※[集中線ダウンロードサイト](https://mangasozai.com/archives/1186)

[^3]: 現状VS2019は``OpenMP 2.0``が標準サポートされています。つまりSIMDを除いて``OpenMP 4.0``はサポートされていません。因みに、``SIMD``を使うには、C/C++コマンドラインに「``/openmp:experimental``」を記入してください。ですが、元から対応しているコンパイラの場合は★１判定ですね。
[^4]: 「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」とエラーを吐きます。教えてください詳しい人...。
[^5]: というのも、C/C++コード・OpenMPコード共にVSCode上でビルド・デバッグも出来て動かせたのですが、どうみてもGPUで動いていないので訳わかめ状態になっています。

#### 5. OpenACC

+ 「Open～」シリーズで分かる通り、OpenMPをGPUに特化させたライブラリになります。なんと、NVIDIAなど系４社が開発に関わっています。そして、ありがたいことに、OpenMPに似た書き方で書くことが可能です。凄いです！
+ 難易度
    + 導入コスト：★★★★（<font color="red">**残念ながら、私の環境では導入は不可能みたいです。**</font>[^6]）
    + 学習コスト：★☆☆（といっても、サンプルを見た感じの評価となっていて、より信憑性が無くなっているのは気のせいという事にしておきます。）
    + 資料数　　：★☆☆（これも、OpenMP（GPU）の半分ぐらいの情報量ですが、こちらはGPU特化なので欲しい資料が割と出てきます。）
+ 参考サイト
    + [OpenACC ディレクティブによるプログラミング](https://hpcworld.jp/archive/SPG/Pgi/OpenACC/002.html)
    + [第1回 FX100システム利用型 OpenMP講習会（初級）](https://icts.nagoya-u.ac.jp/ja/sc/pdf/lecture_20201007_openacc_main.pdf)
    + [OpenACCの紹介・Reedbush-Hお試し](https://www.cspp.cc.u-tokyo.ac.jp/hanawa/class/spc2019a/sp20200114.pdf)
    + [OpenACC・CUDAによるGPUコンピューティング](https://www.r-ccs.riken.jp/wp/wp-content/uploads/2020/09/0719_GPU_computing.pdf)
    + [OpenACC ディレクティブによるプログラミング](https://hpcworld.jp/archive/SPG/Pgi/OpenACC/index.html)
+ [コードへ](#openacc)（といっても、私の環境ではコンパイル・コード実行が出来ないので、サンプルだと思ってください。）

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

#### 6. OpenGL(Compute Shader)

+ 皆さんも聞いた・触った事があるはずのOpenGLです[^7]。OpenCLとは違って名前の通りグラフィック機能がメインの機能です。Ver4.3からCompute Shaderが実装されました。歴史は古く、OpenCLの初版が2008年なのに対し、OpenGLの初版は1992年となっています。（Wiki参照）
+ 難易度
    + 導入コスト：★★☆（OpenGL単体は標準でサポートされており、ですが動かそうとすると他のライブラリを導入する事になると思うのでこの評価です。）
    + 学習コスト：★★★（直接関係無いところも含めて、覚える事がOpenCLより多いです。仕方ないですが、本来は描画目的なので仕方ありません。（そう考えれば、DirectX11よ(ry ））
    + 資料数　　：★★☆（OpenMP(GPU)と同じぐらいの量になりますね。やはり、比較的最近出た機能なので検索数が少ないのでしょう。）
+ 参考サイト
    + [OpenGL Compute Shader を実行するだけのコード](https://kakashibata.hatenablog.jp/entry/2020/08/31/001306)（このサイトのコードを改造して使用しています。kakashibata様ありがとうございます。）
    + [[OpenGL][GLSL] 粒子のレンダリング (2) ポイントの移動](https://marina.sys.wakayama-u.ac.jp/~tokoi/?date=20181018)
    + [GLFWホームページ](https://www.glfw.org/)
    + [GLEWダウンロードサイト](https://sourceforge.net/projects/glew/files/glew/2.1.0/)
    + [Compute Shader](https://www.khronos.org/opengl/wiki/Compute_Shader)（OpenGLの公式サイト）
    + [モダンな OpenGL で頂点モーフ](https://techblog.sega.jp/entry/2016/10/27/140454)
    + [Shader Storage Buffer Object](https://www.khronos.org/opengl/wiki/Shader_Storage_Buffer_Object)（OpenGLの公式サイト）
+ [コードへ](#openglcompute-shader)

[^7]: これもOpenCL同様に[クロノス・グループ](https://ja.wikipedia.org/wiki/%E3%82%AF%E3%83%AD%E3%83%8E%E3%82%B9%E3%83%BB%E3%82%B0%E3%83%AB%E3%83%BC%E3%83%97)が開発を担当しています。

#### 7. SYCL

+ このライブラリは紹介だけします。というのもかなり多くなっているので...[^8]。このライブラリはOpenCLとCUDAをサポートしているフロントエンドライブラリになっています。実は、C++標準規格に選ばれるというのを最終目標にしているようです。
+ 参考サイト
    + [SYCL* と DPC++ を始めよう](https://www.isus.jp/products/oneapi/hello-sycl-and-dpcpp/)
    + [SYCL公式サイト](https://www.khronos.org/sycl/)
    + [SYCL公式Github](https://github.com/KhronosGroup/Khronosdotorg/blob/main/api/sycl/resources.md)

[^8]: あと、導入がとてつもなくめんどく(ry

---

https://qiita.com/aokomoriuta/items/777f00f4d9f2bec887ba

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

# 結論

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

## メモリ管理も含めた場合（確保・転送）
　この欄はメモリ管理を含めて時間を計測しています。GPGPUに置ける大体のオーバーヘッドの部分を含むやり方となっております。

![メモリ管理（確保・転送）.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/31e4a84d-2296-88e4-4712-d019c1f0b692.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](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/cce04f1e-39bb-1f38-85c6-9a4531268ccd.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``に関してはメモリ確保でかなり時間を取られていたことが分かりますね。

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

![処理部分のみ.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/a856763d-a0bf-a9ab-3386-bafe4abbc5bf.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)|

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

## CPU・GPUの平均値
![CPU・GPU.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/3c4fcb29-f8dc-3cd8-2129-3c9cef21d2da.png)
　このグラフは各GPU・CPU処理方法をまとめた上で平均値にしてグラフ化したものになっています。こう見ると、「処理部分のみ」の軽さが際立っています。やはり、GPGPUにおけるメモリ確保・転送などの「オーバーヘッド」はかなり大きのが分かります。
　そして、CPUとGPUを比べてみると、CPUとGPUにおいて中身が無いと、一部を除いてほとんど経過時間は変わらないということも分かりました。

## 各項目の順位
![速度.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/6c2f690f-6bfd-e8b6-7d1a-e5aa6137f935.png)
![導入コスト.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/ab74d849-c339-4075-f4de-e57fb2c3b350.png)
![学習コスト.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/de18395a-d608-5fcb-b636-a06679204635.png)
![資料数.png](https://qiita-image-store.s3.ap-northeast-1.amazonaws.com/0/540257/e204c076-b436-bb66-6e84-ca7ac9673270.png)
　速度に関しては、「[メモリ管理も含めた場合（転送のみ）](#メモリ管理も含めた場合転送のみ)」を参照しています。そして、他の項目に関しては、私の体感・経験を基にしているデータなので、皆さんは参考程度にして下さい。

# コード

```cpp: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
```cpp: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

<details><summary>長くなったので折りたたんでいます</summary><div>

```cpp: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");
```
```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);
}
```
```cuda:cudatest.cuh
void CudaCalc(const dim3& grid, const dim3& block, int* d_mem, const int const_number, const int max_count);
```
```cuda: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
```

</div></details>

## OpenCL
<details><summary>長くなったので折りたたんでいます</summary><div>

```cpp: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");
```
```c: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;
}
```
```c: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
```

</div></details>

## OpenMP(GPU)

<details><summary>少し長くなったので折りたたんでいます</summary><div>
恐らく、下記のようなコードになるとは思います。ですが、<font color="red">**コンパイルと実行が出来ていないので正しいコードとは言えませんし、実際にGPUで動いているのを確認も出来ていない**のでその点だけ注意して下さい。</font>

```cpp: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](https://www.ibm.com/docs/en/xl-c-and-cpp-linux/16.1.0?topic=parallelization-pragma-omp-distribute-parallel)（参照サイト）

```cpp: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];
}
```
</div></details>

## OpenACC

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

```cpp: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 を使って、まず始めてみよう](https://hpcworld.jp/archive/SPG/Pgi/OpenACC/004.html#acc_data)（参照サイト）

```cpp: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);
```
</div></details>

## OpenGL(Compute Shader)

<details><summary>長くなったので折りたたんでいます</summary><div>

```cpp: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");
```

```glsl: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;
}
```

</div></details>

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

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

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

---

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