OpenDLSS: NvidiaのDLSS 5ニューラルレンダリングネットワークのVulkan再実装

OpenDLSS-NR NVIDIAのDLSS 5ニューラルレンダリングネットワークのVulkan再実装、オリジナルとビット単位で一致。 同じ71ブロックのスウィン/ViTネットワークをDLSS-NRビルド310.8.0として、テンソルコア上でFP8で実行します。中間結果も一致しており、最終画像だけでなく、すべての75ブロック境界でバイト単位で一致します。 ports/browser-webgpu/ は2番目の独立した実装です。同じバイトをブラウザで、テンソルコアやFP8なしで実行します。 モデルディレクトリを以下のレイアウトで指定することで、重みを提供します。 ネットワーク グローバルViTを下部に持つシフトウィンドウ・トランスブロックスのU-Net：6つのプーリングレベルにわたる71ブロック、FP8（E4M3）アクティベーションとFP16累積、141 MiBの重み。 これは生成的なニューラルレンダリングネットワーク（NVIDIAの用語）です。エンジンが既に描画したフレームを再レンダリングし、注入されたノイズからディテールを生成し、スタイル設定の下でトーン、構造、肌を調整します。入力と出力は同じ解像度であり、アップスケーラーではありません。 WebGPUポートは2048x1152で、左がNRオフ、右がNRオン。シーン：Cowboy Gramps by Muhammed Ismayil, CC0。 1つのレンダリング済みフレーム（低ダイナミックレンジのプロキシ）、3レーンのガウシアンノイズ、前のフレームの出力の再投影、および5つの条件付けスカラーを入力として受け取り、ピクセルあたり4つのf32チャンネル（RGB残差と1つの時間的ブレンドロジット）を生成します。docs/network.md にグラフ全体が記載されています。 NVIDIAは、そのレポート「DLSS 5: Generative Neural Rendering」（プロジェクトページ）でモデルを説明しています。 ビルドと実行 powershell -File scripts\fetch_tools.ps1 [-Npm] # 一度だけ: tools\ (glslang, Vulkan-Headers, volk, CMake, Ninja) powershell -File scripts\build.ps1 # シェーダー、PTX、build\dlss5vk.exe powershell -File scripts\fetch_filament.ps1 # 一度だけ、デモ用: third_party\filament (+ パッチ) powershell -File scripts\build_filament.ps1 # 一度だけ、デモ用: third_party\filament-install powershell -File scripts\build_demo.ps1 # build\demo\dlss5-demo.exe build\dlss5vk.exe bench --model <dir> --width 768 --height 768 build\dlss5vk.exe profile --model <dir> --width 768 --height 768 # ディスパッチごとのタイミング build\dlss5vk.exe parity --model <dir> --fixture <dir> # フィクスチャに対するビット単位の一致 build\dlss5vk.exe verify --model <dir> --fixture <dir> # ブロック0カーネルごとの確認 python scripts\ptx\test_fast_divmod.py # PTXディバイダー、n < 2^24 の全範囲で (numpy) デモはダブルクリックで実行できます。デモのシーンドロップダウンにbuild\scenes下のすべてのシーンが表示され、最初のシーンから開始するか、コマンドラインで指定されたglTFをロードします。モデルディレクトリは --model <dir>、指定がない場合はDLSS5VK_MODEL、それもなければこのREADMEの隣のmodels\nr です。レンダラー、キー、シーン、view.jsonについては demo/README.md を参照してください。 パフォーマンス RTX 4070 SUPER、フレームあたりのネットワーク全体、40フレームの最小値。 すべての解像度で241ディスパッチ。 解像度 | 時間 -------|------- 768x768 | 2.8 ms 1920x1080 | 7.8 ms 2560x1440 | 12.6 ms 3840x2160 | 29.3 ms GPUは持続的な負荷の下で2つのクロック状態を交互に使用するため、中央値は数パーセント高くなります。最小値と比較してください。 内容 ファイル | 備考 -------|------- src/ (C++20) | Vulkanコンテキスト、モデルローディングと重み再配置、カーネルラッパー、ネットワークグラフ、算術演算のCPUリファレンス、dlss5vkツール GLSLカーネル | shaders/ 参照ルート：cooperative-matrix FP8 GEMMs、フューズド32チャンネルブロック、フューズドQKV + ウィンドウアテンション、エキスパートMLP、グローバルアテンション、要素ごとの演算。それ自体で完全かつ正確。 PTXカーネル | scripts/ptx/ 高速ルート用のPTXを生成するPythonジェネレーター：mma.sync E4M3（f16累積付き）、cp.asyncリング、デバイスカウンターによるバリアフリーチェイニング、スプリットK GEMMs、ストリーミンググローバルアテンション。build/ptxにビルド時に生成されます。 デモ | demo/, third_party/filament.patch Filament（Apache-2.0）フレーム内のネットワーク：オブジェクトごとのモーションベクトルとVulkanインターロップフック用にパッチされたFilament、gltfio経由のglTFシーン、ImGuiコントロール。 WebGPUポート | ports/browser-webgpu/ ブラウザでの同じネットワーク、同じキャプチャとのビット単位の一致、テンソルコア、FP8、ブロック間のフュージョン、チェイニングなし：正確さは仕様にあり、ハードウェアにはありません。512x512で72ms、ここでは2.7ms。 実装されていないもの：DLSS-SR（別のネットワークです）。 時間的パスは実装されていますが、デモ内です：ネットワークの履歴入力レーンとそのピクセルごとのブレンドロジットが再投影されたフィードバックループを駆動します（docs/frame.md）。dlss5vkツールは履歴なしで単一フレームを実行します。これは参照キャプチャが作成された方法です。 要件 Windows、NVIDIA Ada（またはそれ以降）GPU、およびVK_KHR_cooperative_matrix、VK_NV_cooperative_matrix2、VK_EXT_shader_float8、VK_NV_cuda_kernel_launchを公開するドライバー。 C++ x64ツールセット（任意のバージョンまたはBuild Tools、vswhere経由で検索、またはVCVAR64をvcvars64.batに設定）を備えたVisual Studio 2022以降、git、PTXジェネレーター用のPython 3、およびシーンコンバーター用のNode.js + npmとPillow。 tools/下のポータブルツールチェーン（gitで無視）：scripts\fetch_tools.ps1はglslang 16.6.0、Vulkan-Headers v1.4.363、volk（タグ固定）、CMake 3.31、Ninja 1.13をダウンロードします。Vulkan SDKのインストールは不要です。 -Npmはシーンコンバーターのモジュールをtools\gltfにインストールします。 デモ用：Filament v1.77.0。scripts\fetch_filament.ps1によってクローンおよびパッチされ、scripts\build_filament.ps1（両方ともgitで無視、約15分と6GBのビルドツリー、%LOCALAPPDATA%\dlss-vulkanまたはDLSS5_FILAMENT_BUILD_DIRに配置；DLSS5_BUILD_JOBSはMSVCがジョブあたり最大1GBを消費するため、Filamentの並列コンパイルを制限します、デフォルト8）。 モデルディレクトリ nr::Modelはmanifest.jsonを読み取ります：stages配列（各エントリ：id、ディレクトリからの相対パス、packedByteLength、sha256）とtensors配列（各エントリ：name、block、layer、parameter、stage、stageOffset、byteLength）。ステージファイルにはE4M3重みがパックされたバイトとして格納されます。ホストはそれらをカーネルが消費する行列形式に再配置します（src/nr_model.cpp）。このリポジトリには、そのようなディレクトリを生成するものは何も含まれていません。 グラフは310.8.0の71ブロックネットワークのみであり、それ以外はありません。ブロック数が異なるモデルはロード時に拒否されます。 フィクスチャ parityはオリジナルの記録されたキャプチャと比較します。これらはリポジトリに含まれていません。フィクスチャはmanifest.jsonを持つディレクトリです： key | sourceDimensions, fullDimensions | 有効なサイズとパディングされたフィールド proxy or inputFeatures | 入力：RGBA f32画像（条件付け、シード、autoMask付き）、またはf32特徴量 checks | フィクスチャがチェックするゲート、"boundaries"、"head"、"output"のいずれか。必須で空ではない。 blocks, transitions | 比較可能な境界（ブロック0-69、5つのエンコーダー遷移） "boundaries": E4M3参照（ブロック/id、幅、高さ、チャンネル、ファイル） omittedBoundaries | フィクスチャが参照を持たない各比較可能な境界の理由 referenceHead | "head": f32 RGBAヘッド nativeOutput | "output": 合成画像、dtype "f32"（RGBAハーフ、プロキシが必要）または "u8"（8ビットキャプチャ）。 GPUが実行する前にすべてが検証され、いずれかの検証に失敗したフィクスチャは拒否されます：宣言されたチェックに参照がない、参照が欠落している、短い、またはグラフ内の何も参照していない、宣言されたチェックが使用しない参照、または参照も理由もない比較可能な境界（ブロック0-69、5つのエンコーダー遷移）。 判定はビット単位の一致（合格）、ゼロの符号までの等価性（失敗）、1コード以内（8ビットキャプチャのみ、別途報告）、または不一致です。docs/numerics.mdを参照してください。 ヘッドと出力はプロダクションスケジュールで比較され、--repeat回数（デフォルト3）繰り返され、バリア下で同じグラフと一致する必要があります。境界は計測された実行から取得され、そのヘッドはプロダクションのヘッドと一致する必要があります。 verifyはさらにinputFeaturesとブロック0の参照が必要です。 チューニングスイッチ すべてデフォルトで高速かつ正確なルートになります。すべてのスイッチは出力をバイト単位で同一に保ち、各スイッチでのパリティはゲートの一部です。カーネルをGLSLに戻すスイッチは、PTXカーネルのみがチェイニングに参加するため、カウンターチェイニングもオフになります。 DLSS5VK_UNFUSED=1はGLSLリファレンスを実行します。