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

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

NVIDIA/AMDのDirectX 11.x/12対応に関して

Direct3D DirectX 技術・テクノロジー GPU

NVIDIAの場合

GeForce GTX 970のメモリ3.5GB問題で、嘘スペック（なんちゃって4GB他）を公表していたがために訴訟を起こされてしまったNVIDIAですが、実はDirectX (Direct3D) 対応レベルに関してもわりとあこぎなことをやっているというか実際詐欺です。

Fermi/Kepler/Maxwell世代でDirectX 12に対応するとの報道がありますが、少なくともFermi/Keplerに関してはGTX TITAN Zですら機能レベル (Feature Level) 11_0にとどまっており、「DirectX 11.1対応」や「DirectX 11.2対応」を謳っている製品でも実際にフルサポートしているのはDirectX 11.0までの機能となっています。

ASUSやEVGAはきちんとFeature Levelに言及していて正直なんですが、NVIDIA Japan公式サイトの780/750/Titan Blackはアウトですね。980のDirectX対応表記も「12 API」となっていますが、肝心の機能レベルへの言及がありません。ELSAは「Shader Model 5.0」と表記しており、たぶんDirectX 11.0までの対応だと言いたかったんだと思いますが、実はShader Model 5.1仕様に関してMicrosoftからの公式情報はありません（MSDNのDirect3D APIヘルプには一切書かれていないし、Windows SDK 8.0/8.1付属のシェーダーコンパイラーfxc.exeにもvs_5_1やps_5_1のようなプロファイルは存在しない）。DirectX 11.1/11.2であってもShader Model自体は5.0のままであり、ELSAが想定しているであろうShader Model 5.1というのは幻です。

なんでこんなわけの分からないことになっているのかというと、Fermi/KeplerはドライバーレベルでDirectX 11.1やDirectX 11.2のAPI自体はサポートしていて、またDirectX 11.1の機能のうち実際にいくつかはサポートしているんですが、すべてではないため、Direct3D 11 APIを使ってデバイスの機能レベルを取得すると11_1 (D3D_FEATURE_LEVEL_11_1) ではなく11_0 (D3D_FEATURE_LEVEL_11_0) となります。これに関しては実際に手元にあるKepler世代のGTX 760や770で試したことがあり、実証済みです。

DirectX 11.1のCapsは下記構造体です。

DirectX 11.2のCapsは下記構造体です。末尾の数字が2じゃなくて1なのがまぎらわしくてウザい感じです。

ちなみにD3D_FEATURE_LEVEL_11_2はありませんのであしからず。なおDirect3D 11.1/11.2の各機能対応状況は、Windows SDK 8.1付属のDirectX Caps Viewer (dxcapsviewer.exe) で確認することもできます。またGPU-Zの最新版を使うことで、DirectXのフルサポートバージョンを簡単に調べることができます。

DirectX 10世代の古いカード（G80）とかでも最新版の対応ドライバーをインストールすることで、Windows Vista SP2 Platform Update/Windows 7以降であればDirectX 11.0 API自体はFeature Level 10_0で使えますよね？　それと一緒です。下記の公式ブログにおける「all Fermi, Kepler, and Maxwell GPUs will fully support the DX12 API」という表現も、そういう意味を含めたうえでの発言でしょう（おそらくDirectX 12によるマルチスレッドレンダリングやコンピュートタスクの効率向上はFermi/Keplerでも享受できるはず）。

DirectX 12およびDirectX 11.3に関しては先行プログラムに参加しているベンダーや開発者のみが詳細を知っている状態で、まだ一般公開されていないステータスであるためはっきりとしたことは言えませんが、GTX 900シリーズ（第2世代Maxwell）に関しては少なくともDirectX 12／DirectX 11.3の新機能のうちいくつかはサポートしているよ、とNVIDIAの中の人がフォーラムで言及しているようです。機能レベルがどうなるのかは不明です。

2015-03-23追記：
GDC 2015で発表された情報によると、DirectX 12では12_0および12_1という機能レベルが実装されるそうです（スライドではFeature Level 12.0および12.1と表記されています）。

個人的に特に残念なのは、ROV (Rasterizer Ordered View) が確実に使えるのは12_1のほうで、12_0以下では使えない（オプション扱い）という点でしょうか。早くも格差が生まれてしまったという悲壮感が漂います。

ちなみに英語版Wikipediaの「Direct3D」の記事には、「AMD GCN 1.1/1.2では機能レベル12_0となる」「NVIDIA GeForce 900シリーズ（TITAN Xも含む）では機能レベル12_1となる」と2015年3月時点で記載されているんですが、DirectX 11.3/DirectX 12およびWindows 10の正式版はまだリリースされておらず、MSDNライブラリの正式ドキュメントすら準備されていない現状、機能レベルに関する引用はChannel 9のビデオだけで、かなりフライングしています。
マイクロソフトおよびハードウェアベンダー（メーカー）からの最終的な正式情報（特に文字情報の検索が可能なドキュメント）を待つべきだと思うんですが……

2015-04-08追記：
Windows 10 Technical Preview上でVisual Studio 2015 Preview付属のWindows SDK 10を用いることで、一般の開発者でもDirect3D 11.3およびDirect3D 12を使えるようになっているようで、各々の暫定ドキュメントがMSDNライブラリにて公開されています。ただし実行にはWDDM 2.0対応ドライバーを導入する必要があります。最終的な仕様はまた変更される可能性はありますが、大枠はすでに固まっているものと思われます。

DirectX 11.3のCapsは下記構造体です。

DirectX 12のCapsは下記構造体です。

今回もD3D_FEATURE_LEVEL列挙体のメンバーに変動はなさそうです。

2015-05-12追記：
2015年4月末にリリースされたVisual Studio 2015 RC付属のWindows SDK 10では、D3D_FEATURE_LEVEL_12_0およびD3D_FEATURE_LEVEL_12_1が追加されているようです。また、同SDK付属のシェーダーコンパイラーfxc.exeには、幻と思われていたシェーダーモデル5.1のプロファイル（cs_5_1, ds_5_1, gs_5_1, hs_5_1, ps_5_1, vs_5_1）が実装されているようです。シェーダーモデル5.1自体はDirect3D 12 APIをドライバーレベルでサポートするハードウェアであれば、D3D12/D3D11 APIを通して使用可能らしいですが、5.1で新しく追加されたROVに関するオブジェクト（RasterizerOrderedBufferほか）は、ROVに対応するハードウェアのみで使用可能となるはずです。

AMDの場合

AMDに関しては、GCNでFeature Level 11_1をフルサポートするという話だったんですが、結局どうなっているんでしょうか？

また、Radeon HD 7000シリーズにおけるDirectX 11.2に関しても、ドライバーレベルで対応可能という話だったんですが、これも実際のところどうなっているんでしょうか？

ASUSによると「DirectX 11.1対応」「DirectX 11.2対応」とかいいつつ実はFeature Level 11_0どまりとなっています。ASUS JAPANのサイトが更新をサボっているだけ？

SAPPHIREのサイトでは「GCNベース製品はすべてDirectX 12対応予定」と書かれているだけで、Feature Levelへの言及はありません。
NVIDIA同様、最終的にDirectX 12 API自体の対応がなされるだけで、旧製品に関してはFeature Level 12_0/12_1ではなくなるのでは？　新型Radeonに搭載される予定のFijiあたりは12_1にフル対応してきそうですが……

結論

問題は何をもって「11.1対応」「11.2対応」「12対応」と称するべきか、なんですが、個人的にはまともにフル機能をサポートしていない中途半端な製品なのに「DirectX ○○対応!!」とかいう詐欺商法は即刻やめて欲しいです。エンドユーザーがガッカリするのはもちろんのこと、開発者もCaps絡みで苦しむだけです。個別の機能が使えるかどうかいちいちフラグをチェックする必要があるなんていう三流仕様はOpenGLだけにしてください。ていうかDirectX機能レベルというものの存在に関して、ベンダーはもっと説明責任を果たすべきでしょう。かつてDirect3D 9全盛期の頃、グラフィックスカードのベンダー各社はDirectX 9.0c対応を謳いつつも、浮動小数バッファのアンチエイリアスはATIだけがサポート、VTFはNVIDIAだけがサポート、というようなひどい状況になっていたことがありましたが、そういった細かい機能のサポート有無は非常に重要な情報であるにもかかわらずスペックシートには一切書かれていませんでした。

2015-06-03追記：
現行のGeForceシリーズにおけるDirectX機能レベル対応度に関しては、NVIDIAからGeForce GTX 980 Tiとともに公式の解説が発表されました。Maxwell第2世代以降でFeature Level 12_0と12_1にフル対応するようです。当該ハードウェアを所持していて、Windows 10 Preview BuildとWDDM 2.0対応ドライバーをインストールできる人は実際に試してみるとよいでしょう。

AMDも新型Radeonシリーズの発表の際には、GCNアーキテクチャの世代とDirectX機能レベルに関して言及してくるものと思われます。

※2015-06-21追記：
The Era of PC Gamingで発表されたように、新しいRadeonはR9 Furyシリーズ（R9 Fury X、R9 Fury、R9 Nano）のみがFiji（GCN 1.2? 1.3?）を搭載するようです。R9 380/390/390Xのほうはすべてリネーム（リブランド）で、HawaiiやTongaなど（GCN 1.1/1.2）を採用するものと思われます。機能レベルに関する言及はありません。

科学技術英文を翻訳するコツ

技術・テクノロジー 英語

仕事柄、ネイティブ英語圏の開発者と英文メールをやりとりすることがあるんですが、ときにはローカライズ作業などで英文を日本語に翻訳しなければならないこともあります。もちろんその逆もあります。

自分は技術者なので、先方とやりとりする英文の文法自体は特段難しいものではなく、専門用語や技術用語（テクニカルターム）が多数出現することを除けばむしろヘタなビジネス英語よりも易しいくらいですが、技術文書には英語・日本語を問わず特有のクセや言い回しがあります。
日本語として自然な文章になるよう翻訳したいときは、下記のようなサイトが参考になると思います。読み物としても面白く、特に翻訳の際に受動態・能動態をどう使い分けるか、また無生物主語の扱いをどうするかという話は必見です*1。

翻訳の泉 - 教室

逆にいえば、自然な英文を書く際にも、こういった日本語への翻訳に関する知識は役に立ちます。受験英語から一歩進んだWritingのテクニックを勉強しておくことは大いに意義があるでしょう。英語圏のネイティブ向けに、日本人が英語でマニュアルや仕様書を書かなければならない場合、下記のような資料も参考になります。

ドキュメンテーション（文書化）の基礎知識【講議用テキスト】株式会社 テクライト・ジャパン ドキュメンテーションチーム

*1:日本語ではもともと主語を省略することが多いのですが、そのせいで日本人にとっては無生物主語という存在が奇異に映るのかもしれません。

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などに代表されるように、組み込み分野では徐々に活用が進んでいるようです。

F#でファイル・フォルダーをごみ箱へ移動

.NET Framework F# プログラミングTips Visual Studio

Visual Studioのバージョン2010以降をインストールすると、拡張子.fsxにVisual Studioが関連付けられ*1、さらにコンテキストメニューに「Run with F# interactive...」というコマンドが追加されます。.fsxはF#スクリプトと呼ばれるファイルなのですが、Visual Studioコードエディターを使ってF#コードを（インテリセンス付きで！）書けるようになり、さらにF# interatciveを使ってF#コードをスクリプトとして実行できる環境ができあがるわけです。

今回はF#を使って指定のファイル・フォルダーをまとめてごみ箱に移動する簡単なスクリプトを書いてみました。F# 3.xで動作確認しています。こういったOSのシェル機能を使ったタスクはWindows PowerShellの領域だったり、というかむしろC#で書いてしまったほうが返っててっとり早かったりするのですが、今後の将来性が見込まれるマルチパラダイム関数型言語F#の習得を今年から少しでも進めてみようということであえてF#を使っています。
ちなみに現状のfsxのインテリセンスはIronPython (Python Tools for Visual Studio) よりは使いやすいけれど、Visual C#には及ばない、といったレベルでしょうか。

// http://msdn.microsoft.com/ja-jp/library/vstudio/dd233175.aspx
#if INTERACTIVE
#r @"Microsoft.VisualBasic.dll"
#endif

//#light "off"

open System
open Microsoft.VisualBasic.FileIO

printfn "バックアップに不要なファイルおよびビルド中間生成物をごみ箱に移動します。"

let moveFileFolder fileFolderName moveFunc =
    try
        try
            // HACK: ファイルなのかフォルダーなのか判断して、呼び出す関数を分けるような分岐を実装してもよい。
            printfn "Now moving : <%s>" fileFolderName
            moveFunc(fileFolderName, UIOption.AllDialogs, RecycleOption.SendToRecycleBin)
        with
            | :? OperationCanceledException as ex -> printfn "移動がキャンセルされました。Message=\"%s\"" ex.Message
            | ex -> printfn "Message=\"%s\"" ex.Message
    finally
        ignore()

let moveFile fileName = moveFileFolder fileName FileSystem.DeleteFile
let moveFolder fileName = moveFileFolder fileName FileSystem.DeleteDirectory

let moveBinObjFolders baseDirName =
    moveFolder (baseDirName + @"\bin")
    moveFolder (baseDirName + @"\obj")

moveFile ".\\test.txt"
moveFile @"..\HogeBuild.v11.suo"
moveFolder @"..\.svn"
moveBinObjFolders @"..\FugaLib"

printf "Press any key to exit..."
Console.ReadKey(true) |> ignore

なお日本語表示など、Unicode対応する場合はfsxファイルをUTF-8で保存します。

はてなブログのシンタックスハイライトは最近Swift言語に対応したらしいんですが、Microsoft発の関数型言語F#にも早く対応して欲しいです。とりあえず対応してくれるまではOCamlコードとしてハイライトすることにします……

2017-11-26追記：
2016年にF#やPowerShellにも対応してくれたようです。素晴らしいですね。
ソースコードを色付けして表示する（シンタックスハイライト） - はてなブログ ヘルプ

*1:Windows 8.xではVisual Studioをインストールしてもfsxに関連付けられないことがあるようです。その場合はエクスプローラーなどから関連付けを手動設定してやる必要があります。

先輩、今年もよろしくお願いしますね

アニメ ゲーム イラスト


【Fate】「先輩、今年もよろしくお願いしますね」イラスト/sygh[sai] [pixiv]

年末に体調を崩してなかなか年賀絵を描けなかったんですが、ずっと妄想していた和装の桜を描くことができてひとまず肩の荷が降りました。
本年もぼちぼちコードを書いたりイラストや漫画を描いたりしていく所存ですのでよろしくお願いいたします。

和服の花柄は出来合いのテクスチャを使うのが嫌だったので、Illustratorで自作しました。美遊のサファイアちゃんも去年Illustratorで線画を作成してみましたが、ベクトル画像は拡大縮小や回転にいくらでも耐えてくれるので、プログラムのように部品化しやすいのがメリットです。


【プリズマ☆イリヤ】「黒魔法少女」イラスト/sygh[sai] [pixiv]

CUDA Warpシャッフル命令のエミュレーション

C++ CUDA GPU

今更ですがせっかくCompute Capability 3.0対応のKepler世代グラフィックスカードを手に入れたので、CUDAのWarpシャッフル命令の動作テストを兼ねて、代替機能をエミュレートする関数を書いてみました。
Visual Studio 2012、CUDA 6.5、GeForce GTX 770で動作確認済み。

#define WARP_SIZE (32)

typedef unsigned int uint;

// HLSL の SV_GroupIndex 相当。
static __device__ __forceinline__ int GetLocalThreadGroupFlatIndex()
{
  return threadIdx.x + blockDim.x * threadIdx.y + blockDim.x * blockDim.y * threadIdx.z;
}

static __device__ __forceinline__ int GetGlobalThreadFlatIndex()
{
  const int blockId = blockIdx.x + gridDim.x * blockIdx.y + gridDim.x * gridDim.y * blockIdx.z;
  return blockId * blockDim.x * blockDim.y * blockDim.z + GetLocalThreadGroupFlatIndex();
}

#pragma region // Emulation Functions //

// 書き込み後の同期だけでなく、読み込み後の同期も必要。でないと安全に共有メモリ上の古いデータを上書き／破棄できない。

template<typename T> static __device__ T EmulateShuffle(T x, int srcLane)
{
  extern __shared__ T sharedMem[];
  const int localIndex = GetLocalThreadGroupFlatIndex();
  const int laneBlockId = localIndex / WARP_SIZE;
  const int laneId = srcLane & (WARP_SIZE - 1);
  sharedMem[localIndex] = x;
  __syncthreads();
  const T y = sharedMem[laneBlockId * WARP_SIZE + laneId];
  __syncthreads();
  return y;
}

template<typename T> static __device__ T EmulateShuffleUp(T x, uint delta)
{
  extern __shared__ T sharedMem[];
  const int localIndex = GetLocalThreadGroupFlatIndex();
  const int laneBlockId = localIndex / WARP_SIZE;
  const int laneId = localIndex & (WARP_SIZE - 1);
  const int shiftAmt = laneId - static_cast<int>(delta);
  sharedMem[localIndex] = x;
  __syncthreads();
  const T y = (0 <= shiftAmt) ? sharedMem[laneBlockId * WARP_SIZE + shiftAmt] : sharedMem[localIndex];
  __syncthreads();
  return y;
}

template<typename T> static __device__ T EmulateShuffleDown(T x, uint delta)
{
  extern __shared__ T sharedMem[];
  const int localIndex = GetLocalThreadGroupFlatIndex();
  const int laneBlockId = localIndex / WARP_SIZE;
  const int laneId = localIndex & (WARP_SIZE - 1);
  const int shiftAmt = laneId + static_cast<int>(delta);
  sharedMem[localIndex] = x;
  __syncthreads();
  const T y = (shiftAmt <= (WARP_SIZE - 1)) ? sharedMem[laneBlockId * WARP_SIZE + shiftAmt] : sharedMem[localIndex];
  __syncthreads();
  return y;
}

template<typename T> static __device__ T EmulateShuffleXor(T x, int laneMask)
{
  extern __shared__ T sharedMem[];
  const int localIndex = GetLocalThreadGroupFlatIndex();
  const int laneBlockId = localIndex / WARP_SIZE;
  const int laneId = localIndex & (WARP_SIZE - 1);
  sharedMem[localIndex] = x;
  __syncthreads();
  const T y = sharedMem[laneBlockId * WARP_SIZE + ((laneId ^ laneMask) & (WARP_SIZE - 1))];
  __syncthreads();
  return y;
}

#pragma endregion

エミュレート関数はそれぞれ、下記のCUDA組み込み関数に相当します。当然組み込み関数を使ったほうが手軽だし高速ですが、CC 3.0以降のデバイスでないと使えません。

// 指定したレーン (Warp 内のスレッド番号) の var の値を受け取ります。
T __shfl(T var, int srcLane, int width = warpSize);

// 指定した数 delta だけ離れたスレッドの var の値を受け取ります。
// up は指定した数だけ上の番号のスレッドから、down は下の番号のスレッドから受け取ります。
// Warp の外に出てしまった場合は自身のスレッドでの値が返ってきます。
T __shfl_up(T var, unsigned int delta, int width = warpSize);
T __shfl_down(T var, unsigned int delta, int width = warpSize);

// 自分のレーン番号と指定した laneMask を XOR した結果の番号のスレッドから受け取ります。
// laneMask ごとにレジスタの中身を交換することができます。
T __shfl_xor(T var, int laneMask, int width = warpSize);

CUDAの場合、共有メモリのサイズ指定はコンパイル時だけでなく、下記のように（ハードウェア仕様上限が許す限り）実行時に指定することもできます。今回のエミュレーション関数は共有メモリにexternを指定することで、実行時にサイズ指定する方法を選びました。

const dim3 gridSize(16, 4, 4);
const dim3 blockSize(64, 8, 2);
const size_t elemCount = gridSize.x * gridSize.y * gridSize.z * blockSize.x * blockSize.y * blockSize.z;
const size_t sharedMemSize = blockSize.x * blockSize.y * blockSize.z * sizeof(TargetType);
…
kernelFunc<<<gridSize, blockSize, sharedMemSize>>>(devOutArray, devInArray);

一応動作確認テストはしましたが、完全に同一仕様でエミュレートできているかどうかは未保証です。ご利用の際は自己責任でどうぞ。

CUDAはNVIDIA専用という最大のネックがあるので、個人的にDirectComputeに比べてあまり好きではないのですが、CUDA自体は非常に効率よく最高水準のGPGPUプログラムを記述できる、地球上でもっとも優れた超並列プログラミング環境であることは確かです。C++erとして特に気に入っているのは、Visual C++などのホストコンパイラーにコンパイルさせるホスト関数だけでなく、NVCCコンパイラーにコンパイルさせるデバイス関数でもクラス、テンプレートやC++11 Featureがほとんどそのまま使える点です（CUDA 7ではC++11対応が強化されるそうです）。HLSLでもクラスは使えますが、CUDAのプログラマビリティはHLSLを遥かに凌駕します。GLSLコンピュートシェーダーやOpenCL-Cごときではまったく勝負にならないレベルです。
なおネイティブC++から使えるGPGPU言語拡張としては、Visual C++ 2012以降のC++ AMPも強力です。WindowsやXbox Oneのほか、すでにLinux向けにもオープンソース実装されているので、OpenCLの代わりにC++ AMPを使ってみてもよいでしょう。

ちなみに自分はPC用グラフィックスカードに関しては、G7xからG8x、Fermi、Keplerと乗り換えてきましたが、ずっとNVIDIAオーナーです。AMDはハードウェアに関しては良いモノを作っていることは確かなんですが（ゲーム向けも業務向けもHPC向けも性能に関してはAMDのほうが上）、ドライバーの完成度がお話にならないレベルであったり、ソフトウェアを軽視しているメーカーだと感じます。また、APUを除き、現行のグラフィックスカード単体での省電力性能は、ハイエンドでもNVIDIAに分があります。Keplerを使い倒したら、いずれMaxwellへ、さらにその後はPascalへ移行する予定ですが、AMDが今の劣悪な状況を改善しようとしないかぎり、今後も個人でAMDグラフィックスを購入することはないでしょう（逆に言うとドライバー／ソフトウェア面さえ改善されれば、オープン標準規格のOpenCLやC++ AMPを推進するAMDはNVIDIAよりも良い選択肢になり得ます）。

VS 2008のWPF-MFC相互運用

C++ Visual Studio WPF プログラミングTips MFC

（これは2011-03-18に書いた故OCNブログの記事を移植したものです）

以下を参考に、MFC アプリから WPF を使おうとして、Visual Studio 2008 SP1 のバグに遭遇しました。
WPF ユーザー コントロールを HwndSource 経由でホストする Win32 アプリ……といいつつ、結局 C++/CLI による中継が必要になりますのでご注意。
なお、VS 2010 では C++/CLI のインテリセンスが効かないという致命的なダウングレードがあるため、MFC-WPF の相互運用をするならば、VS 2008/2012/2013 を使うことをお勧めします。

バグ現象

スクリーンショット (a) は、WPF ユーザー コントロール プロジェクトと、それを使う予定の MFC クライアント アプリ プロジェクトを格納したソリューション。

(a)
https://sygh-jp.github.io/content_hosting/my_program_ss/Wpf_Designer_Error_ss_2010_05_31a.png

一見何の変哲も無いが、ここで C# コードのタブ (.xaml.cs) へ移動し、WPF のプロジェクトをリビルドする。そして、XAML デザイナのタブ (.xaml) へ再度移動する。
すると、スクリーンショット (b) のように再読み込みを促す情報バーが出るので、クリックすると、スクリーンショット (c) のように IDE がクラッシュする。

(b)
https://sygh-jp.github.io/content_hosting/my_program_ss/Wpf_Designer_Error_ss_2010_05_31b.png

(c)
https://sygh-jp.github.io/content_hosting/my_program_ss/Wpf_Designer_Error_ss_2010_05_31c.png

再度ソリューションを開くと、スクリーンショット (d) のエラーが待っている。

「オブジェクト参照がオブジェクト インスタンスに設定されていません。」

なん……だと……

いわゆる NullReferenceException のハンドルされない例外が発生している。よくあるアプリケーションが落ちる原因のひとつだが、こいつが WPF デザイナー（Visual Studio IDE プロセス）内部で発生している。

(d)
https://sygh-jp.github.io/content_hosting/my_program_ss/Wpf_Designer_Error_ss_2010_05_31d.png

なんぞこれ……ということで調べてみたところ、MS のサポートページにひっそりと修正プログラムが公開されていた。

上記によると、Win32 プロジェクトがスタートアップ プロジェクトとして設定されているとバグるらしい（ソリューション エクスプローラーにおいてプロジェクト名がボールド体になっているものがスタートアップ）。
確かにスクリーンショット (e) のように、WPF のほうをスタートアップ プロジェクトに設定して、リビルドすると先ほどのエラーは発生しない。

(e)
https://sygh-jp.github.io/content_hosting/my_program_ss/Wpf_Designer_Error_ss_2010_05_31e.png

ともかく、修正プログラム 963035 を適用せよ、とのこと。下記へ飛ぶ。

修正プログラムは下記のサイトにアップロードされている、とのこと。なんというたらいまわし。

このサイトの Downloads タブに目的の修正プログラム「VS90SP1-KB963035-x86.exe」のダウンロード リンクがある。
プログラム適用によって問題が一応修正されていることを確認。
開発者にとっては非常に重要な修正プログラムですが、Windows Update 経由では適用されません。また、2014年現在、修正プログラムへのリンクが切れているようです。代替の修正プログラムが公開されているのかどうかは不明。

しかし Win32 アプリをスタートアップにするとバグるっていう理由が意味不明です。WPF との相互運用を図るとき、普通にやる組み合わせだと思うんですが……最低限そのくらいはテストしてからリリースして欲しいです。

余談：Win32 アプリから WPF を使う3つの方法

  1. /clr を使って C++/CLI でホストコードを書く。
  2. /clr を使って C++/CLI でホストコードを書く。ただし C++/CLI のコードは DLL に押し込む。
  3. WPF アセンブリを COM 公開する。

1番目は自明なので説明は不要ですね。

2番目の方法を具体的に説明すると、/clr を有効にした Win32 DLL、MFC 拡張 DLL もしくは MFC 標準 DLL 内で WPF を使用して、ネイティブ MFC インターフェイスもしくは C 言語形式関数を公開する形にしておけば OK。例えば、ネイティブの HWND ウィンドウ ハンドルを受け取って WPF コントロールのホストとして使う C 言語形式関数を C++/CLI で実装して、WPF コントロールのインスタンスを DLL 内で gcnew するようにラップしてしまえば、ネイティブ C++ で書かれた EXE（厳密に言うと、/clr を無効にした EXE）から普通に呼び出せます。ただし /clr を使う場合は CRT/MFC/ATL を静的リンクではなく動的リンクする必要があるので注意しましょう。

1番目および2番目の方法において、もし混合クラスのメンバーとしてマネージ オブジェクトを保持したい場合は、ラッパークラス テンプレートの gcroot を使います。

ちなみに C++/CLI だと、ネイティブ／マネージのうちどちらか片方のインターフェイスだけでなく、両方のインターフェイスを合わせ持つ混合アセンブリ（COM でいうデュアル I/F に近い）まで作れてしまいます。つまり、ネイティブ Win32/MFC 向けの関数／クラスと、C# or VB.NET 向けのクラス両方をひとつのアセンブリに含めるという変態的なことができてしまう。

C++/CLI を一切使いたくない or 使えない場合、3番目の方法として、WPF アセンブリを COM コンポーネントとして公開する方法があります。COM 公開するには、regasm でレジストリ登録する方法のほか、マニフェストで Side-by-Side アセンブリにする方法があるようです。

ATL で開発した従来の COM コンポーネント DLL に関しては、マニフェストで Side-by-Side アセンブリ化（プライベート アセンブリ化）し、実運用にも十分耐えうることを確認したことがあるんですが、WPF アセンブリに関しては未検証です。

もし純粋なネイティブ Win32 アプリ（従来のネイティブ MFC や VB 6.0 などで開発したアプリ）から直接 WPF アセンブリを使いたい場合は、2番目もしくは3番目の方法を使う必要があります。ただし、境界面のやりとりに手間や労力がかかることは覚悟しておく必要があります。