DX12のD3D12_HEAP_TYPEについて

2015/04/06 大幅に修正

DX12では、HeapやResourceを確保するときに、CD3D12_HEAP_PROPERTIES を指定するのですが、その際に指定する、D3D12_HEAP_TYPEについて考えてみたいと思います。

D3D12_HEAP_TYPE_DEFAULTについて

 D3D12_HEAP_TYPE_DEFAULTフラグは、いわゆるGPU側のメモリ(VidMem)を確保するためのものです。Map()することはできないので、他のResouceからコピーしてデータを書き換えます。CPU側のメモリなのか、GPU側のメモリなのかはわかりませんが、VMMapを使って確保したメモリのPageの状況を観察することである程度推測できると思います。

TYPE_DEFAULT

 上記は、NVIDIAのGeForce GTX980を使用している時に、32MBのTYPE_DEFAULT Bufferを確保した際に作られた仮想アドレス空間です。確保した領域をVMMapで見ると、Private DataでCommitされていますが、Working Setがアサインされていません。VidMemがMappingされていると思われます。Map()することはできませんが、該当領域のProtectionは、Read/Write/WriteCombineとなっており、WriteCombine(後述)が適用されているのがわかります。
 これらの適用状況は、使用しているシステムのコンフィグレーションによって異なります。DX12のランタイムは、システムがIntelのGPUに代表されるUMA(Unified Memory Arch.)で構成されているのか、PCI Expressバスに接続されているDiscrete Graphicsなのかは把握しており、これらに合わせてリソースを確保し、アプリケーション側に提供するようになっているようです。以降は、NVIDIAのDiscrete Graphicsデバイスを用いている場合の話です。

D3D12_HEAP_TYPE_UPLOADについて

 D3D12_HEAP_TYPE_UPLOADフラグは、Heap/Resourceを確保する際に指定するフラグで、CPUからGPUにデータを転送する用途で使用するHeapに付けます。この領域は、Map()可能で、CPUから情報を書き込めるだけでなく、GPUから直接参照することができるようにも設定できます。そのため、DXのアプリケーション内では大変有用で、ConstantBufferをここに書き込んでShaderから参照したり、VB/IBなどを書き込んで、これをそのまま使用したり、または、その内容をD3D12_HEAP_TYPE_DEFAULT領域にコピーをしたりすることができます。
 書き込む際には、Map()を呼び出し、書き込みを行う場所のポインタを取得する必要がありますが、Unmap()することなく、書き込んだ内容をGPU側で参照することができます。アプリケーションは、当然ながらGPUが使用中のHeap領域を正しく保つ必要があり、これらを正しく制御する仕組みがありますが、今回は割愛します。
 こちらも、CPU側のメモリなのか、GPU側のメモリなのかはわかりませんが、VMMapを使って確保したメモリのPageの状況を観察することである程度推測できると思います。

TYPE_UPLOAD
上記は、32MBのTYPE_UPLOAD Bufferを確保した際に作られた仮想アドレス空間です。Private Working Setがアサインされているので、CPUのメモリがMapされているようです。ProtectionはRead/Write/WriteCombineとなっております。

D3D12_HEAP_TYPE_CUSTOMについて

 D3D12_HEAP_TYPE_CUSTOMフラグを用いると、アプリケーション側で明示的にリソースの配置されるメモリがVidMemかSysMemか、Map可能かなどを指定することができます。

TYPE_CUS_L0_NA
上記は、TYPE_CUSTOM/MEMORY_POOL_L0/D3D12_CPU_PAGE_NOT_AVAILABLEで確保した場合です。これはMap()できません。Private Working Setがアサインされているので、SysMemのようです。

TYPE_CUS_L1_NA
上記は、TYPE_CUSTOM/MEMORY_POOL_L1/D3D12_CPU_PAGE_NOT_AVAILABLEで確保した場合です。これはMap()できません。Private Working Setがアサインされていないので、VidMemのようです。

TYPE_CUS_L0_WC
上記は、TYPE_CUSTOM/MEMORY_POOL_L0/D3D12_CPU_PAGE_WRITE_COMBINEで確保した場合です。Map()可能です。SysMemのようです。

TYPE_CUS_L1_WC
上記は、TYPE_CUSTOM/MEMORY_POOL_L1/D3D12_CPU_PAGE_WRITE_COMBINEで確保した場合です。Map()可能です。VidMemのようです。

TYPE_CUS_L0_WB
上記は、TYPE_CUSTOM/MEMORY_POOL_L0/D3D12_CPU_PAGE_WRITE_BACKで確保した場合です。Map()可能です。SysMemのようです。Write Back(Write Combineが指定されていない)となっているので、通常のmalloc等で取得するメモリと全く同じ状態です。

ちなみに、TYPE_CUSTOM/MEMORY_POOL_L1/D3D12_CPU_PAGE_WRITE_BACKは確保できませんでした。
これらのことから、MEMORY_POOL_L0はSysMem,MEMORY_POOL_L1はVidMemを取得するために使えるようです。

PAGE_WRITECOMBINE属性について

 現在、NVIDIAのGPUを使用して、GPUで読み出し可能な、D3D12_HEAP_TYPE_UPLOADを指定したHeapを確保して、Map()した際に返されるポインタが指すアドレスは、PAGE_WRITECOMBINEという属性が適用されているようです。TYPE_DEFAULTでもWrite Combineが指定されていますが、通常このリソースを指すCPU側のポインタは参照できません。
 CPUがアクセスするメモリにWrite Combineが適用されているかは、下記のようにWindowsのAPIを用いて簡単に確認することができます。

void CheckPageProperties(const void * const addr)
{
  HANDLE hProc = GetCurrentProcess();
  MEMORY_BASIC_INFORMATION memInfo = {};

  VirtualQueryEx(hProc, addr, &memInfo, sizeof(memInfo));

  if (memInfo.Protect & PAGE_WRITECOMBINE) {
    printf("WriteCombine\n");
  }
}

 PAGE_WRITECOMBINE属性が指定されている仮想アドレス領域は、その領域にCPUが書き込むと、CPU内にあるWrite Combine Buffer(WCB)に一時的に内容が蓄積されます。そして、CPUがWCB以外の領域にアクセスすること、それまでWCBに蓄積された内容を実際のアドレスに対して書き出します。WCBは、近年のCPUでは複数(~10)の64Byteのバッファで構成されているそうです。したがって、CPUが該当のアドレス領域に、連続して書き込むと、自然にWBCにバッファリングされて、64Byteごとに実際の書き込み処理が行われることになります。
 なぜ、このようなことが行われるかというと、2通りの考え方ができると思います。
ひとつは、該当の仮想アドレス空間にPCIeのMemoryMapped I/OがMappingされている場合です。これらのアドレスへの書き込みは、最終的にはPCI express busに対する転送処理になります。 PCI express busに対する転送処理の方法は、いくつかの方法があり、WCBによってバッファリングされた書き込みは、可能な限りPCIeのバスに対して効率的に転送処理が行われるように変換されます。具体的には、可能な限りburst transfer形式を用いるように変換されます。こうすることで、CPU側のアプリケーションコードからは、単なるメモリの書き込みと区別できないような処理が、効率的なPCIeバスによる転送処理となります。
 二つ目は、該当の仮想アドレス空間にSytem MemoryがMappingされている場合です。これに関しては、確認が取れてないので、確証をもって言えないのですが、WCBを介したSytem Memoryへの書き込みは、通常のWrite Backのに比べて、早い段階で実際のメモリに書き出されると思われます。WBCのFlushは他のメモリ領域のアクセスによって自動的に引き起こされるのに対して、Write Back, いわゆる通常のCPU のCache構造では、明示的なFlush処理をしない限りは、書き込んだ内容が、いつ実際のメモリに書き出されるかが不明瞭です。PCIe上のDeviceのGPUはCPUのCacheをSnoopingできないかも知れません。そのため、早い段階で実際のメモリに書き出される必要があるため、Write Combineのフラグが適用されているのかもしれません。

PCIeバスに対して、確実にBurst Transferを行うためには

 先に記したとおり、PAGE_WRITECOMBINE属性の付いたメモリ空間への書き込みは、WCBに蓄積されますが、その動作はきわめて透過的で、APIを用いて明示的にBurst Transferを行えるわけではありません。しかし、IntelからPCIeバスに対して確実にBurst Transferを行う方法のドキュメントが出ています。

How to Implement a 64B PCIe Burst Transfer on Intel Architecture

詳細は割愛しますが、ドキュメント内には、16ByteのBurstTransferを行う方法と、64ByteのBurstTransferを行う方法が記されています。

To ensure 16B burst write.
_mm128_store_si128(pcie_memory_address, xmm0);
_mm_mfence(); 

To enusre 64B burst write.
_mm256_store_si256(pcie_memory_address, ymm0);
_mm256_store_si256(pcie_memory_address+32, ymm1);
_mm_mfence();

 どちらも単純で、対象のメモリ空間にMMX/AVXレジスタの内容をストアして、mfenceを呼び出しているだけです。mfenceを用いることで、明示的にWCBをFlushすることになり、対象のアドレス空間への書き込み順序が保たれるようになりますが、一般的なGPUリソースの転送では必要ないと思われます。また、書き込み対象のGPU側のメモリ空間は、64Byte転送の際には、64ByteAlignする必要があり、16Byte転送の際は16ByteAlignする必要があります。また、転送元の内容をレジスタにロードすることを考えると、転送元のアドレスも64ByteAlignされたメモリに配置したほうが好ましいと思われます。
 これらのコードはPCIeバスに対するBurst Transferを確実に行うためのものですが、先のようにSytem MemoryにもWRITE_COMBINE属性が適用されている場合もあります。このようなケースにも上記のようなコードによってSytem Memoryに対しても効率的に転送処理が行われるかは、機会があれば調べてみたいと思います。

まとめ

 以前このブログでOpenGLのImmutable data sotreについてという題で、記事を書いたことがありますが、DX12のTYPE_UPLOADや、TYPE_CUSTOMによるリソース確保は、まさにこの機能に該当するものだと思われます。この機能は、CPU側からGPUリソースの更新が、簡単なコードで実現可能で、リソース更新に伴うAPIオーバーヘッドが存在しないという、DX12の目玉機能の一つと言っても過言ではないと思います。これを用いることで、ConstantBufferの転送だけでなく、CPUによる高度なSkinningの結果を転送するなどの用途も考えられると思います。一方で、このResourceがSysMemに確保されるのか、VidMemに確保されるのかで、アプリケーション側の意識も大きく変える必要があると思います。SysMem上に確保されるならば、データの書き込みによる負荷は、メモリのコピーと同程度と考えてよいと思いますが、GPU上で該当リソースを使用すれば、PCIeバスを介してデータを参照することになると思われるので、あまり大きなリソースの配置は望ましくないと思われます。
対して、VidMemにリソースが確保される場合は、データの書き込みがPCIeバスの帯域とGPUのMemory帯域を占有することを意識する必要があると思われますが、使用する際は、いわゆるVidMem上のリソースと同様に扱えると思われます。