syghの新フラグメント置き場

プログラミングTipsやコード断片の保管場所です。お絵描きもときどき載せます。

NVIDIAのOpenCL対応状況

C++ OpenCL プログラミングTips GPU CUDA Direct3D DirectX

NVIDIAのOpenCLドライバーはIntelやAMDと比べて規格への対応が遅く、2015年2月時点でもOpenCL 1.1どまりです。NVIDIAのOpenCL SDKとなるCUDA Toolkitも、7.0時点でOpenCL 1.1までしか対応していません（cl.hに定義されているのはCL_VERSION_1_1まで）。例えばOpenCL 1.2のimage3d書き込み拡張機能や、OpenCL 2.0のread_write修飾子には対応していません。

英語版 Wikipedia (2015年3月閲覧) には、Kepler/Maxwell 世代の GeForce は OpenCL 1.2 に対応していると記載されているんですが、どうもこれは Linux に限定した話のように思われます。

いくつかのドライバーバージョンを Fermi/Kepler 世代の GPU とともに試してみましたが、少なくとも Windows の 340.62、344.11、347.25、347.88 ドライバーでは、NVIDIA の OpenCL プラットフォーム バージョン (CL_PLATFORM_VERSION) は 1.1 どまりです。なお NVIDIA のドライバーにはベータ ドライバー、開発者向けドライバー、エンドユーザー向け安定ドライバーなど、いくつかのブランチがあるため、単純に数字が大きいほうが新しいとは限りません。
ちなみに Linux 用のドライバーは Windows 版と違ってまともなリリースノートがなく、したがって対応する OpenGL/OpenCL バージョンが記載されているようなまともなドキュメントが存在しない模様。

英語版 Wikipedia で引用されている下記のデータベース サイトも信憑性が疑われます（個人的には DirectX の対応バージョンに Feature Level すなわち機能レベルが記載されていない時点で、すでにもう信用ならない）。

※2015-05-05追記：
Windows版ドライバーもバージョン350.12でOpenCL 1.2に正式対応した模様です。ただしKeplerおよびMaxwell世代のGeForceのみで、QuadroおよびTeslaはまだ対応していません。またFermiはDirectX 11対応ですがOpenCL 1.1どまりとなるようです。なお、OpenCL 1.2拡張のcl_khr_spirおよびcl_khr_3d_image_writesはサポートされないようです。ちなみにCUDA Toolkitのほうは7.0時点でOpenCL 1.2に対応していません。OpenCL 1.2のヘッダーはKhronosのサイトからダウンロードするとして、新しい関数エントリポイントはどうやって取得するんでしょうか？　LoadLibrary()+GetProcAddress()ではないだろうし、clGetExtensionFunctionAddress()も違うだろうし……OpenCL.dllにエクスポートされているのはOpenCL 1.0の関数のみであり、またOpenCLにはInstallable Client Driver (ICD) Loaderによって、OpenCL APIコールを特定のプラットフォーム実装にディスパッチする仕組みがあるんですが、Khronosが公開しているICD Loaderのtgzパッケージにはdefファイルが含まれていました。これを使ってlibファイルを作って遅延バインドしろということなんでしょうか……OpenCL 1.2を使いたい場合、ターゲット環境がNVIDIAであってもIntelあるいはAMDのSDKを使って開発したほうがいいかもしれません。

※2015-07-10追記：
Quadro/Teslaも353.06ドライバーでOpenCL 1.2に正式対応したようです。こちらもKeplerおよびMaxwellのみとなります。

※2015-12-30追記：
CUDA Toolkitも7.5でOpenCL 1.2に正式対応したようです。リリースノートには何も書かれていませんが、cl.hにはCL_VERSION_1_2が定義されるようになったので、OpenCL 1.1までしか対応していないFermi世代以前の旧製品や旧ドライバーでも実行可能なOpenCLアプリケーションを開発する場合は注意が必要です。なお、CUDA Toolkit 7.5.18付属のOpenCL 1.2ドライバーには致命的なバグがあり、clRetainDeviceなどのOpenCL 1.2関数エントリポイントをICDローダー経由で正常に取得できない現象が発生します。

OpenCL-C コンパイラーの挙動

OpenCL 対応が 1.1 までというのは GeForce ドライバーでも Quadro ドライバーでも同じなのですが、今回調査してみた結果、ドライバーによって OpenCL-C コンパイル エラー発生時の挙動がずいぶん異なるようです。

具体的には、例えば下記のカーネルをコンパイルする際、Quadro ドライバーだと read_write 周辺できちんとコンパイル エラーになるものの、GeForce ドライバーだとコンパイル エラーというよりは PTX アセンブル エラー*1になるようです。

const sampler_t myLinearSamplerObj = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR;

__kernel void myOclkDoProc(
  __read_write image2d_t inoutImage
)
{
  // For Compile Test
  const float4 oldValue = read_imagef(inoutImage, myLinearSamplerObj, (float2)(0, 0));
  write_imagef(inoutImage, (int2)(0, 0), oldValue + (float4)(1, 0, 0, 0));
}

結果（CL_PROGRAM_BUILD_LOG）：

[Driver 340.62, on Quadro 2000]
:RRR:CCC: error: unknown type name '__read_write'

Quadro では、clBuildProgram() は CL_BUILD_PROGRAM_FAILURE で失敗します。これは想定範囲内。

一方、GeForce のほうは OpenCL-C 自体のコンパイルは通ってしまいます。とはいえ、最終的にアセンブル エラーが発生してしまうため、カーネルを動作させることはできません。しかも read_write を指定したオブジェクトに対する read_imagef() 呼び出しがアセンブル エラーのトリガーになるようで、write_imagef() はエラーにならないようです。

[Driver 347.25, on GeForce GTX 770]
ptxas application ptx input, line XXX; error   : Argument 1 of instruction 'tex':
 .texref or .u64 register expected
ptxas fatal   : Ptx assembly aborted due to errors

clBuildProgram() は CL_INVALID_BINARY で失敗します。

OpenCL 1.1 コンパイラーとして正しい挙動なのは Quadro ドライバー実装のほうです。GeForce はひどすぎです。ハードウェアやプラットフォームの違いを吸収するのが OpenCL のような標準 API に与えられた役目であるはずなのに、いきなり PTX アセンブル エラーとか言われてもなんだよそんなの知らねーよって感じです。

こういったドライバーの出来不出来の差というのは何も OpenCL ドライバーに始まったわけではなく、OpenGL ドライバーにも古くから存在します。中にはそもそも GeForce ではサポートされず、Quadro でしかサポートされない GL ベンダー拡張もあったりするのでさらに厄介です。自分が現行の OpenCL や OpenGL/GLSL を毛嫌いしているのは、このようにハードウェア ベンダーに好き勝手を許してしまっている無責任でいい加減な規格だからです。よく言われるポータビリティや互換性というのは無知な開発者をだます甘言でしかありません。

そもそも公式のオフライン コンパイラーもまともに用意されていない三流の仕様を、アクセラレータ系 API の標準規格として認めている時点でどうかしています（nvcc や fxc に相当するものがないのは致命的すぎる）。なぜ OpenCL-C 規格を作る前にバイトコード規格（中間言語）を作らないのかと。最初にバイトコード規格を作っておけば C 以外でもカーネルコードを生成できる独自言語を自由に開発できる可能性も出てくるんですが、いきなり高級言語の規格を作って実装をベンダーに投げようとするあたり本当に頭がおかしいです。ちなみに OpenCL の中間言語 SPIR は後出しの規格で、SPIR 1.0 の暫定規格は 2012-08-24 に発行されています。現行規格では OpenCL 1.2 の SPIR 1.2、そして OpenCL 2.0 の SPIR 2.0 が存在しますが、それぞれ必須サポート機能ではなく拡張機能扱い（cl_khr_spir）でしかないのが非常に残念です。もともと OpenCL を提唱したのは Apple なんですが、OpenCL 1.0 時点で相当に見積りの甘い API 設計だったとしか言いようがありません。CUDA や DirectCompute と比較するとおもちゃレベルです。

ハードウェアを抽象化してソフトウェアを開発しやすくするのが標準規格とAPIの最大の役目なのに、ベンダーに拡張を許したりシェーダーやカーネルのコンパイラを個別に実装させたりするなど、愚行にもほどがあります。時代錯誤も甚だしい。

OpenCL 1.2, 2.0 対応に関して

OpenCL 1.2 の image3d 書き込みと OpenCL 2.0 の read_write 修飾子は、CUDA の Surface や Direct3D/HLSL (DirectCompute) の RWTexture に近い機能で、DirectX 11 世代の GPU であれば普通にハードウェア レベルでサポートしているはずの機能ですが、ドライバーが対応していなければ意味がありません。この write_image(image3d_t) と read_write がないと、CUDA プログラムや HLSL コンピュート シェーダープログラムを OpenCL に移植するのに結構困ります*2。しかも image3d 書き込みに至っては標準機能ですらなく、ただの拡張機能扱いです（cl_khr_3d_image_writes）。GPGPU においては基本中の基本と思われる重要機能を拡張扱いにするなど、Khronos の連中がいったい何を考えているのかもうさっぱり分かりません。

OpenCL 拡張に関する不満はとりあえずおいておくとして、現時点では少なくとも OpenCL 2.0 が使えるかどうかというのは移植性の観点からも非常に重要になってくるんですが、2015年2月現時点での最新規格である OpenCL 2.0 をサポートしているのは「AMD APP SDK 3.0 Beta」と「Intel SDK for OpenCL Applications 2014」くらいです。
AMD は GCN アーキテクチャ（GCN 1.1世代以降）の GPU/APU で OpenCL 2.0 をハードウェア サポートしているものの、対応ドライバーの RC 版がリリースされたのは2014年9月になってからで、さらに SDK が例のごとく未だベータ ステータスということからも分かるように、正直ドライバー品質に関しては信頼できません。OpenCL 2.0 規格発表は 2013-11-18 なんですが、同規格にようやく対応した AMD APP SDK 3.0 Beta の発表は 2014-12-09 です。正直対応が遅すぎです。1年もかけていながら、お前らいったい何をやっていたのかと。
http://developer.amd.com/community/blog/2014/12/09/amd-app-sdk-3-0-beta/

なお OpenCL 2.0 対応の AMD ドライバーは、Windows に関しては現状 Win8.1 x64 のみサポートです。Win7 向けには提供されていません。マルチ GPU もサポートしてないらしいです（CrossFire 無効で AMD グラフィックスカードを2基接続しても、GPGPU に使えるのは1基だけ？）。
http://support.amd.com/en-us/kb-articles/Pages/OpenCL2-Driver.aspx

ちなみに AMD は OpenCL のソフトウェア エミュレーターをオープンソース開発しているようですが、1.2 どまりです。2.0 には対応していません。開発も2012年10月を最後に停滞していて、相変わらず仕事が中途半端な印象の AMD です。目の付けどころは悪くないんですが……
http://developer.amd.com/tools-and-sdks/opencl-zone/opencl-emulator-debugger/

Intel も OpenCL 2.0 をハードウェア サポートしているのは最新世代（Broadwell アーキテクチャ）のモバイル プロセッサーである Core M などのみとなっています。
http://www.isus.jp/article/intel-software-dev-products/opencl-sdk/
https://software.intel.com/en-us/intel-opencl

そして哀しいかな、HPC 分野でも、すでに OpenCL には見切りをつけて冷ややかな目を向け始めているとのこと。
http://news.mynavi.jp/articles/2014/12/16/gpu_mic/003.html

正直な話、現時点で GPGPU を本気でやろうとするのであれば、OpenCL より CUDA や DirectCompute を選択することをお勧めします*3。ベンダーロックイン？　知ったこっちゃないですね。低品質のドライバーのデバッグに付き合わされるよりは遥かにマシでしょう。とはいえ個人的には、NVIDIA の OpenCL に対するスタンスが気に食わないのは確かです。ベンダーロックインによる囲い込みばかりを重視するあまり、標準規格の推進・強化・準拠がおろそかになっています。邪推ですが OpenCL のサポートは意図的に手を抜いているんじゃないかと。CUDA が NVIDIA Cg (C for Graphics) と同じ末路をたどるとは思いませんが、CUDA を選択するということはすなわち NVIDIA と命運をともにすることになる、ということくらいは肝に銘じておく必要があります。

*1:PTX (Parallel Thread Execution) はNVIDIA独自の中間命令体系で、CUDAでも使われています。

*2:読み書きをする場合は Image でなく Buffer を使わなければならず、つまり読み取り時に補間サンプリングをしたい場合は、ハードウェアに例え高速なサンプラーが用意されていても、自前でバイリニア フィルタリングなどのコードをわざわざ書かないといけなくなる。ピンポン処理で回避するにしても無駄にメモリを食う羽目になる。

*3:OpenCLはGPGPU/HPC分野においては頼りないのですが、AlteraやXilinxのFPGAなどに代表されるように、組み込み分野では徐々に活用が進んでいるようです。