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

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

CUDA

VulkanシェーダーでSub-group命令を使う

NVIDIAはKepler (Compute Capability 3.0) 世代のハードウェアにおいて、Warp Shuffle命令を実装しました。WarpシャッフルはCUDAから組み込み関数 (intrinsic function) の形で利用できるSIMD命令の一種で、Warpと呼ばれるスレッドグループ内での並列データ…

GPGPU戦争の歴史を紐解く ―勃発から現在に至るまで―

GDC 2015でついにOpenCL 2.1が発表されました。SPIR-V中間表現(中間言語、バイトコード)でプログラミング基盤をVulkan API (OpenGL Next Generation, glNext) と共有することができるそうです。これでカーネルをオフラインコンパイルできるようになります…

NVIDIAのOpenCL対応状況

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まで)。例え…

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

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

CUDAでテクスチャ

前回の記事では CUDA の線形メモリ(Linear Memory, cudaMalloc() で生成)を使って FP16 浮動小数点数値を格納しました。 今回は CUDA の Array と Texture を使って、0x0000 ~ 0xFFFF の範囲の 16bit 整数値(ushort)を正規化して 0.0 ~ 1.0 の範囲の 32b…

CUDAでhalf型

VRAM 使用量や帯域の節約目的で、CUDA や OpenCL のカーネルに FP16 半精度浮動小数点数(half)型のデータを渡す場合の話です。 通例 GPGPU で使われる実数は FP32 単精度浮動小数点数(float)型なんですが、half だとその半分のデータ量で済むため、帯域…