C++ AMP (C++ Accelerated Massive Parallelism) C++ AMP (C++ Accelerated Massive Parallelism)とは、マイクロソフトが中心になって開発していた、演算アクセラレータを活用した並列プログラミングのためのC++言語拡張およびライブラリである。バックエンドとしてDirectCompute (DirectX) やOpenCLなどを使用した実装が存在している。C++ AMPによって、DirectComputeやOpenCLよりも高レベルの並列プログラミングAPIが提供される。

概要 マイクロソフトはWindows上においてGPGPUを標準化するために、DirectX 11にてDirectComputeを導入した。しかし、DirectComputeはDirect3D APIを直接操作する必要があるため、OpenMPを利用した一般的なCPU向けのC++並列プログラミングなどと比較してハードルが高い4Gamer.net ― 見えてきたAMDの次世代GPUアーキテクチャ。なぜAMDはVLIWを捨てるのか。マイクロソフトはより抽象化された純粋C++プログラミングに近いGPU並列プログラミングモデルを提供するAPIおよび言語拡張として、2012年にC++ AMPの仕様を公開したMicrosoftがC++ AMP仕様を公開し、データ並列処理の障壁を下げる意向。また、マイクロソフトによるC++ AMPの公式プロプライエタリ実装として、Microsoft Visual C++ 2012はDirectComputeをバックエンドとした最初のC++ AMP対応コンパイラとなっている。

C++ AMPのライブラリおよびC++ AMPに対応したVisual C++コンパイラが付属しているのはVisual Studio 2012以降であるが、Visual Studio 2013ではC++ AMP v1.2規格に対応し、さらにコンピュートシェーダー以外のシェーダーステージとの相互運用やテクスチャフィルタリング(サンプラー)などの機能が強化されているGraphics (C++ AMP) | Microsoft Learn Texture Copy Improvements in C++ AMP with Visual Studio 2013 - Parallel Programming in Native Code - Site Home - MSDN Blogs C++ AMP Open Spec V1.2 Published - Parallel Programming in Native Code - Site Home - MSDN Blogs。しかしVisual C++ 2022以降、C++ AMPは非推奨となったC++ AMP Overview | Microsoft Learn。

なお、C++ AMP自体はDirectCompute/Windows専用というわけではなく、またGPU専用というわけでもないaccelerator Class。AMDとマイクロソフトは、OpenCLなどをバックエンドとするLinuxやOS X向けのC++ AMP v1.2オープン実装を発表したAMDとMS,GPU演算用途向けのコンパイラ「C++ AMP v1.2」を発表 - 4Gamer.net。その後、同コンパイラはMulticoreWareを経てRadeon Open Compute (ROCm) プラットフォームに移管され、Clamp, Kalmar, HCCと名前を変えながら開発が続けられていたが、2019年6月予定の最終リリースをもってHCCの開発が終了することになったRadeonOpenCompute/hcc: HCC is an Open Source, Optimizing C++ Compiler for Heterogeneous Compute currently for the ROCm GPU Computing Platform。HCCでは、OpenCLのカーネル記述言語"OpenCL C"や中間表現"SPIR"SPIR - The first open standard intermediate language for parallel compute and graphics、およびHSA (Heterogeneous System Architecture) の中間言語"HSAIL"HSA Developer Tools - HSA Foundationへの変換層が提供され、Visual C++版よりもさらにヘテロジニアス環境向けを意識した実装となっている。

MSDNブログではCUDA、OpenCL、DirectComputeプログラマー向けに、各APIに対応するC++ AMP相当機能の比較資料が公開されていた。 Microsoft Learnにおけるアーカイブ C++ AMP for the CUDA Programmer - Parallel Programming in Native Code - Site Home - MSDN Blogs C++ AMP for the OpenCL Programmer - Parallel Programming in Native Code - Site Home - MSDN Blogs C++ AMP for the DirectCompute Programmer - Parallel Programming in Native Code - Site Home - MSDN Blogs Internet Archive C++ AMP for the CUDA Programmer | Microsoft Docs C++ AMP for the OpenCL Programmer | Microsoft Docs C++ AMP for the DirectCompute Programmer | Microsoft Docs

コード例 C++ AMPのコード例を示す。OpenMPの#pragma omp parallel forのように、並列処理するコード単位すなわちカーネルを、C++11のラムダ式を使ってインラインで記述している部分が特徴的である。デフォルトのアクセラレーターを使用する場合、OpenCLやDirectComputeのような冗長なデバイス準備処理の記述も不要となる。

なおVisual C++実装では、実行環境にDirectX 11対応のGPUが搭載されている場合、restrict(amp)句が指定されたコードの実行にはGPUが使用されるC++ AMP Overview。対応GPUが搭載されていない場合は、Direct3D WARPデバイスやREFデバイスなどによるCPUエミュレーション実行(フォールバック)が提供されるUsing accelerator and accelerator_view Objects。

等差数列を作成する例 #include <amp.h> using namespace concurrency;

int main() { const int ElemCount = 100; int ary[ElemCount]; array_view<int, 1> av(ElemCount, ary);

parallel_for_each(extent<1>(ElemCount), [=](index<1> idx) restrict(amp)

int i = idx; av[i] = i; }); av.synchronize(); }

多次元配列の例 #include <cstdio> #include <amp.h> #include <amp_graphics.h>

using namespace concurrency; using namespace concurrency::graphics;

int main() { const int width = 10; const int height = 3; const int depth = 2; const int elemCount = width * height * depth;

std::vector<float_3> srcHost(elemCount); std::vector<float_3> dstHost(elemCount);

for (int z = 0; z < depth; ++z) for (int y = 0; y < height; ++y) for (int x = 0; x < width; ++x)

srcHost[z * height * width + y * width + x] = float_3(float(x), float(y), float(z));

array_view<float_3, 3> srcView(depth, height, width, srcHost); array_view<float_3, 3> dstView(depth, height, width, dstHost);

parallel_for_each( dstView.extent, [=](index<3> idx) restrict(amp)

const auto globalx = idx; const auto globaly = idx; const auto globalz = idx; auto temp = srcView[idx]; if (globaly == 0)

temp += 0.5f;

if (globalz == 0)

temp *= -1;

dstView[idx] = temp; });

dstView.synchronize(); for (int z = 0; z < depth; ++z) for (int y = 0; y < height; ++y) for (int x = 0; x < width; ++x)

const auto val = dstHost[z * height * width + y * width + x]; printf("[%d][%d][%d] = (%.1f, %.1f, %.1f)\n", z, y, x, val.x, val.y, val.z);

関連項目 DirectCompute OpenCL CUDA Microsoft Visual C++ OpenACC

参照

外部リンク C++ AMP (C++ Accelerated Massive Parallelism) - MSDN (英語)

multicoreware / hcc / wiki / Home — Bitbucket - オープン実装のC++ AMPコンパイラ