feat(prototypes): étape 2 gpu_interop, import Vulkan vers CUDA (non exécutée)
Code Vulkan commun factorisé dans vk_common (étape 1 inchangée et verte). Nouvel exécutable gpu_interop_vk_cuda : sélection du device CUDA par deviceUUID, cudaImportExternalMemory OPAQUE_WIN32 dédié, mappage en mipmapped array, copie GPU->GPU mesurée, vérification octet à octet. Construit seulement si le CUDA Toolkit est présent ; pas encore compilé contre les vrais en-têtes ni exécuté.
This commit is contained in:
Ange-Marie MAURIN - A2M Systèmes committed
1 parent
65f9b2fab0
commit
32d052cb45
6 files changed
+1553
-890
No files matched your search
@@ -1,14 +1,22 @@
|
||||
# SPDX-License-Identifier: MPL-2.0
|
||||
#
|
||||
# Prototype jetable : export/import d'une texture Vulkan via un HANDLE Win32.
|
||||
# Aucune dependance hors Vulkan SDK.
|
||||
# Prototype jetable : export/import d'une texture Vulkan via un HANDLE Win32,
|
||||
# puis import de la meme allocation dans CUDA (etape 2).
|
||||
|
||||
find_package(Vulkan REQUIRED)
|
||||
|
||||
# Code Vulkan commun aux etapes 1 et 2.
|
||||
add_library(gpu_interop_vk_common STATIC
|
||||
vk_common.cpp)
|
||||
|
||||
target_include_directories(gpu_interop_vk_common PUBLIC "${CMAKE_CURRENT_SOURCE_DIR}")
|
||||
target_link_libraries(gpu_interop_vk_common PUBLIC Vulkan::Vulkan)
|
||||
sequa_target_defaults(gpu_interop_vk_common)
|
||||
|
||||
add_executable(gpu_interop_vk_export
|
||||
main.cpp)
|
||||
|
||||
target_link_libraries(gpu_interop_vk_export PRIVATE Vulkan::Vulkan)
|
||||
target_link_libraries(gpu_interop_vk_export PRIVATE gpu_interop_vk_common)
|
||||
|
||||
sequa_target_defaults(gpu_interop_vk_export)
|
||||
|
||||
@@ -16,3 +24,36 @@ if(SEQUA_BUILD_TESTS)
|
||||
add_test(NAME gpu_interop_vk_export COMMAND gpu_interop_vk_export)
|
||||
set_tests_properties(gpu_interop_vk_export PROPERTIES LABELS "gpu;prototype")
|
||||
endif()
|
||||
|
||||
# Etape 2 : facultative. Sans CUDA Toolkit, la configuration et les autres
|
||||
# tests doivent continuer de passer.
|
||||
find_package(CUDAToolkit QUIET)
|
||||
|
||||
if(CUDAToolkit_FOUND)
|
||||
add_executable(gpu_interop_vk_cuda
|
||||
vk_cuda.cpp)
|
||||
|
||||
target_link_libraries(gpu_interop_vk_cuda PRIVATE gpu_interop_vk_common CUDA::cudart)
|
||||
|
||||
# Les en-tetes CUDA sont du code tiers : SYSTEM pour ne pas declencher
|
||||
# /W4 /WX sur leurs avertissements. Le niveau d'avertissement global reste
|
||||
# inchange.
|
||||
if(CUDAToolkit_INCLUDE_DIRS)
|
||||
target_include_directories(gpu_interop_vk_cuda SYSTEM PRIVATE ${CUDAToolkit_INCLUDE_DIRS})
|
||||
endif()
|
||||
|
||||
sequa_target_defaults(gpu_interop_vk_cuda)
|
||||
|
||||
if(SEQUA_BUILD_TESTS)
|
||||
add_test(NAME gpu_interop_vk_cuda COMMAND gpu_interop_vk_cuda)
|
||||
set_tests_properties(gpu_interop_vk_cuda PROPERTIES LABELS "gpu;prototype;cuda")
|
||||
# cudart est une DLL : ctest doit la trouver meme si le shell appelant
|
||||
# n'a pas encore herite du PATH mis a jour par l'installateur CUDA.
|
||||
if(CUDAToolkit_BIN_DIR)
|
||||
set_tests_properties(gpu_interop_vk_cuda PROPERTIES
|
||||
ENVIRONMENT_MODIFICATION "PATH=path_list_prepend:${CUDAToolkit_BIN_DIR}")
|
||||
endif()
|
||||
endif()
|
||||
else()
|
||||
message(STATUS "CUDA Toolkit absent : gpu_interop_vk_cuda non construit")
|
||||
endif()
|
||||
@@ -2,7 +2,7 @@
|
||||
> Ce dossier sert uniquement à valider une hypothèse technique. Il n'a ni les garanties, ni l'API,
|
||||
> ni la qualité d'ingénierie attendues du reste du dépôt, et il sera supprimé ou réécrit.
|
||||
|
||||
# gpu_interop — export/import d'une texture Vulkan (étape 1)
|
||||
# gpu_interop — partage d'une texture Vulkan vers Vulkan et CUDA (étapes 1 et 2)
|
||||
|
||||
## SEQUA Engine en bref
|
||||
|
||||
@@ -24,21 +24,26 @@ abstractions au passage GPU rendu → inférence** :
|
||||
> les abstractions fuient le plus ; il est testé en tout premier. »
|
||||
|
||||
Ce prototype vérifie donc la brique la plus basse et la plus décisive : une texture produite par un
|
||||
contexte Vulkan peut-elle être **exportée puis réutilisée par un autre contexte GPU** sans jamais
|
||||
transiter par la mémoire système ? Si l'étape 1 échoue, le contrat `NeuralFrameContext` (§5.6) doit
|
||||
prévoir des copies CPU, ce qui change l'architecture et les budgets. L'étape 1 ne dépend que du
|
||||
Vulkan SDK et ne touche ni à CUDA ni à ONNX Runtime.
|
||||
contexte Vulkan peut-elle être **exportée puis réutilisée par un autre contexte GPU** (Vulkan, puis
|
||||
CUDA) sans jamais transiter par la mémoire système ? Si l'étape 1 ou 2 échoue, le contrat
|
||||
`NeuralFrameContext` (§5.6) doit prévoir des copies CPU, ce qui change l'architecture et les budgets.
|
||||
|
||||
## Étapes prévues
|
||||
|
||||
| Étape | Contenu | État |
|
||||
|---|---|---|
|
||||
| 1 | Export d'une `VkImage` en `HANDLE` Win32 (`OPAQUE_WIN32`), import dans un second `VkDevice`, comparaison octet à octet, mesures | **fait (ce programme)** |
|
||||
| 2 | Import de la même allocation dans CUDA via `cudaImportExternalMemory` / `cudaExternalMemoryGetMappedMipmappedArray` | à faire — **nécessite le CUDA Toolkit, non installé sur la machine de référence** |
|
||||
| 1 | Export d'une `VkImage` en `HANDLE` Win32 (`OPAQUE_WIN32`), import dans un second `VkDevice`, comparaison octet à octet, mesures (`gpu_interop_vk_export`) | **fait et mesuré** |
|
||||
| 2 | Import de la même allocation dans CUDA via `cudaImportExternalMemory` / `cudaExternalMemoryGetMappedMipmappedArray`, copie GPU→GPU mesurée et vérification (`gpu_interop_vk_cuda`) | **code écrit, non encore compilé ni exécuté** — CUDA Toolkit 13.5 absent de la machine de référence au moment de la rédaction |
|
||||
| 3 | Inférence ONNX Runtime (ou TensorRT) directement sur le tampon importé, sans copie CPU | à faire |
|
||||
| 4 | Synchronisation fine par sémaphores externes (`VK_KHR_external_semaphore_win32` / `cudaExternalSemaphore`), suppression des `vkQueueWaitIdle` | à faire — l'extension est **disponible** sur le GPU de test, vérifiée et affichée par l'étape 1 |
|
||||
|
||||
## Ce que fait le programme
|
||||
Le code Vulkan commun aux deux étapes (instance, sélection du GPU, device headless, image exportable,
|
||||
staging, motif déterministe, export du `HANDLE`, RAII, aller-retour de référence) est factorisé dans
|
||||
`vk_common.hpp` / `vk_common.cpp`. `main.cpp` (étape 1) et `vk_cuda.cpp` (étape 2) ne gardent que leur
|
||||
logique propre ; **l'étape 1 se comporte exactement comme avant la factorisation** (même sortie, mêmes
|
||||
mesures).
|
||||
|
||||
## Étape 1 — `gpu_interop_vk_export`
|
||||
|
||||
1. Crée une instance Vulkan headless (aucune fenêtre, aucune surface) sur API 1.2.
|
||||
2. Sélectionne un GPU avec file graphique et `VK_KHR_external_memory_win32`, puis crée **deux
|
||||
@@ -58,31 +63,79 @@ Vulkan SDK et ne touche ni à CUDA ni à ONNX Runtime.
|
||||
Code de retour : `0` si les 8 294 400 octets sont identiques, `2` si le contenu diffère, `1` sur
|
||||
erreur Vulkan ou d'initialisation.
|
||||
|
||||
## Étape 2 — `gpu_interop_vk_cuda`
|
||||
|
||||
C++ pur : **aucun `.cu`, aucun noyau CUDA**, uniquement l'API runtime (`CUDA::cudart`). Le programme
|
||||
réutilise `vk_common.hpp` / `vk_common.cpp` (instance, device headless, image exportable, motif,
|
||||
staging, export du `HANDLE`, RAII).
|
||||
|
||||
1. Crée l'image Vulkan exportable remplie du motif et exporte son `HANDLE` `OPAQUE_WIN32`.
|
||||
2. Énumère les devices CUDA (`cudaGetDeviceCount` / `cudaGetDeviceProperties`) et choisit celui dont
|
||||
`cudaDeviceProp::uuid` est **égal au `deviceUUID` Vulkan** (`VkPhysicalDeviceIDProperties`).
|
||||
Échec explicite, avec la liste des UUID vus, si aucun ne correspond.
|
||||
3. Importe la mémoire avec `cudaImportExternalMemory` (`cudaExternalMemoryHandleTypeOpaqueWin32`) en
|
||||
passant `size =` taille de l'allocation Vulkan et `cudaExternalMemoryDedicated` **uniquement**
|
||||
parce que l'allocation Vulkan est dédiée (`VkMemoryDedicatedAllocateInfo` dans
|
||||
`create_filled_exportable_image` ; le drapeau est propagé depuis `ExportableImage::dedicated`).
|
||||
4. Mappe l'image (tiling `OPTIMAL`) avec `cudaExternalMemoryGetMappedMipmappedArray` : format de
|
||||
canal `uchar4` (8 bits × 4 non signé), extent 1920×1080 (depth 0 = tableau 2D), 1 niveau, puis
|
||||
`cudaGetMipmappedArrayLevel(…, 0)`.
|
||||
5. Copie le niveau 0 vers un tampon device avec `cudaMemcpy2DFromArray` — **c'est le chemin mesuré
|
||||
GPU → GPU** : `std::chrono` autour de l'appel et `cudaEvent` pour le temps GPU.
|
||||
6. Pour la **vérification uniquement**, copie le tampon device vers l'hôte et compare octet à octet
|
||||
au motif (cette copie est hors du chemin mesuré).
|
||||
7. Mesure séparément l'import, le mappage et la copie, puis rappelle l'aller-retour GPU → CPU → GPU
|
||||
de l'étape 1 comme référence.
|
||||
8. Libère tout : `cudaFree`, `cudaFreeMipmappedArray`, `cudaDestroyExternalMemory`, `CloseHandle`.
|
||||
|
||||
Les ressources CUDA sont libérées dans une portée dédiée **avant** l'aller-retour de référence Vulkan :
|
||||
Vulkan réécrit l'image ensuite, donc plus aucun objet CUDA ne doit la référencer. Le `CloseHandle`
|
||||
intervient en dernier, après la destruction de l'objet `cudaExternalMemory` (CUDA garde sa propre
|
||||
référence sur le handle ; il ne prend pas possession de celui de l'application).
|
||||
|
||||
Synchronisation : **`vkQueueWaitIdle` avant l'export suffit pour cette étape**. Les sémaphores
|
||||
externes (`VK_KHR_external_semaphore_win32` / `cudaExternalSemaphore`) sont l'étape 4 ; le programme
|
||||
ne prétend donc pas mesurer un pipeline asynchrone de production.
|
||||
|
||||
## Construire et lancer
|
||||
|
||||
Prérequis : Windows 10/11, Visual Studio 2022 avec les outils C++ x64, Vulkan SDK (testé avec
|
||||
Prérequis étape 1 : Windows 10/11, Visual Studio 2022 avec les outils C++ x64, Vulkan SDK (testé avec
|
||||
1.4.341.1), un GPU exposant `VK_KHR_external_memory_win32`.
|
||||
|
||||
Prérequis étape 2 : CUDA Toolkit (13.5 sur la machine de référence) et un GPU dont l'UUID CUDA
|
||||
correspond à celui exposé par Vulkan. **Sans CUDA Toolkit, la cible est simplement ignorée** et le
|
||||
reste du dépôt continue de se configurer et de se tester (`find_package(CUDAToolkit QUIET)` puis
|
||||
message `STATUS`). Quand la cible est construite, `ctest` reçoit automatiquement le dossier `bin` du
|
||||
Toolkit dans son `PATH` (`ENVIRONMENT_MODIFICATION`), afin de trouver `cudart64_*.dll` même si le shell
|
||||
appelant n'a pas encore hérité du `PATH` mis à jour par l'installateur.
|
||||
|
||||
Depuis la racine du dépôt :
|
||||
|
||||
```powershell
|
||||
powershell.exe -NoProfile -ExecutionPolicy Bypass -File scripts\build.ps1 -Preset windows-msvc-prototypes
|
||||
```
|
||||
|
||||
Le script configure, compile et lance `ctest`. Pour exécuter seulement l'essai du prototype :
|
||||
Le script configure, compile et lance `ctest`. Pour exécuter seulement les prototypes :
|
||||
|
||||
```powershell
|
||||
ctest --preset windows-msvc-prototypes -L prototype --output-on-failure
|
||||
```
|
||||
|
||||
Étape 2 seulement (si elle a été construite) :
|
||||
|
||||
```powershell
|
||||
ctest --preset windows-msvc-prototypes -L cuda --output-on-failure
|
||||
```
|
||||
|
||||
Ou directement :
|
||||
|
||||
```powershell
|
||||
.\build\windows-msvc-prototypes\prototypes\gpu_interop\gpu_interop_vk_export.exe
|
||||
.\build\windows-msvc-prototypes\prototypes\gpu_interop\gpu_interop_vk_cuda.exe
|
||||
```
|
||||
|
||||
Sous Linux/macOS, `prototypes/CMakeLists.txt` n'ajoute pas la cible et affiche un message `STATUS` :
|
||||
le code de cette étape est Windows uniquement.
|
||||
Sous Linux/macOS, `prototypes/CMakeLists.txt` n'ajoute pas les cibles et affiche un message `STATUS` :
|
||||
le code de ces étapes est Windows uniquement.
|
||||
|
||||
## Résultats mesurés
|
||||
|
||||
@@ -91,27 +144,31 @@ Machine de référence : Windows 11, NVIDIA GeForce RTX 4090 Laptop GPU (16 Go),
|
||||
|
||||
Image : 1920×1080 `R8G8B8A8_UNORM`, 8 294 400 octets utiles, allocation de 8 847 360 octets.
|
||||
|
||||
### Étape 1 (mesurée)
|
||||
|
||||
| Mesure | Valeur |
|
||||
|---|---|
|
||||
| Export `vkGetMemoryWin32HandleKHR` | **0,001 ms** |
|
||||
| Import dans le second `VkDevice` (création image + `vkAllocateMemory` + `vkBindImageMemory`) | **0,066 – 0,135 ms** (selon l'exécution) |
|
||||
| Aller-retour GPU → CPU → GPU de référence (médiane de 5 essais) | **30,7 – 37,0 ms** |
|
||||
| Export `vkGetMemoryWin32HandleKHR` | **0,000 – 0,001 ms** |
|
||||
| Import dans le second `VkDevice` (création image + `vkAllocateMemory` + `vkBindImageMemory`) | **0,045 – 0,135 ms** (selon l'exécution) |
|
||||
| Aller-retour GPU → CPU → GPU de référence (médiane de 5 essais) | **30,5 – 37,0 ms** |
|
||||
| Comparaison octet à octet après import | **identique** (8 294 400 / 8 294 400 octets) |
|
||||
|
||||
Extrait d'exécution :
|
||||
|
||||
```text
|
||||
[5] Export : HANDLE Win32 0x00000000000005F4 en 0.001 ms
|
||||
[6] Import dans le device B : type 1, 0.075 ms (creation image + alloc + bind)
|
||||
[3] Image A : memoire DEVICE_LOCAL exportable (type 1, 8847360 octets)
|
||||
[4] Motif deterministe uploade (staging buffer + vkCmdCopyBufferToImage)
|
||||
[5] Export : HANDLE Win32 0x00000000000005E4 en 0.000 ms
|
||||
[6] Import dans le device B : type 1, 0.045 ms (creation image + alloc + bind)
|
||||
[7] Comparaison octet a octet apres import : IDENTIQUE
|
||||
[8] CloseHandle du HANDLE exporte : OK
|
||||
[9] Aller-retour GPU->CPU->GPU de reference (5 essais)
|
||||
min 30.612 ms, mediane 30.725 ms, contenu identique
|
||||
min 30.048 ms, mediane 30.528 ms, contenu identique
|
||||
|
||||
=== RESULTATS ===
|
||||
Export HANDLE Win32 : 0.001 ms
|
||||
Import dans un 2e device : 0.075 ms
|
||||
Aller-retour GPU->CPU->GPU : 30.725 ms (mediane de 5 essais)
|
||||
Export HANDLE Win32 : 0.000 ms
|
||||
Import dans un 2e device : 0.045 ms
|
||||
Aller-retour GPU->CPU->GPU : 30.528 ms (mediane de 5 essais)
|
||||
Verification octet a octet : OK
|
||||
```
|
||||
|
||||
@@ -120,35 +177,58 @@ temps négligeables devant le round-trip CPU, qui déplace réellement 2 × 8,3
|
||||
deux synchronisations `vkQueueWaitIdle` (latence WDDM). C'est exactement l'écart que l'étape 2 doit
|
||||
rendre inutile.
|
||||
|
||||
### Étape 2 (non encore exécutée)
|
||||
|
||||
**Aucune mesure CUDA disponible.** Le CUDA Toolkit 13.5 n'était pas installé au moment de la
|
||||
rédaction (`nvcc.exe` absent, `CUDA_PATH` vide) : la cible n'a donc été **ni compilée avec les vrais
|
||||
en-têtes 13.5, ni exécutée**. Le CMake a seulement été validé dans le cas « CUDA absent » (message
|
||||
`STATUS`, configuration, build et tests existants inchangés).
|
||||
|
||||
Seule une vérification de syntaxe a été faite localement, avec un `cuda_runtime.h` factice hors dépôt
|
||||
et les mêmes options `/W4 /WX` : `vk_cuda.cpp` compile sans avertissement contre cette déclaration
|
||||
approximative, ce qui élimine les fautes de C++ évidentes mais ne prouve ni la conformité aux vrais
|
||||
en-têtes, ni la signature exacte des fonctions, ni le comportement à l'exécution. Les valeurs à
|
||||
relever après installation sont : import CUDA, mappage mipmapped array, copie
|
||||
`cudaMemcpy2DFromArray` (chrono + cudaEvent, avec débit), et comparaison octet à octet.
|
||||
|
||||
## Limites et points non vérifiés
|
||||
|
||||
- Import d'un handle `OPAQUE_WIN32` : `vkGetMemoryWin32HandlePropertiesKHR` ne doit **pas** être
|
||||
appelé pour ce type de handle (VUID-vkGetMemoryWin32HandlePropertiesKHR-handleType-00666 ; la
|
||||
première version l'appelait et obtenait `VK_ERROR_INITIALIZATION_FAILED`). La spec impose
|
||||
d'importer avec le même `memoryTypeIndex` et la même taille qu'à l'export ; le programme vérifie
|
||||
seulement que l'image importée accepte ce type. Non retesté avec les validation layers.
|
||||
- La synchronisation inter-devices repose sur un `vkQueueWaitIdle` avant l'export : suffisant pour un
|
||||
prototype mono-thread, **pas** un mécanisme de production. L'étape 4 doit le remplacer par des
|
||||
sémaphores externes.
|
||||
- Les deux `VkDevice` vivent dans le **même processus**. Un partage inter-processus demanderait
|
||||
`DuplicateHandle`.
|
||||
- La vérification `deviceUUID` compare deux requêtes sur le même `VkPhysicalDevice` : elle est donc
|
||||
structurellement vraie ici. Elle deviendra significative quand le device d'inférence sera créé par
|
||||
un autre runtime (CUDA) ou un autre processus.
|
||||
appelé pour ce type de handle (VUID-vkGetMemoryWin32HandlePropertiesKHR-handleType-00666). La spec
|
||||
impose d'importer avec le même `memoryTypeIndex` et la même taille qu'à l'export ; le programme
|
||||
vérifie seulement que l'image importée accepte ce type. Non retesté avec les validation layers.
|
||||
- **Étape 2 non compilée ni exécutée** (CUDA Toolkit absent). Les signatures CUDA utilisées sont
|
||||
celles de l'API runtime publique 13.5 (relues dans la documentation NVIDIA), mais rien n'a été
|
||||
confronté aux vrais en-têtes ni au pilote.
|
||||
- Le comportement exact du pilote pour `cudaExternalMemoryGetMappedMipmappedArray` sur une image
|
||||
Vulkan `OPTIMAL` **n'a pas été observé** (contrainte de taille d'allocation, compatibilité exacte
|
||||
du format `R8G8B8A8_UNORM` ↔ `uchar4`, cohérence de cache après écriture Vulkan ultérieure).
|
||||
- La synchronisation repose sur un `vkQueueWaitIdle` avant l'export : suffisant pour un prototype
|
||||
mono-thread, **pas** un mécanisme de production. L'étape 4 doit le remplacer par des sémaphores
|
||||
externes, côté Vulkan et côté CUDA.
|
||||
- Les deux `VkDevice` de l'étape 1 vivent dans le **même processus**. Un partage inter-processus
|
||||
demanderait `DuplicateHandle`.
|
||||
- La vérification `deviceUUID` de l'étape 1 compare deux requêtes sur le même `VkPhysicalDevice` :
|
||||
elle est structurellement vraie. L'étape 2 la rend réellement significative en la confrontant à
|
||||
l'UUID du device CUDA.
|
||||
- L'aller-retour de référence inclut deux soumissions, deux attentes de file et les copies hôte ;
|
||||
il mesure donc surtout la latence de synchronisation Windows, pas la bande passante PCIe.
|
||||
- Les temps export/import sont des mesures uniques par exécution (ils ne peuvent pas être répétés sans
|
||||
réexporter) ; seule la variabilité de l'aller-retour est établie sur 5 essais.
|
||||
- Aucune couche de validation Vulkan n'est activée : le prototype n'a pas été passé sous
|
||||
`VK_LAYER_KHRONOS_validation`.
|
||||
- Étape 2 non commencée : aucun CUDA Toolkit n'est installé, aucune mesure CUDA n'est donc disponible.
|
||||
- Si les en-têtes CUDA déclenchent des avertissements sous `/W4`, la cible les inclut en `SYSTEM` ;
|
||||
le niveau d'avertissement global n'est pas abaissé. Ce point n'a pas pu être vérifié sans Toolkit.
|
||||
|
||||
## Licences
|
||||
|
||||
- Ce prototype est publié sous **MPL 2.0**, comme le reste de SEQUA (voir `LICENSE` à la racine).
|
||||
Chaque fichier source porte l'en-tête `SPDX-License-Identifier: MPL-2.0`.
|
||||
- Dépendances de l'étape 1 : **Vulkan SDK / Vulkan-Loader** (Apache 2.0) uniquement. Aucune autre
|
||||
bibliothèque n'est liée : pas de glslang, pas de validation layer, pas de framework de test dédié.
|
||||
- Dépendances de l'étape 1 : **Vulkan SDK / Vulkan-Loader** (Apache 2.0) uniquement.
|
||||
- Dépendance de l'étape 2 : **CUDA Toolkit**, propriétaire NVIDIA, distribué sous l'EULA NVIDIA.
|
||||
C'est une dépendance **facultative du prototype uniquement** : elle n'entre jamais dans le cœur,
|
||||
la cible n'est construite que si `CUDAToolkit` est trouvé, et l'engine doit rester compilable et
|
||||
utilisable sans CUDA.
|
||||
- **CUDA et TensorRT sont propriétaires (NVIDIA)**. Ils n'entreront jamais dans le cœur : ce seront
|
||||
des plugins facultatifs derrière `IComputeBackend`, conformément au §10.2 du document d'architecture,
|
||||
et ils ne seront pas requis pour compiler ou exécuter l'engine.
|
||||
des plugins facultatifs derrière `IComputeBackend`, conformément au §10.2 du document
|
||||
d'architecture, et ils ne seront pas requis pour compiler ou exécuter l'engine.
|
||||
+27
-849
@@ -17,715 +17,23 @@
|
||||
//
|
||||
// Windows uniquement (VK_KHR_external_memory_win32). Code non produit.
|
||||
|
||||
#ifndef NOMINMAX
|
||||
#define NOMINMAX
|
||||
#endif
|
||||
#ifndef WIN32_LEAN_AND_MEAN
|
||||
#define WIN32_LEAN_AND_MEAN
|
||||
#endif
|
||||
#include "vk_common.hpp"
|
||||
|
||||
#include <windows.h>
|
||||
|
||||
#include <vulkan/vulkan.h>
|
||||
#include <vulkan/vulkan_win32.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <array>
|
||||
#include <chrono>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cstring>
|
||||
#include <iterator>
|
||||
#include <stdexcept>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
namespace {
|
||||
|
||||
using Clock = std::chrono::steady_clock;
|
||||
using namespace gpu_interop;
|
||||
|
||||
constexpr uint32_t kWidth = 1920;
|
||||
constexpr uint32_t kHeight = 1080;
|
||||
constexpr VkDeviceSize kBytesPerPixel = 4;
|
||||
constexpr VkDeviceSize kImageBytes = static_cast<VkDeviceSize>(kWidth) * kHeight * kBytesPerPixel;
|
||||
constexpr VkFormat kFormat = VK_FORMAT_R8G8B8A8_UNORM;
|
||||
constexpr int kRoundtripIterations = 5;
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Erreurs et formatage
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
const char* result_name(VkResult result)
|
||||
{
|
||||
switch (result) {
|
||||
case VK_SUCCESS:
|
||||
return "VK_SUCCESS";
|
||||
case VK_NOT_READY:
|
||||
return "VK_NOT_READY";
|
||||
case VK_TIMEOUT:
|
||||
return "VK_TIMEOUT";
|
||||
case VK_EVENT_SET:
|
||||
return "VK_EVENT_SET";
|
||||
case VK_EVENT_RESET:
|
||||
return "VK_EVENT_RESET";
|
||||
case VK_INCOMPLETE:
|
||||
return "VK_INCOMPLETE";
|
||||
case VK_ERROR_OUT_OF_HOST_MEMORY:
|
||||
return "VK_ERROR_OUT_OF_HOST_MEMORY";
|
||||
case VK_ERROR_OUT_OF_DEVICE_MEMORY:
|
||||
return "VK_ERROR_OUT_OF_DEVICE_MEMORY";
|
||||
case VK_ERROR_INITIALIZATION_FAILED:
|
||||
return "VK_ERROR_INITIALIZATION_FAILED";
|
||||
case VK_ERROR_DEVICE_LOST:
|
||||
return "VK_ERROR_DEVICE_LOST";
|
||||
case VK_ERROR_MEMORY_MAP_FAILED:
|
||||
return "VK_ERROR_MEMORY_MAP_FAILED";
|
||||
case VK_ERROR_LAYER_NOT_PRESENT:
|
||||
return "VK_ERROR_LAYER_NOT_PRESENT";
|
||||
case VK_ERROR_EXTENSION_NOT_PRESENT:
|
||||
return "VK_ERROR_EXTENSION_NOT_PRESENT";
|
||||
case VK_ERROR_FEATURE_NOT_PRESENT:
|
||||
return "VK_ERROR_FEATURE_NOT_PRESENT";
|
||||
case VK_ERROR_INCOMPATIBLE_DRIVER:
|
||||
return "VK_ERROR_INCOMPATIBLE_DRIVER";
|
||||
case VK_ERROR_TOO_MANY_OBJECTS:
|
||||
return "VK_ERROR_TOO_MANY_OBJECTS";
|
||||
case VK_ERROR_FORMAT_NOT_SUPPORTED:
|
||||
return "VK_ERROR_FORMAT_NOT_SUPPORTED";
|
||||
case VK_ERROR_INVALID_EXTERNAL_HANDLE:
|
||||
return "VK_ERROR_INVALID_EXTERNAL_HANDLE";
|
||||
case VK_ERROR_OUT_OF_DATE_KHR:
|
||||
return "VK_ERROR_OUT_OF_DATE_KHR";
|
||||
default:
|
||||
return "VkResult inconnu";
|
||||
}
|
||||
}
|
||||
|
||||
[[noreturn]] void throw_vk(const char* what, VkResult result)
|
||||
{
|
||||
char buffer[256];
|
||||
std::snprintf(buffer, sizeof(buffer), "%s : %s (%d)", what, result_name(result), static_cast<int>(result));
|
||||
throw std::runtime_error(buffer);
|
||||
}
|
||||
|
||||
void check(VkResult result, const char* what)
|
||||
{
|
||||
if (result != VK_SUCCESS) {
|
||||
throw_vk(what, result);
|
||||
}
|
||||
}
|
||||
|
||||
double elapsed_ms(Clock::time_point start)
|
||||
{
|
||||
return std::chrono::duration<double, std::milli>(Clock::now() - start).count();
|
||||
}
|
||||
|
||||
void print_external_format_support(VkPhysicalDevice physicalDevice)
|
||||
{
|
||||
VkPhysicalDeviceExternalImageFormatInfo externalInfo{};
|
||||
externalInfo.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTERNAL_IMAGE_FORMAT_INFO;
|
||||
externalInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
VkPhysicalDeviceImageFormatInfo2 formatInfo{};
|
||||
formatInfo.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGE_FORMAT_INFO_2;
|
||||
formatInfo.pNext = &externalInfo;
|
||||
formatInfo.format = kFormat;
|
||||
formatInfo.type = VK_IMAGE_TYPE_2D;
|
||||
formatInfo.tiling = VK_IMAGE_TILING_OPTIMAL;
|
||||
formatInfo.usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT;
|
||||
|
||||
VkExternalImageFormatProperties externalProperties{};
|
||||
externalProperties.sType = VK_STRUCTURE_TYPE_EXTERNAL_IMAGE_FORMAT_PROPERTIES;
|
||||
VkImageFormatProperties2 properties{};
|
||||
properties.sType = VK_STRUCTURE_TYPE_IMAGE_FORMAT_PROPERTIES_2;
|
||||
properties.pNext = &externalProperties;
|
||||
|
||||
const VkResult result = vkGetPhysicalDeviceImageFormatProperties2(physicalDevice, &formatInfo, &properties);
|
||||
if (result != VK_SUCCESS) {
|
||||
std::printf(" VkExternalImageFormatProperties 2D/R8G8B8A8_UNORM/OPAQUE_WIN32 : %s\n", result_name(result));
|
||||
return;
|
||||
}
|
||||
|
||||
const VkExternalMemoryFeatureFlags features = externalProperties.externalMemoryProperties.externalMemoryFeatures;
|
||||
std::printf(" VkExternalImageFormatProperties 2D/R8G8B8A8_UNORM/OPAQUE_WIN32 : exportable=%s importable=%s "
|
||||
"dedicated_only=%s\n",
|
||||
(features & VK_EXTERNAL_MEMORY_FEATURE_EXPORTABLE_BIT) != 0u ? "oui" : "non",
|
||||
(features & VK_EXTERNAL_MEMORY_FEATURE_IMPORTABLE_BIT) != 0u ? "oui" : "non",
|
||||
(features & VK_EXTERNAL_MEMORY_FEATURE_DEDICATED_ONLY_BIT) != 0u ? "oui" : "non");
|
||||
}
|
||||
|
||||
const char* device_type_name(VkPhysicalDeviceType type)
|
||||
{
|
||||
switch (type) {
|
||||
case VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU:
|
||||
return "discrete";
|
||||
case VK_PHYSICAL_DEVICE_TYPE_INTEGRATED_GPU:
|
||||
return "integrated";
|
||||
case VK_PHYSICAL_DEVICE_TYPE_VIRTUAL_GPU:
|
||||
return "virtual";
|
||||
case VK_PHYSICAL_DEVICE_TYPE_CPU:
|
||||
return "cpu";
|
||||
default:
|
||||
return "other";
|
||||
}
|
||||
}
|
||||
|
||||
std::string uuid_to_string(const std::array<uint8_t, VK_UUID_SIZE>& uuid)
|
||||
{
|
||||
std::string text;
|
||||
text.reserve(uuid.size() * 2);
|
||||
char byte[4];
|
||||
for (uint8_t value : uuid) {
|
||||
std::snprintf(byte, sizeof(byte), "%02x", static_cast<unsigned int>(value));
|
||||
text += byte;
|
||||
}
|
||||
return text;
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// RAII minimal
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
struct Win32Handle {
|
||||
HANDLE value = nullptr;
|
||||
|
||||
Win32Handle() = default;
|
||||
Win32Handle(const Win32Handle&) = delete;
|
||||
Win32Handle& operator=(const Win32Handle&) = delete;
|
||||
~Win32Handle()
|
||||
{
|
||||
if (value != nullptr) {
|
||||
(void)CloseHandle(value);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct InstanceContext {
|
||||
VkInstance instance = VK_NULL_HANDLE;
|
||||
|
||||
InstanceContext() = default;
|
||||
InstanceContext(const InstanceContext&) = delete;
|
||||
InstanceContext& operator=(const InstanceContext&) = delete;
|
||||
~InstanceContext()
|
||||
{
|
||||
if (instance != VK_NULL_HANDLE) {
|
||||
vkDestroyInstance(instance, nullptr);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct DeviceContext {
|
||||
VkPhysicalDevice physicalDevice = VK_NULL_HANDLE;
|
||||
VkDevice device = VK_NULL_HANDLE;
|
||||
VkQueue queue = VK_NULL_HANDLE;
|
||||
uint32_t queueFamily = 0;
|
||||
VkCommandPool commandPool = VK_NULL_HANDLE;
|
||||
VkCommandBuffer commandBuffer = VK_NULL_HANDLE;
|
||||
|
||||
DeviceContext() = default;
|
||||
DeviceContext(const DeviceContext&) = delete;
|
||||
DeviceContext& operator=(const DeviceContext&) = delete;
|
||||
~DeviceContext()
|
||||
{
|
||||
if (device != VK_NULL_HANDLE) {
|
||||
(void)vkDeviceWaitIdle(device);
|
||||
if (commandPool != VK_NULL_HANDLE) {
|
||||
vkDestroyCommandPool(device, commandPool, nullptr);
|
||||
}
|
||||
vkDestroyDevice(device, nullptr);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct ImageResource {
|
||||
VkDevice device = VK_NULL_HANDLE;
|
||||
VkImage image = VK_NULL_HANDLE;
|
||||
VkDeviceMemory memory = VK_NULL_HANDLE;
|
||||
|
||||
ImageResource() = default;
|
||||
ImageResource(const ImageResource&) = delete;
|
||||
ImageResource& operator=(const ImageResource&) = delete;
|
||||
~ImageResource()
|
||||
{
|
||||
if (device != VK_NULL_HANDLE) {
|
||||
if (image != VK_NULL_HANDLE) {
|
||||
vkDestroyImage(device, image, nullptr);
|
||||
}
|
||||
if (memory != VK_NULL_HANDLE) {
|
||||
vkFreeMemory(device, memory, nullptr);
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct BufferResource {
|
||||
VkDevice device = VK_NULL_HANDLE;
|
||||
VkBuffer buffer = VK_NULL_HANDLE;
|
||||
VkDeviceMemory memory = VK_NULL_HANDLE;
|
||||
VkDeviceSize size = 0;
|
||||
|
||||
BufferResource() = default;
|
||||
BufferResource(const BufferResource&) = delete;
|
||||
BufferResource& operator=(const BufferResource&) = delete;
|
||||
~BufferResource()
|
||||
{
|
||||
if (device != VK_NULL_HANDLE) {
|
||||
if (buffer != VK_NULL_HANDLE) {
|
||||
vkDestroyBuffer(device, buffer, nullptr);
|
||||
}
|
||||
if (memory != VK_NULL_HANDLE) {
|
||||
vkFreeMemory(device, memory, nullptr);
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Enumeration / selection
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
std::vector<VkExtensionProperties> enumerate_instance_extensions()
|
||||
{
|
||||
uint32_t count = 0;
|
||||
check(vkEnumerateInstanceExtensionProperties(nullptr, &count, nullptr), "vkEnumerateInstanceExtensionProperties");
|
||||
std::vector<VkExtensionProperties> extensions(count);
|
||||
if (count != 0) {
|
||||
check(vkEnumerateInstanceExtensionProperties(nullptr, &count, extensions.data()),
|
||||
"vkEnumerateInstanceExtensionProperties");
|
||||
}
|
||||
return extensions;
|
||||
}
|
||||
|
||||
std::vector<VkExtensionProperties> enumerate_device_extensions(VkPhysicalDevice physicalDevice)
|
||||
{
|
||||
uint32_t count = 0;
|
||||
check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, nullptr),
|
||||
"vkEnumerateDeviceExtensionProperties");
|
||||
std::vector<VkExtensionProperties> extensions(count);
|
||||
if (count != 0) {
|
||||
check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, extensions.data()),
|
||||
"vkEnumerateDeviceExtensionProperties");
|
||||
}
|
||||
return extensions;
|
||||
}
|
||||
|
||||
bool has_extension(const std::vector<VkExtensionProperties>& extensions, const char* name)
|
||||
{
|
||||
for (const VkExtensionProperties& extension : extensions) {
|
||||
if (std::strcmp(extension.extensionName, name) == 0) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
void create_instance(InstanceContext& out)
|
||||
{
|
||||
uint32_t loaderVersion = 0;
|
||||
check(vkEnumerateInstanceVersion(&loaderVersion), "vkEnumerateInstanceVersion");
|
||||
if (loaderVersion < VK_API_VERSION_1_2) {
|
||||
char buffer[160];
|
||||
std::snprintf(buffer, sizeof(buffer), "Vulkan loader trop ancien : %u.%u.%u (1.2 minimum requis)",
|
||||
static_cast<unsigned int>(VK_API_VERSION_MAJOR(loaderVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_MINOR(loaderVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_PATCH(loaderVersion)));
|
||||
throw std::runtime_error(buffer);
|
||||
}
|
||||
|
||||
const std::vector<VkExtensionProperties> available = enumerate_instance_extensions();
|
||||
std::vector<const char*> enabled;
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME);
|
||||
}
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME);
|
||||
}
|
||||
|
||||
VkApplicationInfo application{};
|
||||
application.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;
|
||||
application.pApplicationName = "gpu_interop_vk_export";
|
||||
application.applicationVersion = VK_MAKE_API_VERSION(0, 0, 1, 0);
|
||||
application.pEngineName = "SEQUA (prototype)";
|
||||
application.engineVersion = VK_MAKE_API_VERSION(0, 0, 1, 0);
|
||||
application.apiVersion = VK_API_VERSION_1_2;
|
||||
|
||||
VkInstanceCreateInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;
|
||||
info.pApplicationInfo = &application;
|
||||
info.enabledExtensionCount = static_cast<uint32_t>(enabled.size());
|
||||
info.ppEnabledExtensionNames = enabled.data();
|
||||
|
||||
check(vkCreateInstance(&info, nullptr, &out.instance), "vkCreateInstance");
|
||||
}
|
||||
|
||||
std::array<uint8_t, VK_UUID_SIZE> query_device_uuid(VkPhysicalDevice physicalDevice)
|
||||
{
|
||||
VkPhysicalDeviceIDProperties idProperties{};
|
||||
idProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES;
|
||||
VkPhysicalDeviceProperties2 properties2{};
|
||||
properties2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2;
|
||||
properties2.pNext = &idProperties;
|
||||
vkGetPhysicalDeviceProperties2(physicalDevice, &properties2);
|
||||
|
||||
std::array<uint8_t, VK_UUID_SIZE> uuid{};
|
||||
std::copy(std::begin(idProperties.deviceUUID), std::end(idProperties.deviceUUID), uuid.begin());
|
||||
return uuid;
|
||||
}
|
||||
|
||||
struct DeviceInfo {
|
||||
VkPhysicalDevice physicalDevice = VK_NULL_HANDLE;
|
||||
uint32_t queueFamily = 0;
|
||||
VkPhysicalDeviceProperties properties{};
|
||||
std::array<uint8_t, VK_UUID_SIZE> uuid{};
|
||||
bool externalMemory = false;
|
||||
bool externalMemoryWin32 = false;
|
||||
bool externalSemaphore = false;
|
||||
bool externalSemaphoreWin32 = false;
|
||||
};
|
||||
|
||||
DeviceInfo pick_physical_device(VkInstance instance)
|
||||
{
|
||||
uint32_t count = 0;
|
||||
check(vkEnumeratePhysicalDevices(instance, &count, nullptr), "vkEnumeratePhysicalDevices");
|
||||
if (count == 0) {
|
||||
throw std::runtime_error("Aucun GPU Vulkan detecte.");
|
||||
}
|
||||
std::vector<VkPhysicalDevice> devices(count);
|
||||
check(vkEnumeratePhysicalDevices(instance, &count, devices.data()), "vkEnumeratePhysicalDevices");
|
||||
|
||||
std::string rejected;
|
||||
for (VkPhysicalDevice physicalDevice : devices) {
|
||||
DeviceInfo info{};
|
||||
info.physicalDevice = physicalDevice;
|
||||
|
||||
VkPhysicalDeviceProperties2 properties2{};
|
||||
properties2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2;
|
||||
vkGetPhysicalDeviceProperties2(physicalDevice, &properties2);
|
||||
info.properties = properties2.properties;
|
||||
info.uuid = query_device_uuid(physicalDevice);
|
||||
|
||||
const std::vector<VkExtensionProperties> extensions = enumerate_device_extensions(physicalDevice);
|
||||
info.externalMemory = has_extension(extensions, VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME);
|
||||
info.externalMemoryWin32 = has_extension(extensions, VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME);
|
||||
info.externalSemaphore = has_extension(extensions, VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME);
|
||||
info.externalSemaphoreWin32 = has_extension(extensions, VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME);
|
||||
|
||||
uint32_t familyCount = 0;
|
||||
vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &familyCount, nullptr);
|
||||
std::vector<VkQueueFamilyProperties> families(familyCount);
|
||||
if (familyCount != 0) {
|
||||
vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &familyCount, families.data());
|
||||
}
|
||||
bool hasGraphicsQueue = false;
|
||||
for (uint32_t index = 0; index < familyCount; ++index) {
|
||||
if ((families[index].queueFlags & VK_QUEUE_GRAPHICS_BIT) != 0u) {
|
||||
info.queueFamily = index;
|
||||
hasGraphicsQueue = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
std::printf(" GPU : %s (%s, API %u.%u.%u) - external_memory_win32=%s\n", info.properties.deviceName,
|
||||
device_type_name(info.properties.deviceType),
|
||||
static_cast<unsigned int>(VK_API_VERSION_MAJOR(info.properties.apiVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_MINOR(info.properties.apiVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_PATCH(info.properties.apiVersion)),
|
||||
info.externalMemoryWin32 ? "oui" : "non");
|
||||
|
||||
if (hasGraphicsQueue && info.externalMemory && info.externalMemoryWin32) {
|
||||
return info;
|
||||
}
|
||||
rejected += " - ignore : ";
|
||||
rejected += info.properties.deviceName;
|
||||
rejected += hasGraphicsQueue ? "" : " (pas de file graphique)";
|
||||
rejected += info.externalMemoryWin32 ? "" : " (VK_KHR_external_memory_win32 absent)";
|
||||
rejected += '\n';
|
||||
}
|
||||
|
||||
throw std::runtime_error("Aucun GPU ne remplit les conditions (file graphique + VK_KHR_external_memory_win32).\n" +
|
||||
rejected);
|
||||
}
|
||||
|
||||
void create_device(const DeviceInfo& info, bool enableExternalSemaphore, DeviceContext& out, const char* label)
|
||||
{
|
||||
const std::vector<VkExtensionProperties> available = enumerate_device_extensions(info.physicalDevice);
|
||||
std::vector<const char*> enabled;
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME);
|
||||
}
|
||||
enabled.push_back(VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME);
|
||||
if (enableExternalSemaphore) {
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME);
|
||||
}
|
||||
if (info.externalSemaphoreWin32) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME);
|
||||
}
|
||||
}
|
||||
|
||||
const float priority = 1.0f;
|
||||
VkDeviceQueueCreateInfo queueInfo{};
|
||||
queueInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;
|
||||
queueInfo.queueFamilyIndex = info.queueFamily;
|
||||
queueInfo.queueCount = 1;
|
||||
queueInfo.pQueuePriorities = &priority;
|
||||
|
||||
VkDeviceCreateInfo createInfo{};
|
||||
createInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;
|
||||
createInfo.queueCreateInfoCount = 1;
|
||||
createInfo.pQueueCreateInfos = &queueInfo;
|
||||
createInfo.enabledExtensionCount = static_cast<uint32_t>(enabled.size());
|
||||
createInfo.ppEnabledExtensionNames = enabled.data();
|
||||
|
||||
check(vkCreateDevice(info.physicalDevice, &createInfo, nullptr, &out.device), label);
|
||||
out.physicalDevice = info.physicalDevice;
|
||||
out.queueFamily = info.queueFamily;
|
||||
vkGetDeviceQueue(out.device, info.queueFamily, 0, &out.queue);
|
||||
|
||||
VkCommandPoolCreateInfo poolInfo{};
|
||||
poolInfo.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;
|
||||
poolInfo.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
|
||||
poolInfo.queueFamilyIndex = info.queueFamily;
|
||||
check(vkCreateCommandPool(out.device, &poolInfo, nullptr, &out.commandPool), "vkCreateCommandPool");
|
||||
|
||||
VkCommandBufferAllocateInfo bufferInfo{};
|
||||
bufferInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;
|
||||
bufferInfo.commandPool = out.commandPool;
|
||||
bufferInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
|
||||
bufferInfo.commandBufferCount = 1;
|
||||
check(vkAllocateCommandBuffers(out.device, &bufferInfo, &out.commandBuffer), "vkAllocateCommandBuffers");
|
||||
}
|
||||
|
||||
uint32_t find_memory_type(VkPhysicalDevice physicalDevice, uint32_t typeBits, VkMemoryPropertyFlags required,
|
||||
const char* what)
|
||||
{
|
||||
VkPhysicalDeviceMemoryProperties properties{};
|
||||
vkGetPhysicalDeviceMemoryProperties(physicalDevice, &properties);
|
||||
for (uint32_t index = 0; index < properties.memoryTypeCount; ++index) {
|
||||
const bool allowed = (typeBits & (1u << index)) != 0u;
|
||||
const bool suitable = (properties.memoryTypes[index].propertyFlags & required) == required;
|
||||
if (allowed && suitable) {
|
||||
return index;
|
||||
}
|
||||
}
|
||||
throw std::runtime_error(std::string("Aucun type de memoire compatible pour : ") + what);
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Objets GPU
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
void create_external_image(ImageResource& out, VkDevice device, const char* label)
|
||||
{
|
||||
out.device = device;
|
||||
|
||||
VkExternalMemoryImageCreateInfo externalInfo{};
|
||||
externalInfo.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO;
|
||||
externalInfo.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
VkImageCreateInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO;
|
||||
info.pNext = &externalInfo;
|
||||
info.imageType = VK_IMAGE_TYPE_2D;
|
||||
info.format = kFormat;
|
||||
info.extent.width = kWidth;
|
||||
info.extent.height = kHeight;
|
||||
info.extent.depth = 1;
|
||||
info.mipLevels = 1;
|
||||
info.arrayLayers = 1;
|
||||
info.samples = VK_SAMPLE_COUNT_1_BIT;
|
||||
info.tiling = VK_IMAGE_TILING_OPTIMAL;
|
||||
info.usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT;
|
||||
info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;
|
||||
info.initialLayout = VK_IMAGE_LAYOUT_UNDEFINED;
|
||||
|
||||
check(vkCreateImage(device, &info, nullptr, &out.image), label);
|
||||
}
|
||||
|
||||
void create_host_buffer(BufferResource& out, DeviceContext& context, VkDeviceSize size, const char* label)
|
||||
{
|
||||
out.device = context.device;
|
||||
out.size = size;
|
||||
|
||||
VkBufferCreateInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
|
||||
info.size = size;
|
||||
info.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT;
|
||||
info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;
|
||||
check(vkCreateBuffer(context.device, &info, nullptr, &out.buffer), label);
|
||||
|
||||
VkMemoryRequirements requirements{};
|
||||
vkGetBufferMemoryRequirements(context.device, out.buffer, &requirements);
|
||||
const uint32_t memoryType = find_memory_type(context.physicalDevice, requirements.memoryTypeBits,
|
||||
VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT |
|
||||
VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
|
||||
label);
|
||||
|
||||
VkMemoryAllocateInfo allocateInfo{};
|
||||
allocateInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
|
||||
allocateInfo.allocationSize = requirements.size;
|
||||
allocateInfo.memoryTypeIndex = memoryType;
|
||||
check(vkAllocateMemory(context.device, &allocateInfo, nullptr, &out.memory), label);
|
||||
check(vkBindBufferMemory(context.device, out.buffer, out.memory, 0), label);
|
||||
}
|
||||
|
||||
void* map_memory(VkDevice device, VkDeviceMemory memory, VkDeviceSize size, const char* label)
|
||||
{
|
||||
void* pointer = nullptr;
|
||||
check(vkMapMemory(device, memory, 0, size, 0, &pointer), label);
|
||||
return pointer;
|
||||
}
|
||||
|
||||
void flush_mapped_memory(VkDevice device, VkDeviceMemory memory)
|
||||
{
|
||||
VkMappedMemoryRange range{};
|
||||
range.sType = VK_STRUCTURE_TYPE_MAPPED_MEMORY_RANGE;
|
||||
range.memory = memory;
|
||||
range.offset = 0;
|
||||
range.size = VK_WHOLE_SIZE;
|
||||
check(vkFlushMappedMemoryRanges(device, 1, &range), "vkFlushMappedMemoryRanges");
|
||||
}
|
||||
|
||||
void invalidate_mapped_memory(VkDevice device, VkDeviceMemory memory)
|
||||
{
|
||||
VkMappedMemoryRange range{};
|
||||
range.sType = VK_STRUCTURE_TYPE_MAPPED_MEMORY_RANGE;
|
||||
range.memory = memory;
|
||||
range.offset = 0;
|
||||
range.size = VK_WHOLE_SIZE;
|
||||
check(vkInvalidateMappedMemoryRanges(device, 1, &range), "vkInvalidateMappedMemoryRanges");
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Commandes
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
void begin_commands(DeviceContext& context)
|
||||
{
|
||||
check(vkResetCommandPool(context.device, context.commandPool, 0), "vkResetCommandPool");
|
||||
VkCommandBufferBeginInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
|
||||
info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
|
||||
check(vkBeginCommandBuffer(context.commandBuffer, &info), "vkBeginCommandBuffer");
|
||||
}
|
||||
|
||||
void submit_and_wait(DeviceContext& context)
|
||||
{
|
||||
check(vkEndCommandBuffer(context.commandBuffer), "vkEndCommandBuffer");
|
||||
|
||||
VkSubmitInfo submit{};
|
||||
submit.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
|
||||
submit.commandBufferCount = 1;
|
||||
submit.pCommandBuffers = &context.commandBuffer;
|
||||
check(vkQueueSubmit(context.queue, 1, &submit, VK_NULL_HANDLE), "vkQueueSubmit");
|
||||
check(vkQueueWaitIdle(context.queue), "vkQueueWaitIdle");
|
||||
}
|
||||
|
||||
void cmd_image_barrier(VkCommandBuffer commandBuffer, VkImage image, VkImageLayout oldLayout, VkImageLayout newLayout,
|
||||
VkPipelineStageFlags sourceStage, VkAccessFlags sourceAccess, VkPipelineStageFlags dstStage,
|
||||
VkAccessFlags dstAccess)
|
||||
{
|
||||
VkImageMemoryBarrier barrier{};
|
||||
barrier.sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER;
|
||||
barrier.srcAccessMask = sourceAccess;
|
||||
barrier.dstAccessMask = dstAccess;
|
||||
barrier.oldLayout = oldLayout;
|
||||
barrier.newLayout = newLayout;
|
||||
barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.image = image;
|
||||
barrier.subresourceRange.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT;
|
||||
barrier.subresourceRange.baseMipLevel = 0;
|
||||
barrier.subresourceRange.levelCount = 1;
|
||||
barrier.subresourceRange.baseArrayLayer = 0;
|
||||
barrier.subresourceRange.layerCount = 1;
|
||||
|
||||
vkCmdPipelineBarrier(commandBuffer, sourceStage, dstStage, 0, 0, nullptr, 0, nullptr, 1, &barrier);
|
||||
}
|
||||
|
||||
void cmd_buffer_barrier(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize size,
|
||||
VkPipelineStageFlags sourceStage, VkAccessFlags sourceAccess, VkPipelineStageFlags dstStage,
|
||||
VkAccessFlags dstAccess)
|
||||
{
|
||||
VkBufferMemoryBarrier barrier{};
|
||||
barrier.sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER;
|
||||
barrier.srcAccessMask = sourceAccess;
|
||||
barrier.dstAccessMask = dstAccess;
|
||||
barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.buffer = buffer;
|
||||
barrier.offset = 0;
|
||||
barrier.size = size;
|
||||
|
||||
vkCmdPipelineBarrier(commandBuffer, sourceStage, dstStage, 0, 0, nullptr, 1, &barrier, 0, nullptr);
|
||||
}
|
||||
|
||||
VkBufferImageCopy full_image_copy()
|
||||
{
|
||||
VkBufferImageCopy region{};
|
||||
region.bufferOffset = 0;
|
||||
region.bufferRowLength = 0; // 0 = largeur de l'image, donnees serrees
|
||||
region.bufferImageHeight = 0; // 0 = hauteur de l'image
|
||||
region.imageSubresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT;
|
||||
region.imageSubresource.mipLevel = 0;
|
||||
region.imageSubresource.baseArrayLayer = 0;
|
||||
region.imageSubresource.layerCount = 1;
|
||||
region.imageOffset.x = 0;
|
||||
region.imageOffset.y = 0;
|
||||
region.imageOffset.z = 0;
|
||||
region.imageExtent.width = kWidth;
|
||||
region.imageExtent.height = kHeight;
|
||||
region.imageExtent.depth = 1;
|
||||
return region;
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Motif deterministe
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
std::vector<uint8_t> make_pattern()
|
||||
{
|
||||
std::vector<uint8_t> pixels(static_cast<size_t>(kImageBytes));
|
||||
for (uint32_t y = 0; y < kHeight; ++y) {
|
||||
for (uint32_t x = 0; x < kWidth; ++x) {
|
||||
const size_t offset = (static_cast<size_t>(y) * kWidth + x) * 4;
|
||||
pixels[offset + 0] = static_cast<uint8_t>(x & 0xFFu); // R
|
||||
pixels[offset + 1] = static_cast<uint8_t>(y & 0xFFu); // G
|
||||
pixels[offset + 2] = static_cast<uint8_t>((x ^ y) & 0xFFu); // B
|
||||
pixels[offset + 3] = 0xFFu; // A
|
||||
}
|
||||
}
|
||||
return pixels;
|
||||
}
|
||||
|
||||
size_t first_mismatch(const std::vector<uint8_t>& expected, const uint8_t* actual, size_t size)
|
||||
{
|
||||
for (size_t index = 0; index < size; ++index) {
|
||||
if (expected[index] != actual[index]) {
|
||||
return index;
|
||||
}
|
||||
}
|
||||
return size;
|
||||
}
|
||||
|
||||
void print_pattern_mismatch(size_t offset, size_t total)
|
||||
{
|
||||
std::printf(" PREMIERE DIFFERENCE a l'octet %llu sur %llu\n", static_cast<unsigned long long>(offset),
|
||||
static_cast<unsigned long long>(total));
|
||||
std::printf(" pixel (%llu, %llu), composante %llu\n",
|
||||
static_cast<unsigned long long>(offset / (4 * kWidth)),
|
||||
static_cast<unsigned long long>((offset % (4 * kWidth)) / 4),
|
||||
static_cast<unsigned long long>(offset % 4));
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Programme
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
int run()
|
||||
{
|
||||
std::printf("=== SEQUA - prototype gpu_interop_vk_export (etape 1) ===\n");
|
||||
std::printf("Image %ux%u R8G8B8A8_UNORM, %llu octets, DEVICE_LOCAL exportable\n\n", kWidth, kHeight,
|
||||
std::printf("Image %ux%u R8G8B8A8_UNORM, %llu octets, DEVICE_LOCAL exportable\n\n", kImageWidth, kImageHeight,
|
||||
static_cast<unsigned long long>(kImageBytes));
|
||||
|
||||
// --- Instance -----------------------------------------------------------
|
||||
@@ -757,86 +65,17 @@ int run()
|
||||
}
|
||||
std::printf(" correspondance deviceUUID : OK\n");
|
||||
|
||||
// --- Motif + staging ----------------------------------------------------
|
||||
// --- Motif + image exportable remplie -----------------------------------
|
||||
const std::vector<uint8_t> pattern = make_pattern();
|
||||
|
||||
BufferResource stagingA;
|
||||
create_host_buffer(stagingA, deviceA, kImageBytes, "staging A");
|
||||
{
|
||||
void* pointer = map_memory(deviceA.device, stagingA.memory, kImageBytes, "vkMapMemory (staging A)");
|
||||
std::memcpy(pointer, pattern.data(), pattern.size());
|
||||
flush_mapped_memory(deviceA.device, stagingA.memory);
|
||||
vkUnmapMemory(deviceA.device, stagingA.memory);
|
||||
}
|
||||
ExportableImage exportable;
|
||||
create_filled_exportable_image(exportable, deviceA, pattern);
|
||||
|
||||
// --- Image A exportable -------------------------------------------------
|
||||
ImageResource imageA;
|
||||
create_external_image(imageA, deviceA.device, "vkCreateImage (A, export)");
|
||||
|
||||
VkMemoryRequirements requirementsA{};
|
||||
vkGetImageMemoryRequirements(deviceA.device, imageA.image, &requirementsA);
|
||||
const uint32_t memoryTypeIndex =
|
||||
find_memory_type(deviceA.physicalDevice, requirementsA.memoryTypeBits,
|
||||
VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, "image A DEVICE_LOCAL exportable");
|
||||
|
||||
VkMemoryDedicatedAllocateInfo dedicatedA{};
|
||||
dedicatedA.sType = VK_STRUCTURE_TYPE_MEMORY_DEDICATED_ALLOCATE_INFO;
|
||||
dedicatedA.image = imageA.image;
|
||||
|
||||
VkExportMemoryAllocateInfo exportAllocateInfo{};
|
||||
exportAllocateInfo.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO;
|
||||
exportAllocateInfo.pNext = &dedicatedA;
|
||||
exportAllocateInfo.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
VkMemoryAllocateInfo allocateA{};
|
||||
allocateA.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
|
||||
allocateA.pNext = &exportAllocateInfo;
|
||||
allocateA.allocationSize = requirementsA.size;
|
||||
allocateA.memoryTypeIndex = memoryTypeIndex;
|
||||
check(vkAllocateMemory(deviceA.device, &allocateA, nullptr, &imageA.memory), "vkAllocateMemory (A, export)");
|
||||
check(vkBindImageMemory(deviceA.device, imageA.image, imageA.memory, 0), "vkBindImageMemory (A)");
|
||||
|
||||
std::printf("[3] Image A : memoire DEVICE_LOCAL exportable (type %u, %llu octets)\n", memoryTypeIndex,
|
||||
static_cast<unsigned long long>(requirementsA.size));
|
||||
|
||||
// --- Remplissage CPU -> GPU --------------------------------------------
|
||||
{
|
||||
const VkBufferImageCopy region = full_image_copy();
|
||||
begin_commands(deviceA);
|
||||
cmd_image_barrier(deviceA.commandBuffer, imageA.image, VK_IMAGE_LAYOUT_UNDEFINED,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, 0,
|
||||
VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_WRITE_BIT);
|
||||
vkCmdCopyBufferToImage(deviceA.commandBuffer, stagingA.buffer, imageA.image,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, ®ion);
|
||||
cmd_image_barrier(deviceA.commandBuffer, imageA.image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_READ_BIT);
|
||||
submit_and_wait(deviceA);
|
||||
}
|
||||
check(vkDeviceWaitIdle(deviceA.device), "vkDeviceWaitIdle (A)");
|
||||
std::printf("[3] Image A : memoire DEVICE_LOCAL exportable (type %u, %llu octets)\n", exportable.memoryTypeIndex,
|
||||
static_cast<unsigned long long>(exportable.allocationSize));
|
||||
std::printf("[4] Motif deterministe uploade (staging buffer + vkCmdCopyBufferToImage)\n");
|
||||
|
||||
// --- Export -------------------------------------------------------------
|
||||
const auto getWin32Handle = reinterpret_cast<PFN_vkGetMemoryWin32HandleKHR>(
|
||||
vkGetDeviceProcAddr(deviceA.device, "vkGetMemoryWin32HandleKHR"));
|
||||
if (getWin32Handle == nullptr) {
|
||||
throw std::runtime_error("vkGetMemoryWin32HandleKHR introuvable (extension non exposee).");
|
||||
}
|
||||
|
||||
VkMemoryGetWin32HandleInfoKHR handleInfo{};
|
||||
handleInfo.sType = VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR;
|
||||
handleInfo.memory = imageA.memory;
|
||||
handleInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
Win32Handle exported;
|
||||
const Clock::time_point exportStart = Clock::now();
|
||||
check(getWin32Handle(deviceA.device, &handleInfo, &exported.value), "vkGetMemoryWin32HandleKHR");
|
||||
const double exportMs = elapsed_ms(exportStart);
|
||||
|
||||
if (exported.value == nullptr) {
|
||||
throw std::runtime_error("vkGetMemoryWin32HandleKHR a retourne un HANDLE nul.");
|
||||
}
|
||||
std::printf("[5] Export : HANDLE Win32 0x%p en %.3f ms\n", static_cast<void*>(exported.value), exportMs);
|
||||
std::printf("[5] Export : HANDLE Win32 0x%p en %.3f ms\n",
|
||||
static_cast<void*>(exportable.handle.value), exportable.exportMs);
|
||||
|
||||
// --- Import dans le device B -------------------------------------------
|
||||
ImageResource imageB;
|
||||
@@ -849,13 +88,13 @@ int run()
|
||||
// on verifie seulement que l'image B accepte ce type.
|
||||
VkMemoryRequirements requirementsB{};
|
||||
vkGetImageMemoryRequirements(deviceB.device, imageB.image, &requirementsB);
|
||||
if ((requirementsB.memoryTypeBits & (1u << memoryTypeIndex)) == 0) {
|
||||
if ((requirementsB.memoryTypeBits & (1u << exportable.memoryTypeIndex)) == 0) {
|
||||
char message[192];
|
||||
std::snprintf(message, sizeof(message),
|
||||
"L'image B n'accepte pas le type de memoire exporte (%u).", memoryTypeIndex);
|
||||
std::snprintf(message, sizeof(message), "L'image B n'accepte pas le type de memoire exporte (%u).",
|
||||
exportable.memoryTypeIndex);
|
||||
throw std::runtime_error(message);
|
||||
}
|
||||
const uint32_t importMemoryType = memoryTypeIndex;
|
||||
const uint32_t importMemoryType = exportable.memoryTypeIndex;
|
||||
|
||||
VkMemoryDedicatedAllocateInfo dedicatedB{};
|
||||
dedicatedB.sType = VK_STRUCTURE_TYPE_MEMORY_DEDICATED_ALLOCATE_INFO;
|
||||
@@ -865,13 +104,13 @@ int run()
|
||||
importInfo.sType = VK_STRUCTURE_TYPE_IMPORT_MEMORY_WIN32_HANDLE_INFO_KHR;
|
||||
importInfo.pNext = &dedicatedB;
|
||||
importInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
importInfo.handle = exported.value;
|
||||
importInfo.handle = exportable.handle.value;
|
||||
importInfo.name = nullptr;
|
||||
|
||||
VkMemoryAllocateInfo allocateB{};
|
||||
allocateB.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
|
||||
allocateB.pNext = &importInfo;
|
||||
allocateB.allocationSize = requirementsA.size;
|
||||
allocateB.allocationSize = exportable.allocationSize;
|
||||
allocateB.memoryTypeIndex = importMemoryType;
|
||||
check(vkAllocateMemory(deviceB.device, &allocateB, nullptr, &imageB.memory), "vkAllocateMemory (B, import)");
|
||||
check(vkBindImageMemory(deviceB.device, imageB.image, imageB.memory, 0), "vkBindImageMemory (B)");
|
||||
@@ -899,7 +138,7 @@ int run()
|
||||
submit_and_wait(deviceB);
|
||||
}
|
||||
|
||||
size_t mismatch = kImageBytes;
|
||||
size_t mismatch = static_cast<size_t>(kImageBytes);
|
||||
{
|
||||
void* pointer = map_memory(deviceB.device, stagingB.memory, kImageBytes, "vkMapMemory (staging B)");
|
||||
invalidate_mapped_memory(deviceB.device, stagingB.memory);
|
||||
@@ -908,93 +147,32 @@ int run()
|
||||
}
|
||||
|
||||
const bool identical = mismatch == static_cast<size_t>(kImageBytes);
|
||||
std::printf("[7] Comparaison octet a octet apres import : %s\n",
|
||||
identical ? "IDENTIQUE" : "DIFFERENT");
|
||||
std::printf("[7] Comparaison octet a octet apres import : %s\n", identical ? "IDENTIQUE" : "DIFFERENT");
|
||||
if (!identical) {
|
||||
print_pattern_mismatch(mismatch, static_cast<size_t>(kImageBytes));
|
||||
}
|
||||
|
||||
// Le HANDLE NT a ete duplique par le driver lors de vkAllocateMemory : on peut le fermer.
|
||||
const bool handleClosed = CloseHandle(exported.value) != FALSE;
|
||||
exported.value = nullptr;
|
||||
const bool handleClosed = exportable.handle.close();
|
||||
std::printf("[8] CloseHandle du HANDLE exporte : %s\n", handleClosed ? "OK" : "ECHEC");
|
||||
|
||||
// --- Reference : aller-retour GPU -> CPU -> GPU ------------------------
|
||||
std::vector<uint8_t> hostPixels(static_cast<size_t>(kImageBytes));
|
||||
std::vector<double> roundtripSamples;
|
||||
bool roundtripOk = true;
|
||||
|
||||
for (int iteration = 0; iteration < kRoundtripIterations; ++iteration) {
|
||||
const Clock::time_point roundtripStart = Clock::now();
|
||||
|
||||
// GPU -> CPU
|
||||
{
|
||||
const VkBufferImageCopy region = full_image_copy();
|
||||
begin_commands(deviceA);
|
||||
cmd_image_barrier(deviceA.commandBuffer, imageA.image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, 0,
|
||||
VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_READ_BIT);
|
||||
vkCmdCopyImageToBuffer(deviceA.commandBuffer, imageA.image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL,
|
||||
stagingA.buffer, 1, ®ion);
|
||||
cmd_buffer_barrier(deviceA.commandBuffer, stagingA.buffer, kImageBytes, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_HOST_BIT, VK_ACCESS_HOST_READ_BIT);
|
||||
submit_and_wait(deviceA);
|
||||
}
|
||||
|
||||
{
|
||||
void* pointer = map_memory(deviceA.device, stagingA.memory, kImageBytes, "vkMapMemory (relecture A)");
|
||||
invalidate_mapped_memory(deviceA.device, stagingA.memory);
|
||||
std::memcpy(hostPixels.data(), pointer, hostPixels.size());
|
||||
vkUnmapMemory(deviceA.device, stagingA.memory);
|
||||
}
|
||||
|
||||
if (iteration == 0) {
|
||||
roundtripOk = first_mismatch(pattern, hostPixels.data(), hostPixels.size()) == hostPixels.size();
|
||||
}
|
||||
|
||||
// CPU -> GPU
|
||||
{
|
||||
void* pointer = map_memory(deviceA.device, stagingA.memory, kImageBytes, "vkMapMemory (renvoi A)");
|
||||
std::memcpy(pointer, hostPixels.data(), hostPixels.size());
|
||||
flush_mapped_memory(deviceA.device, stagingA.memory);
|
||||
vkUnmapMemory(deviceA.device, stagingA.memory);
|
||||
|
||||
const VkBufferImageCopy region = full_image_copy();
|
||||
begin_commands(deviceA);
|
||||
cmd_buffer_barrier(deviceA.commandBuffer, stagingA.buffer, kImageBytes, VK_PIPELINE_STAGE_HOST_BIT,
|
||||
VK_ACCESS_HOST_WRITE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_READ_BIT);
|
||||
cmd_image_barrier(deviceA.commandBuffer, imageA.image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, VK_PIPELINE_STAGE_TRANSFER_BIT, 0,
|
||||
VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_WRITE_BIT);
|
||||
vkCmdCopyBufferToImage(deviceA.commandBuffer, stagingA.buffer, imageA.image,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, ®ion);
|
||||
cmd_image_barrier(deviceA.commandBuffer, imageA.image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_READ_BIT);
|
||||
submit_and_wait(deviceA);
|
||||
}
|
||||
|
||||
roundtripSamples.push_back(elapsed_ms(roundtripStart));
|
||||
}
|
||||
|
||||
std::sort(roundtripSamples.begin(), roundtripSamples.end());
|
||||
const double roundtripMin = roundtripSamples.front();
|
||||
const double roundtripMedian = roundtripSamples[roundtripSamples.size() / 2];
|
||||
const RoundtripResult roundtrip =
|
||||
measure_cpu_roundtrip(deviceA, exportable.image.image, exportable.staging, pattern, kRoundtripIterations);
|
||||
|
||||
std::printf("[9] Aller-retour GPU->CPU->GPU de reference (%d essais)\n", kRoundtripIterations);
|
||||
for (size_t index = 0; index < roundtripSamples.size(); ++index) {
|
||||
for (size_t index = 0; index < roundtrip.samplesMs.size(); ++index) {
|
||||
std::printf(" echantillon %llu (trie) : %.3f ms\n", static_cast<unsigned long long>(index + 1),
|
||||
roundtripSamples[index]);
|
||||
roundtrip.samplesMs[index]);
|
||||
}
|
||||
std::printf(" min %.3f ms, mediane %.3f ms, contenu %s\n", roundtripMin, roundtripMedian,
|
||||
roundtripOk ? "identique" : "DIFFERENT");
|
||||
std::printf(" min %.3f ms, mediane %.3f ms, contenu %s\n", roundtrip.minMs, roundtrip.medianMs,
|
||||
roundtrip.contentIdentical ? "identique" : "DIFFERENT");
|
||||
|
||||
// --- Resume -------------------------------------------------------------
|
||||
std::printf("\n=== RESULTATS ===\n");
|
||||
std::printf(" Export HANDLE Win32 : %8.3f ms\n", exportMs);
|
||||
std::printf(" Export HANDLE Win32 : %8.3f ms\n", exportable.exportMs);
|
||||
std::printf(" Import dans un 2e device : %8.3f ms\n", importMs);
|
||||
std::printf(" Aller-retour GPU->CPU->GPU : %8.3f ms (mediane de %d essais)\n", roundtripMedian,
|
||||
std::printf(" Aller-retour GPU->CPU->GPU : %8.3f ms (mediane de %d essais)\n", roundtrip.medianMs,
|
||||
kRoundtripIterations);
|
||||
std::printf(" Verification octet a octet : %s\n", identical ? "OK" : "ECHEC");
|
||||
|
||||
|
||||
@@ -0,0 +1,806 @@
|
||||
// SPDX-License-Identifier: MPL-2.0
|
||||
//
|
||||
// PROTOTYPE JETABLE - pas du code produit.
|
||||
//
|
||||
// Implementation du code Vulkan commun aux etapes 1 et 2 du prototype gpu_interop.
|
||||
|
||||
#include "vk_common.hpp"
|
||||
|
||||
#include <algorithm>
|
||||
#include <cstdio>
|
||||
#include <cstring>
|
||||
#include <iterator>
|
||||
#include <stdexcept>
|
||||
|
||||
namespace gpu_interop {
|
||||
|
||||
namespace {
|
||||
|
||||
std::vector<VkExtensionProperties> enumerate_instance_extensions()
|
||||
{
|
||||
uint32_t count = 0;
|
||||
check(vkEnumerateInstanceExtensionProperties(nullptr, &count, nullptr), "vkEnumerateInstanceExtensionProperties");
|
||||
std::vector<VkExtensionProperties> extensions(count);
|
||||
if (count != 0) {
|
||||
check(vkEnumerateInstanceExtensionProperties(nullptr, &count, extensions.data()),
|
||||
"vkEnumerateInstanceExtensionProperties");
|
||||
}
|
||||
return extensions;
|
||||
}
|
||||
|
||||
std::vector<VkExtensionProperties> enumerate_device_extensions(VkPhysicalDevice physicalDevice)
|
||||
{
|
||||
uint32_t count = 0;
|
||||
check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, nullptr),
|
||||
"vkEnumerateDeviceExtensionProperties");
|
||||
std::vector<VkExtensionProperties> extensions(count);
|
||||
if (count != 0) {
|
||||
check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, extensions.data()),
|
||||
"vkEnumerateDeviceExtensionProperties");
|
||||
}
|
||||
return extensions;
|
||||
}
|
||||
|
||||
bool has_extension(const std::vector<VkExtensionProperties>& extensions, const char* name)
|
||||
{
|
||||
for (const VkExtensionProperties& extension : extensions) {
|
||||
if (std::strcmp(extension.extensionName, name) == 0) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Erreurs et formatage
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
const char* result_name(VkResult result)
|
||||
{
|
||||
switch (result) {
|
||||
case VK_SUCCESS:
|
||||
return "VK_SUCCESS";
|
||||
case VK_NOT_READY:
|
||||
return "VK_NOT_READY";
|
||||
case VK_TIMEOUT:
|
||||
return "VK_TIMEOUT";
|
||||
case VK_EVENT_SET:
|
||||
return "VK_EVENT_SET";
|
||||
case VK_EVENT_RESET:
|
||||
return "VK_EVENT_RESET";
|
||||
case VK_INCOMPLETE:
|
||||
return "VK_INCOMPLETE";
|
||||
case VK_ERROR_OUT_OF_HOST_MEMORY:
|
||||
return "VK_ERROR_OUT_OF_HOST_MEMORY";
|
||||
case VK_ERROR_OUT_OF_DEVICE_MEMORY:
|
||||
return "VK_ERROR_OUT_OF_DEVICE_MEMORY";
|
||||
case VK_ERROR_INITIALIZATION_FAILED:
|
||||
return "VK_ERROR_INITIALIZATION_FAILED";
|
||||
case VK_ERROR_DEVICE_LOST:
|
||||
return "VK_ERROR_DEVICE_LOST";
|
||||
case VK_ERROR_MEMORY_MAP_FAILED:
|
||||
return "VK_ERROR_MEMORY_MAP_FAILED";
|
||||
case VK_ERROR_LAYER_NOT_PRESENT:
|
||||
return "VK_ERROR_LAYER_NOT_PRESENT";
|
||||
case VK_ERROR_EXTENSION_NOT_PRESENT:
|
||||
return "VK_ERROR_EXTENSION_NOT_PRESENT";
|
||||
case VK_ERROR_FEATURE_NOT_PRESENT:
|
||||
return "VK_ERROR_FEATURE_NOT_PRESENT";
|
||||
case VK_ERROR_INCOMPATIBLE_DRIVER:
|
||||
return "VK_ERROR_INCOMPATIBLE_DRIVER";
|
||||
case VK_ERROR_TOO_MANY_OBJECTS:
|
||||
return "VK_ERROR_TOO_MANY_OBJECTS";
|
||||
case VK_ERROR_FORMAT_NOT_SUPPORTED:
|
||||
return "VK_ERROR_FORMAT_NOT_SUPPORTED";
|
||||
case VK_ERROR_INVALID_EXTERNAL_HANDLE:
|
||||
return "VK_ERROR_INVALID_EXTERNAL_HANDLE";
|
||||
case VK_ERROR_OUT_OF_DATE_KHR:
|
||||
return "VK_ERROR_OUT_OF_DATE_KHR";
|
||||
default:
|
||||
return "VkResult inconnu";
|
||||
}
|
||||
}
|
||||
|
||||
[[noreturn]] void throw_vk(const char* what, VkResult result)
|
||||
{
|
||||
char buffer[256];
|
||||
std::snprintf(buffer, sizeof(buffer), "%s : %s (%d)", what, result_name(result), static_cast<int>(result));
|
||||
throw std::runtime_error(buffer);
|
||||
}
|
||||
|
||||
void check(VkResult result, const char* what)
|
||||
{
|
||||
if (result != VK_SUCCESS) {
|
||||
throw_vk(what, result);
|
||||
}
|
||||
}
|
||||
|
||||
double elapsed_ms(Clock::time_point start)
|
||||
{
|
||||
return std::chrono::duration<double, std::milli>(Clock::now() - start).count();
|
||||
}
|
||||
|
||||
const char* device_type_name(VkPhysicalDeviceType type)
|
||||
{
|
||||
switch (type) {
|
||||
case VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU:
|
||||
return "discrete";
|
||||
case VK_PHYSICAL_DEVICE_TYPE_INTEGRATED_GPU:
|
||||
return "integrated";
|
||||
case VK_PHYSICAL_DEVICE_TYPE_VIRTUAL_GPU:
|
||||
return "virtual";
|
||||
case VK_PHYSICAL_DEVICE_TYPE_CPU:
|
||||
return "cpu";
|
||||
default:
|
||||
return "other";
|
||||
}
|
||||
}
|
||||
|
||||
std::string uuid_to_string(const std::array<uint8_t, VK_UUID_SIZE>& uuid)
|
||||
{
|
||||
std::string text;
|
||||
text.reserve(uuid.size() * 2);
|
||||
char byte[4];
|
||||
for (uint8_t value : uuid) {
|
||||
std::snprintf(byte, sizeof(byte), "%02x", static_cast<unsigned int>(value));
|
||||
text += byte;
|
||||
}
|
||||
return text;
|
||||
}
|
||||
|
||||
void print_external_format_support(VkPhysicalDevice physicalDevice)
|
||||
{
|
||||
VkPhysicalDeviceExternalImageFormatInfo externalInfo{};
|
||||
externalInfo.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_EXTERNAL_IMAGE_FORMAT_INFO;
|
||||
externalInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
VkPhysicalDeviceImageFormatInfo2 formatInfo{};
|
||||
formatInfo.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGE_FORMAT_INFO_2;
|
||||
formatInfo.pNext = &externalInfo;
|
||||
formatInfo.format = kImageFormat;
|
||||
formatInfo.type = VK_IMAGE_TYPE_2D;
|
||||
formatInfo.tiling = VK_IMAGE_TILING_OPTIMAL;
|
||||
formatInfo.usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT;
|
||||
|
||||
VkExternalImageFormatProperties externalProperties{};
|
||||
externalProperties.sType = VK_STRUCTURE_TYPE_EXTERNAL_IMAGE_FORMAT_PROPERTIES;
|
||||
VkImageFormatProperties2 properties{};
|
||||
properties.sType = VK_STRUCTURE_TYPE_IMAGE_FORMAT_PROPERTIES_2;
|
||||
properties.pNext = &externalProperties;
|
||||
|
||||
const VkResult result = vkGetPhysicalDeviceImageFormatProperties2(physicalDevice, &formatInfo, &properties);
|
||||
if (result != VK_SUCCESS) {
|
||||
std::printf(" VkExternalImageFormatProperties 2D/R8G8B8A8_UNORM/OPAQUE_WIN32 : %s\n", result_name(result));
|
||||
return;
|
||||
}
|
||||
|
||||
const VkExternalMemoryFeatureFlags features = externalProperties.externalMemoryProperties.externalMemoryFeatures;
|
||||
std::printf(" VkExternalImageFormatProperties 2D/R8G8B8A8_UNORM/OPAQUE_WIN32 : exportable=%s importable=%s "
|
||||
"dedicated_only=%s\n",
|
||||
(features & VK_EXTERNAL_MEMORY_FEATURE_EXPORTABLE_BIT) != 0u ? "oui" : "non",
|
||||
(features & VK_EXTERNAL_MEMORY_FEATURE_IMPORTABLE_BIT) != 0u ? "oui" : "non",
|
||||
(features & VK_EXTERNAL_MEMORY_FEATURE_DEDICATED_ONLY_BIT) != 0u ? "oui" : "non");
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// RAII minimal
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
Win32Handle::~Win32Handle()
|
||||
{
|
||||
if (value != nullptr) {
|
||||
(void)CloseHandle(value);
|
||||
}
|
||||
}
|
||||
|
||||
bool Win32Handle::close()
|
||||
{
|
||||
if (value == nullptr) {
|
||||
return false;
|
||||
}
|
||||
const bool closed = CloseHandle(value) != FALSE;
|
||||
value = nullptr;
|
||||
return closed;
|
||||
}
|
||||
|
||||
InstanceContext::~InstanceContext()
|
||||
{
|
||||
if (instance != VK_NULL_HANDLE) {
|
||||
vkDestroyInstance(instance, nullptr);
|
||||
}
|
||||
}
|
||||
|
||||
DeviceContext::~DeviceContext()
|
||||
{
|
||||
if (device != VK_NULL_HANDLE) {
|
||||
(void)vkDeviceWaitIdle(device);
|
||||
if (commandPool != VK_NULL_HANDLE) {
|
||||
vkDestroyCommandPool(device, commandPool, nullptr);
|
||||
}
|
||||
vkDestroyDevice(device, nullptr);
|
||||
}
|
||||
}
|
||||
|
||||
ImageResource::~ImageResource()
|
||||
{
|
||||
if (device != VK_NULL_HANDLE) {
|
||||
if (image != VK_NULL_HANDLE) {
|
||||
vkDestroyImage(device, image, nullptr);
|
||||
}
|
||||
if (memory != VK_NULL_HANDLE) {
|
||||
vkFreeMemory(device, memory, nullptr);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
BufferResource::~BufferResource()
|
||||
{
|
||||
if (device != VK_NULL_HANDLE) {
|
||||
if (buffer != VK_NULL_HANDLE) {
|
||||
vkDestroyBuffer(device, buffer, nullptr);
|
||||
}
|
||||
if (memory != VK_NULL_HANDLE) {
|
||||
vkFreeMemory(device, memory, nullptr);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Enumeration / creation
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
void create_instance(InstanceContext& out)
|
||||
{
|
||||
uint32_t loaderVersion = 0;
|
||||
check(vkEnumerateInstanceVersion(&loaderVersion), "vkEnumerateInstanceVersion");
|
||||
if (loaderVersion < VK_API_VERSION_1_2) {
|
||||
char buffer[160];
|
||||
std::snprintf(buffer, sizeof(buffer), "Vulkan loader trop ancien : %u.%u.%u (1.2 minimum requis)",
|
||||
static_cast<unsigned int>(VK_API_VERSION_MAJOR(loaderVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_MINOR(loaderVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_PATCH(loaderVersion)));
|
||||
throw std::runtime_error(buffer);
|
||||
}
|
||||
|
||||
const std::vector<VkExtensionProperties> available = enumerate_instance_extensions();
|
||||
std::vector<const char*> enabled;
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME);
|
||||
}
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME);
|
||||
}
|
||||
|
||||
VkApplicationInfo application{};
|
||||
application.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO;
|
||||
application.pApplicationName = "gpu_interop";
|
||||
application.applicationVersion = VK_MAKE_API_VERSION(0, 0, 2, 0);
|
||||
application.pEngineName = "SEQUA (prototype)";
|
||||
application.engineVersion = VK_MAKE_API_VERSION(0, 0, 1, 0);
|
||||
application.apiVersion = VK_API_VERSION_1_2;
|
||||
|
||||
VkInstanceCreateInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO;
|
||||
info.pApplicationInfo = &application;
|
||||
info.enabledExtensionCount = static_cast<uint32_t>(enabled.size());
|
||||
info.ppEnabledExtensionNames = enabled.data();
|
||||
|
||||
check(vkCreateInstance(&info, nullptr, &out.instance), "vkCreateInstance");
|
||||
}
|
||||
|
||||
std::array<uint8_t, VK_UUID_SIZE> query_device_uuid(VkPhysicalDevice physicalDevice)
|
||||
{
|
||||
VkPhysicalDeviceIDProperties idProperties{};
|
||||
idProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES;
|
||||
VkPhysicalDeviceProperties2 properties2{};
|
||||
properties2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2;
|
||||
properties2.pNext = &idProperties;
|
||||
vkGetPhysicalDeviceProperties2(physicalDevice, &properties2);
|
||||
|
||||
std::array<uint8_t, VK_UUID_SIZE> uuid{};
|
||||
std::copy(std::begin(idProperties.deviceUUID), std::end(idProperties.deviceUUID), uuid.begin());
|
||||
return uuid;
|
||||
}
|
||||
|
||||
DeviceInfo pick_physical_device(VkInstance instance)
|
||||
{
|
||||
uint32_t count = 0;
|
||||
check(vkEnumeratePhysicalDevices(instance, &count, nullptr), "vkEnumeratePhysicalDevices");
|
||||
if (count == 0) {
|
||||
throw std::runtime_error("Aucun GPU Vulkan detecte.");
|
||||
}
|
||||
std::vector<VkPhysicalDevice> devices(count);
|
||||
check(vkEnumeratePhysicalDevices(instance, &count, devices.data()), "vkEnumeratePhysicalDevices");
|
||||
|
||||
std::string rejected;
|
||||
for (VkPhysicalDevice physicalDevice : devices) {
|
||||
DeviceInfo info{};
|
||||
info.physicalDevice = physicalDevice;
|
||||
|
||||
VkPhysicalDeviceProperties2 properties2{};
|
||||
properties2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2;
|
||||
vkGetPhysicalDeviceProperties2(physicalDevice, &properties2);
|
||||
info.properties = properties2.properties;
|
||||
info.uuid = query_device_uuid(physicalDevice);
|
||||
|
||||
const std::vector<VkExtensionProperties> extensions = enumerate_device_extensions(physicalDevice);
|
||||
info.externalMemory = has_extension(extensions, VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME);
|
||||
info.externalMemoryWin32 = has_extension(extensions, VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME);
|
||||
info.externalSemaphore = has_extension(extensions, VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME);
|
||||
info.externalSemaphoreWin32 = has_extension(extensions, VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME);
|
||||
|
||||
uint32_t familyCount = 0;
|
||||
vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &familyCount, nullptr);
|
||||
std::vector<VkQueueFamilyProperties> families(familyCount);
|
||||
if (familyCount != 0) {
|
||||
vkGetPhysicalDeviceQueueFamilyProperties(physicalDevice, &familyCount, families.data());
|
||||
}
|
||||
bool hasGraphicsQueue = false;
|
||||
for (uint32_t index = 0; index < familyCount; ++index) {
|
||||
if ((families[index].queueFlags & VK_QUEUE_GRAPHICS_BIT) != 0u) {
|
||||
info.queueFamily = index;
|
||||
hasGraphicsQueue = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
std::printf(" GPU : %s (%s, API %u.%u.%u) - external_memory_win32=%s\n", info.properties.deviceName,
|
||||
device_type_name(info.properties.deviceType),
|
||||
static_cast<unsigned int>(VK_API_VERSION_MAJOR(info.properties.apiVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_MINOR(info.properties.apiVersion)),
|
||||
static_cast<unsigned int>(VK_API_VERSION_PATCH(info.properties.apiVersion)),
|
||||
info.externalMemoryWin32 ? "oui" : "non");
|
||||
|
||||
if (hasGraphicsQueue && info.externalMemory && info.externalMemoryWin32) {
|
||||
return info;
|
||||
}
|
||||
rejected += " - ignore : ";
|
||||
rejected += info.properties.deviceName;
|
||||
rejected += hasGraphicsQueue ? "" : " (pas de file graphique)";
|
||||
rejected += info.externalMemoryWin32 ? "" : " (VK_KHR_external_memory_win32 absent)";
|
||||
rejected += '\n';
|
||||
}
|
||||
|
||||
throw std::runtime_error("Aucun GPU ne remplit les conditions (file graphique + VK_KHR_external_memory_win32).\n" +
|
||||
rejected);
|
||||
}
|
||||
|
||||
void create_device(const DeviceInfo& info, bool enableExternalSemaphore, DeviceContext& out, const char* label)
|
||||
{
|
||||
const std::vector<VkExtensionProperties> available = enumerate_device_extensions(info.physicalDevice);
|
||||
std::vector<const char*> enabled;
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME);
|
||||
}
|
||||
enabled.push_back(VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME);
|
||||
if (enableExternalSemaphore) {
|
||||
if (has_extension(available, VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME)) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME);
|
||||
}
|
||||
if (info.externalSemaphoreWin32) {
|
||||
enabled.push_back(VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME);
|
||||
}
|
||||
}
|
||||
|
||||
const float priority = 1.0f;
|
||||
VkDeviceQueueCreateInfo queueInfo{};
|
||||
queueInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO;
|
||||
queueInfo.queueFamilyIndex = info.queueFamily;
|
||||
queueInfo.queueCount = 1;
|
||||
queueInfo.pQueuePriorities = &priority;
|
||||
|
||||
VkDeviceCreateInfo createInfo{};
|
||||
createInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO;
|
||||
createInfo.queueCreateInfoCount = 1;
|
||||
createInfo.pQueueCreateInfos = &queueInfo;
|
||||
createInfo.enabledExtensionCount = static_cast<uint32_t>(enabled.size());
|
||||
createInfo.ppEnabledExtensionNames = enabled.data();
|
||||
|
||||
check(vkCreateDevice(info.physicalDevice, &createInfo, nullptr, &out.device), label);
|
||||
out.physicalDevice = info.physicalDevice;
|
||||
out.queueFamily = info.queueFamily;
|
||||
vkGetDeviceQueue(out.device, info.queueFamily, 0, &out.queue);
|
||||
|
||||
VkCommandPoolCreateInfo poolInfo{};
|
||||
poolInfo.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO;
|
||||
poolInfo.flags = VK_COMMAND_POOL_CREATE_RESET_COMMAND_BUFFER_BIT;
|
||||
poolInfo.queueFamilyIndex = info.queueFamily;
|
||||
check(vkCreateCommandPool(out.device, &poolInfo, nullptr, &out.commandPool), "vkCreateCommandPool");
|
||||
|
||||
VkCommandBufferAllocateInfo bufferInfo{};
|
||||
bufferInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO;
|
||||
bufferInfo.commandPool = out.commandPool;
|
||||
bufferInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY;
|
||||
bufferInfo.commandBufferCount = 1;
|
||||
check(vkAllocateCommandBuffers(out.device, &bufferInfo, &out.commandBuffer), "vkAllocateCommandBuffers");
|
||||
}
|
||||
|
||||
uint32_t find_memory_type(VkPhysicalDevice physicalDevice, uint32_t typeBits, VkMemoryPropertyFlags required,
|
||||
const char* what)
|
||||
{
|
||||
VkPhysicalDeviceMemoryProperties properties{};
|
||||
vkGetPhysicalDeviceMemoryProperties(physicalDevice, &properties);
|
||||
for (uint32_t index = 0; index < properties.memoryTypeCount; ++index) {
|
||||
const bool allowed = (typeBits & (1u << index)) != 0u;
|
||||
const bool suitable = (properties.memoryTypes[index].propertyFlags & required) == required;
|
||||
if (allowed && suitable) {
|
||||
return index;
|
||||
}
|
||||
}
|
||||
throw std::runtime_error(std::string("Aucun type de memoire compatible pour : ") + what);
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Objets GPU
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
void create_external_image(ImageResource& out, VkDevice device, const char* label)
|
||||
{
|
||||
out.device = device;
|
||||
|
||||
VkExternalMemoryImageCreateInfo externalInfo{};
|
||||
externalInfo.sType = VK_STRUCTURE_TYPE_EXTERNAL_MEMORY_IMAGE_CREATE_INFO;
|
||||
externalInfo.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
VkImageCreateInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO;
|
||||
info.pNext = &externalInfo;
|
||||
info.imageType = VK_IMAGE_TYPE_2D;
|
||||
info.format = kImageFormat;
|
||||
info.extent.width = kImageWidth;
|
||||
info.extent.height = kImageHeight;
|
||||
info.extent.depth = 1;
|
||||
info.mipLevels = 1;
|
||||
info.arrayLayers = 1;
|
||||
info.samples = VK_SAMPLE_COUNT_1_BIT;
|
||||
info.tiling = VK_IMAGE_TILING_OPTIMAL;
|
||||
info.usage = VK_IMAGE_USAGE_TRANSFER_SRC_BIT | VK_IMAGE_USAGE_TRANSFER_DST_BIT;
|
||||
info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;
|
||||
info.initialLayout = VK_IMAGE_LAYOUT_UNDEFINED;
|
||||
|
||||
check(vkCreateImage(device, &info, nullptr, &out.image), label);
|
||||
}
|
||||
|
||||
void create_host_buffer(BufferResource& out, DeviceContext& context, VkDeviceSize size, const char* label)
|
||||
{
|
||||
out.device = context.device;
|
||||
out.size = size;
|
||||
|
||||
VkBufferCreateInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO;
|
||||
info.size = size;
|
||||
info.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT;
|
||||
info.sharingMode = VK_SHARING_MODE_EXCLUSIVE;
|
||||
check(vkCreateBuffer(context.device, &info, nullptr, &out.buffer), label);
|
||||
|
||||
VkMemoryRequirements requirements{};
|
||||
vkGetBufferMemoryRequirements(context.device, out.buffer, &requirements);
|
||||
const uint32_t memoryType = find_memory_type(context.physicalDevice, requirements.memoryTypeBits,
|
||||
VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT |
|
||||
VK_MEMORY_PROPERTY_HOST_COHERENT_BIT,
|
||||
label);
|
||||
|
||||
VkMemoryAllocateInfo allocateInfo{};
|
||||
allocateInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
|
||||
allocateInfo.allocationSize = requirements.size;
|
||||
allocateInfo.memoryTypeIndex = memoryType;
|
||||
check(vkAllocateMemory(context.device, &allocateInfo, nullptr, &out.memory), label);
|
||||
check(vkBindBufferMemory(context.device, out.buffer, out.memory, 0), label);
|
||||
}
|
||||
|
||||
void* map_memory(VkDevice device, VkDeviceMemory memory, VkDeviceSize size, const char* label)
|
||||
{
|
||||
void* pointer = nullptr;
|
||||
check(vkMapMemory(device, memory, 0, size, 0, &pointer), label);
|
||||
return pointer;
|
||||
}
|
||||
|
||||
void flush_mapped_memory(VkDevice device, VkDeviceMemory memory)
|
||||
{
|
||||
VkMappedMemoryRange range{};
|
||||
range.sType = VK_STRUCTURE_TYPE_MAPPED_MEMORY_RANGE;
|
||||
range.memory = memory;
|
||||
range.offset = 0;
|
||||
range.size = VK_WHOLE_SIZE;
|
||||
check(vkFlushMappedMemoryRanges(device, 1, &range), "vkFlushMappedMemoryRanges");
|
||||
}
|
||||
|
||||
void invalidate_mapped_memory(VkDevice device, VkDeviceMemory memory)
|
||||
{
|
||||
VkMappedMemoryRange range{};
|
||||
range.sType = VK_STRUCTURE_TYPE_MAPPED_MEMORY_RANGE;
|
||||
range.memory = memory;
|
||||
range.offset = 0;
|
||||
range.size = VK_WHOLE_SIZE;
|
||||
check(vkInvalidateMappedMemoryRanges(device, 1, &range), "vkInvalidateMappedMemoryRanges");
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Commandes
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
void begin_commands(DeviceContext& context)
|
||||
{
|
||||
check(vkResetCommandPool(context.device, context.commandPool, 0), "vkResetCommandPool");
|
||||
VkCommandBufferBeginInfo info{};
|
||||
info.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO;
|
||||
info.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT;
|
||||
check(vkBeginCommandBuffer(context.commandBuffer, &info), "vkBeginCommandBuffer");
|
||||
}
|
||||
|
||||
void submit_and_wait(DeviceContext& context)
|
||||
{
|
||||
check(vkEndCommandBuffer(context.commandBuffer), "vkEndCommandBuffer");
|
||||
|
||||
VkSubmitInfo submit{};
|
||||
submit.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
|
||||
submit.commandBufferCount = 1;
|
||||
submit.pCommandBuffers = &context.commandBuffer;
|
||||
check(vkQueueSubmit(context.queue, 1, &submit, VK_NULL_HANDLE), "vkQueueSubmit");
|
||||
check(vkQueueWaitIdle(context.queue), "vkQueueWaitIdle");
|
||||
}
|
||||
|
||||
void cmd_image_barrier(VkCommandBuffer commandBuffer, VkImage image, VkImageLayout oldLayout, VkImageLayout newLayout,
|
||||
VkPipelineStageFlags sourceStage, VkAccessFlags sourceAccess, VkPipelineStageFlags dstStage,
|
||||
VkAccessFlags dstAccess)
|
||||
{
|
||||
VkImageMemoryBarrier barrier{};
|
||||
barrier.sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER;
|
||||
barrier.srcAccessMask = sourceAccess;
|
||||
barrier.dstAccessMask = dstAccess;
|
||||
barrier.oldLayout = oldLayout;
|
||||
barrier.newLayout = newLayout;
|
||||
barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.image = image;
|
||||
barrier.subresourceRange.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT;
|
||||
barrier.subresourceRange.baseMipLevel = 0;
|
||||
barrier.subresourceRange.levelCount = 1;
|
||||
barrier.subresourceRange.baseArrayLayer = 0;
|
||||
barrier.subresourceRange.layerCount = 1;
|
||||
|
||||
vkCmdPipelineBarrier(commandBuffer, sourceStage, dstStage, 0, 0, nullptr, 0, nullptr, 1, &barrier);
|
||||
}
|
||||
|
||||
void cmd_buffer_barrier(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize size,
|
||||
VkPipelineStageFlags sourceStage, VkAccessFlags sourceAccess, VkPipelineStageFlags dstStage,
|
||||
VkAccessFlags dstAccess)
|
||||
{
|
||||
VkBufferMemoryBarrier barrier{};
|
||||
barrier.sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER;
|
||||
barrier.srcAccessMask = sourceAccess;
|
||||
barrier.dstAccessMask = dstAccess;
|
||||
barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED;
|
||||
barrier.buffer = buffer;
|
||||
barrier.offset = 0;
|
||||
barrier.size = size;
|
||||
|
||||
vkCmdPipelineBarrier(commandBuffer, sourceStage, dstStage, 0, 0, nullptr, 1, &barrier, 0, nullptr);
|
||||
}
|
||||
|
||||
VkBufferImageCopy full_image_copy()
|
||||
{
|
||||
VkBufferImageCopy region{};
|
||||
region.bufferOffset = 0;
|
||||
region.bufferRowLength = 0; // 0 = largeur de l'image, donnees serrees
|
||||
region.bufferImageHeight = 0; // 0 = hauteur de l'image
|
||||
region.imageSubresource.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT;
|
||||
region.imageSubresource.mipLevel = 0;
|
||||
region.imageSubresource.baseArrayLayer = 0;
|
||||
region.imageSubresource.layerCount = 1;
|
||||
region.imageOffset.x = 0;
|
||||
region.imageOffset.y = 0;
|
||||
region.imageOffset.z = 0;
|
||||
region.imageExtent.width = kImageWidth;
|
||||
region.imageExtent.height = kImageHeight;
|
||||
region.imageExtent.depth = 1;
|
||||
return region;
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Motif deterministe
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
std::vector<uint8_t> make_pattern()
|
||||
{
|
||||
std::vector<uint8_t> pixels(static_cast<size_t>(kImageBytes));
|
||||
for (uint32_t y = 0; y < kImageHeight; ++y) {
|
||||
for (uint32_t x = 0; x < kImageWidth; ++x) {
|
||||
const size_t offset = (static_cast<size_t>(y) * kImageWidth + x) * 4;
|
||||
pixels[offset + 0] = static_cast<uint8_t>(x & 0xFFu); // R
|
||||
pixels[offset + 1] = static_cast<uint8_t>(y & 0xFFu); // G
|
||||
pixels[offset + 2] = static_cast<uint8_t>((x ^ y) & 0xFFu); // B
|
||||
pixels[offset + 3] = 0xFFu; // A
|
||||
}
|
||||
}
|
||||
return pixels;
|
||||
}
|
||||
|
||||
size_t first_mismatch(const std::vector<uint8_t>& expected, const uint8_t* actual, size_t size)
|
||||
{
|
||||
for (size_t index = 0; index < size; ++index) {
|
||||
if (expected[index] != actual[index]) {
|
||||
return index;
|
||||
}
|
||||
}
|
||||
return size;
|
||||
}
|
||||
|
||||
void print_pattern_mismatch(size_t offset, size_t total)
|
||||
{
|
||||
std::printf(" PREMIERE DIFFERENCE a l'octet %llu sur %llu\n", static_cast<unsigned long long>(offset),
|
||||
static_cast<unsigned long long>(total));
|
||||
std::printf(" pixel (%llu, %llu), composante %llu\n",
|
||||
static_cast<unsigned long long>(offset / (4 * kImageWidth)),
|
||||
static_cast<unsigned long long>((offset % (4 * kImageWidth)) / 4),
|
||||
static_cast<unsigned long long>(offset % 4));
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Image exportable prete a partager
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
void create_filled_exportable_image(ExportableImage& out, DeviceContext& context, const std::vector<uint8_t>& pattern)
|
||||
{
|
||||
if (pattern.size() != static_cast<size_t>(kImageBytes)) {
|
||||
throw std::runtime_error("Motif de taille inattendue.");
|
||||
}
|
||||
|
||||
// 1. Staging + motif.
|
||||
create_host_buffer(out.staging, context, kImageBytes, "staging de l'image exportable");
|
||||
{
|
||||
void* pointer = map_memory(context.device, out.staging.memory, kImageBytes, "vkMapMemory (staging export)");
|
||||
std::memcpy(pointer, pattern.data(), pattern.size());
|
||||
flush_mapped_memory(context.device, out.staging.memory);
|
||||
vkUnmapMemory(context.device, out.staging.memory);
|
||||
}
|
||||
|
||||
// 2. Image + allocation DEVICE_LOCAL exportable dediee.
|
||||
create_external_image(out.image, context.device, "vkCreateImage (export)");
|
||||
|
||||
VkMemoryRequirements requirements{};
|
||||
vkGetImageMemoryRequirements(context.device, out.image.image, &requirements);
|
||||
out.allocationSize = requirements.size;
|
||||
out.memoryTypeIndex = find_memory_type(context.physicalDevice, requirements.memoryTypeBits,
|
||||
VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, "image exportable DEVICE_LOCAL");
|
||||
|
||||
VkMemoryDedicatedAllocateInfo dedicated{};
|
||||
dedicated.sType = VK_STRUCTURE_TYPE_MEMORY_DEDICATED_ALLOCATE_INFO;
|
||||
dedicated.image = out.image.image;
|
||||
|
||||
VkExportMemoryAllocateInfo exportAllocateInfo{};
|
||||
exportAllocateInfo.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO;
|
||||
exportAllocateInfo.pNext = &dedicated;
|
||||
exportAllocateInfo.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
VkMemoryAllocateInfo allocateInfo{};
|
||||
allocateInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO;
|
||||
allocateInfo.pNext = &exportAllocateInfo;
|
||||
allocateInfo.allocationSize = requirements.size;
|
||||
allocateInfo.memoryTypeIndex = out.memoryTypeIndex;
|
||||
check(vkAllocateMemory(context.device, &allocateInfo, nullptr, &out.image.memory),
|
||||
"vkAllocateMemory (export)");
|
||||
check(vkBindImageMemory(context.device, out.image.image, out.image.memory, 0), "vkBindImageMemory (export)");
|
||||
|
||||
// Cette allocation est toujours dediee : cudaImportExternalMemory recevra
|
||||
// cudaExternalMemoryDedicated.
|
||||
out.dedicated = true;
|
||||
|
||||
// 3. Remplissage CPU -> GPU.
|
||||
{
|
||||
const VkBufferImageCopy region = full_image_copy();
|
||||
begin_commands(context);
|
||||
cmd_image_barrier(context.commandBuffer, out.image.image, VK_IMAGE_LAYOUT_UNDEFINED,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, 0,
|
||||
VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_WRITE_BIT);
|
||||
vkCmdCopyBufferToImage(context.commandBuffer, out.staging.buffer, out.image.image,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, 1, ®ion);
|
||||
cmd_image_barrier(context.commandBuffer, out.image.image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_READ_BIT);
|
||||
submit_and_wait(context);
|
||||
}
|
||||
|
||||
// 4. Export du HANDLE. vkQueueWaitIdle ci-dessus suffit pour cette etape :
|
||||
// les semaphores externes sont l'etape 4 du prototype.
|
||||
check(vkDeviceWaitIdle(context.device), "vkDeviceWaitIdle (export)");
|
||||
|
||||
const auto getWin32Handle = reinterpret_cast<PFN_vkGetMemoryWin32HandleKHR>(
|
||||
vkGetDeviceProcAddr(context.device, "vkGetMemoryWin32HandleKHR"));
|
||||
if (getWin32Handle == nullptr) {
|
||||
throw std::runtime_error("vkGetMemoryWin32HandleKHR introuvable (extension non exposee).");
|
||||
}
|
||||
|
||||
VkMemoryGetWin32HandleInfoKHR handleInfo{};
|
||||
handleInfo.sType = VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR;
|
||||
handleInfo.memory = out.image.memory;
|
||||
handleInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT;
|
||||
|
||||
const Clock::time_point exportStart = Clock::now();
|
||||
check(getWin32Handle(context.device, &handleInfo, &out.handle.value), "vkGetMemoryWin32HandleKHR");
|
||||
out.exportMs = elapsed_ms(exportStart);
|
||||
|
||||
if (out.handle.value == nullptr) {
|
||||
throw std::runtime_error("vkGetMemoryWin32HandleKHR a retourne un HANDLE nul.");
|
||||
}
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Reference : aller-retour GPU -> CPU -> GPU
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
RoundtripResult measure_cpu_roundtrip(DeviceContext& context, VkImage image, const BufferResource& staging,
|
||||
const std::vector<uint8_t>& pattern, int iterations)
|
||||
{
|
||||
RoundtripResult result;
|
||||
std::vector<uint8_t> hostPixels(static_cast<size_t>(kImageBytes));
|
||||
bool compared = false;
|
||||
|
||||
for (int iteration = 0; iteration < iterations; ++iteration) {
|
||||
const Clock::time_point roundtripStart = Clock::now();
|
||||
|
||||
// GPU -> CPU
|
||||
{
|
||||
const VkBufferImageCopy region = full_image_copy();
|
||||
begin_commands(context);
|
||||
cmd_image_barrier(context.commandBuffer, image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT, 0,
|
||||
VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_READ_BIT);
|
||||
vkCmdCopyImageToBuffer(context.commandBuffer, image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, staging.buffer,
|
||||
1, ®ion);
|
||||
cmd_buffer_barrier(context.commandBuffer, staging.buffer, kImageBytes, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_HOST_BIT, VK_ACCESS_HOST_READ_BIT);
|
||||
submit_and_wait(context);
|
||||
}
|
||||
|
||||
{
|
||||
void* pointer = map_memory(context.device, staging.memory, kImageBytes, "vkMapMemory (relecture)");
|
||||
invalidate_mapped_memory(context.device, staging.memory);
|
||||
std::memcpy(hostPixels.data(), pointer, hostPixels.size());
|
||||
vkUnmapMemory(context.device, staging.memory);
|
||||
}
|
||||
|
||||
if (!compared) {
|
||||
result.contentIdentical =
|
||||
first_mismatch(pattern, hostPixels.data(), hostPixels.size()) == hostPixels.size();
|
||||
compared = true;
|
||||
}
|
||||
|
||||
// CPU -> GPU
|
||||
{
|
||||
void* pointer = map_memory(context.device, staging.memory, kImageBytes, "vkMapMemory (renvoi)");
|
||||
std::memcpy(pointer, hostPixels.data(), hostPixels.size());
|
||||
flush_mapped_memory(context.device, staging.memory);
|
||||
vkUnmapMemory(context.device, staging.memory);
|
||||
|
||||
const VkBufferImageCopy region = full_image_copy();
|
||||
begin_commands(context);
|
||||
cmd_buffer_barrier(context.commandBuffer, staging.buffer, kImageBytes, VK_PIPELINE_STAGE_HOST_BIT,
|
||||
VK_ACCESS_HOST_WRITE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_READ_BIT);
|
||||
cmd_image_barrier(context.commandBuffer, image, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, VK_PIPELINE_STAGE_TRANSFER_BIT, 0,
|
||||
VK_PIPELINE_STAGE_TRANSFER_BIT, VK_ACCESS_TRANSFER_WRITE_BIT);
|
||||
vkCmdCopyBufferToImage(context.commandBuffer, staging.buffer, image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
|
||||
1, ®ion);
|
||||
cmd_image_barrier(context.commandBuffer, image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL,
|
||||
VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_WRITE_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT,
|
||||
VK_ACCESS_TRANSFER_READ_BIT);
|
||||
submit_and_wait(context);
|
||||
}
|
||||
|
||||
result.samplesMs.push_back(elapsed_ms(roundtripStart));
|
||||
}
|
||||
|
||||
std::sort(result.samplesMs.begin(), result.samplesMs.end());
|
||||
if (!result.samplesMs.empty()) {
|
||||
result.minMs = result.samplesMs.front();
|
||||
result.medianMs = result.samplesMs[result.samplesMs.size() / 2];
|
||||
}
|
||||
return result;
|
||||
}
|
||||
|
||||
} // namespace gpu_interop
|
||||
@@ -0,0 +1,205 @@
|
||||
// SPDX-License-Identifier: MPL-2.0
|
||||
//
|
||||
// PROTOTYPE JETABLE - pas du code produit.
|
||||
//
|
||||
// Code Vulkan commun aux etapes 1 et 2 du prototype gpu_interop
|
||||
// (docs/SEQUA_Engine_resume_complet_v2.md §5.6) : instance headless, device,
|
||||
// image exportable en HANDLE Win32, staging, motif deterministe et RAII.
|
||||
//
|
||||
// Windows uniquement (VK_KHR_external_memory_win32).
|
||||
|
||||
#pragma once
|
||||
|
||||
#ifndef NOMINMAX
|
||||
#define NOMINMAX
|
||||
#endif
|
||||
#ifndef WIN32_LEAN_AND_MEAN
|
||||
#define WIN32_LEAN_AND_MEAN
|
||||
#endif
|
||||
|
||||
#include <windows.h>
|
||||
|
||||
#include <vulkan/vulkan.h>
|
||||
#include <vulkan/vulkan_win32.h>
|
||||
|
||||
#include <array>
|
||||
#include <chrono>
|
||||
#include <cstddef>
|
||||
#include <cstdint>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
namespace gpu_interop {
|
||||
|
||||
using Clock = std::chrono::steady_clock;
|
||||
|
||||
inline constexpr uint32_t kImageWidth = 1920;
|
||||
inline constexpr uint32_t kImageHeight = 1080;
|
||||
inline constexpr VkDeviceSize kBytesPerPixel = 4;
|
||||
inline constexpr VkDeviceSize kImageBytes = static_cast<VkDeviceSize>(kImageWidth) * kImageHeight * kBytesPerPixel;
|
||||
inline constexpr VkFormat kImageFormat = VK_FORMAT_R8G8B8A8_UNORM;
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Erreurs et formatage
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
const char* result_name(VkResult result);
|
||||
[[noreturn]] void throw_vk(const char* what, VkResult result);
|
||||
void check(VkResult result, const char* what);
|
||||
double elapsed_ms(Clock::time_point start);
|
||||
const char* device_type_name(VkPhysicalDeviceType type);
|
||||
std::string uuid_to_string(const std::array<uint8_t, VK_UUID_SIZE>& uuid);
|
||||
void print_external_format_support(VkPhysicalDevice physicalDevice);
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// RAII minimal
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
struct Win32Handle {
|
||||
HANDLE value = nullptr;
|
||||
|
||||
Win32Handle() = default;
|
||||
Win32Handle(const Win32Handle&) = delete;
|
||||
Win32Handle& operator=(const Win32Handle&) = delete;
|
||||
~Win32Handle();
|
||||
|
||||
/// Ferme le handle immediatement et remet value a nullptr.
|
||||
bool close();
|
||||
};
|
||||
|
||||
struct InstanceContext {
|
||||
VkInstance instance = VK_NULL_HANDLE;
|
||||
|
||||
InstanceContext() = default;
|
||||
InstanceContext(const InstanceContext&) = delete;
|
||||
InstanceContext& operator=(const InstanceContext&) = delete;
|
||||
~InstanceContext();
|
||||
};
|
||||
|
||||
struct DeviceContext {
|
||||
VkPhysicalDevice physicalDevice = VK_NULL_HANDLE;
|
||||
VkDevice device = VK_NULL_HANDLE;
|
||||
VkQueue queue = VK_NULL_HANDLE;
|
||||
uint32_t queueFamily = 0;
|
||||
VkCommandPool commandPool = VK_NULL_HANDLE;
|
||||
VkCommandBuffer commandBuffer = VK_NULL_HANDLE;
|
||||
|
||||
DeviceContext() = default;
|
||||
DeviceContext(const DeviceContext&) = delete;
|
||||
DeviceContext& operator=(const DeviceContext&) = delete;
|
||||
~DeviceContext();
|
||||
};
|
||||
|
||||
struct ImageResource {
|
||||
VkDevice device = VK_NULL_HANDLE;
|
||||
VkImage image = VK_NULL_HANDLE;
|
||||
VkDeviceMemory memory = VK_NULL_HANDLE;
|
||||
|
||||
ImageResource() = default;
|
||||
ImageResource(const ImageResource&) = delete;
|
||||
ImageResource& operator=(const ImageResource&) = delete;
|
||||
~ImageResource();
|
||||
};
|
||||
|
||||
struct BufferResource {
|
||||
VkDevice device = VK_NULL_HANDLE;
|
||||
VkBuffer buffer = VK_NULL_HANDLE;
|
||||
VkDeviceMemory memory = VK_NULL_HANDLE;
|
||||
VkDeviceSize size = 0;
|
||||
|
||||
BufferResource() = default;
|
||||
BufferResource(const BufferResource&) = delete;
|
||||
BufferResource& operator=(const BufferResource&) = delete;
|
||||
~BufferResource();
|
||||
};
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Enumeration / creation
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
struct DeviceInfo {
|
||||
VkPhysicalDevice physicalDevice = VK_NULL_HANDLE;
|
||||
uint32_t queueFamily = 0;
|
||||
VkPhysicalDeviceProperties properties{};
|
||||
std::array<uint8_t, VK_UUID_SIZE> uuid{};
|
||||
bool externalMemory = false;
|
||||
bool externalMemoryWin32 = false;
|
||||
bool externalSemaphore = false;
|
||||
bool externalSemaphoreWin32 = false;
|
||||
};
|
||||
|
||||
void create_instance(InstanceContext& out);
|
||||
DeviceInfo pick_physical_device(VkInstance instance);
|
||||
void create_device(const DeviceInfo& info, bool enableExternalSemaphore, DeviceContext& out, const char* label);
|
||||
std::array<uint8_t, VK_UUID_SIZE> query_device_uuid(VkPhysicalDevice physicalDevice);
|
||||
uint32_t find_memory_type(VkPhysicalDevice physicalDevice, uint32_t typeBits, VkMemoryPropertyFlags required,
|
||||
const char* what);
|
||||
|
||||
void create_external_image(ImageResource& out, VkDevice device, const char* label);
|
||||
void create_host_buffer(BufferResource& out, DeviceContext& context, VkDeviceSize size, const char* label);
|
||||
|
||||
void* map_memory(VkDevice device, VkDeviceMemory memory, VkDeviceSize size, const char* label);
|
||||
void flush_mapped_memory(VkDevice device, VkDeviceMemory memory);
|
||||
void invalidate_mapped_memory(VkDevice device, VkDeviceMemory memory);
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Commandes
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
void begin_commands(DeviceContext& context);
|
||||
void submit_and_wait(DeviceContext& context);
|
||||
void cmd_image_barrier(VkCommandBuffer commandBuffer, VkImage image, VkImageLayout oldLayout, VkImageLayout newLayout,
|
||||
VkPipelineStageFlags sourceStage, VkAccessFlags sourceAccess, VkPipelineStageFlags dstStage,
|
||||
VkAccessFlags dstAccess);
|
||||
void cmd_buffer_barrier(VkCommandBuffer commandBuffer, VkBuffer buffer, VkDeviceSize size,
|
||||
VkPipelineStageFlags sourceStage, VkAccessFlags sourceAccess, VkPipelineStageFlags dstStage,
|
||||
VkAccessFlags dstAccess);
|
||||
VkBufferImageCopy full_image_copy();
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Motif deterministe
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
std::vector<uint8_t> make_pattern();
|
||||
size_t first_mismatch(const std::vector<uint8_t>& expected, const uint8_t* actual, size_t size);
|
||||
void print_pattern_mismatch(size_t offset, size_t total);
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Image exportable prete a partager
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
struct ExportableImage {
|
||||
ImageResource image;
|
||||
BufferResource staging;
|
||||
VkDeviceSize allocationSize = 0;
|
||||
uint32_t memoryTypeIndex = 0;
|
||||
/// Vrai si la memoire Vulkan a ete allouee avec VkMemoryDedicatedAllocateInfo
|
||||
/// (c'est toujours le cas ici). A transmettre a cudaExternalMemoryDedicated.
|
||||
bool dedicated = true;
|
||||
Win32Handle handle;
|
||||
double exportMs = 0.0;
|
||||
};
|
||||
|
||||
/// Cree le staging, l'image 1920x1080 R8G8B8A8_UNORM, l'allocation DEVICE_LOCAL
|
||||
/// exportable, remplit le motif puis exporte le HANDLE (OPAQUE_WIN32).
|
||||
/// Ne fait aucune synchronisation implicite autre que vkDeviceWaitIdle avant
|
||||
/// l'export (suffisant pour cette etape ; les semaphores externes sont l'etape 4).
|
||||
void create_filled_exportable_image(ExportableImage& out, DeviceContext& context, const std::vector<uint8_t>& pattern);
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Reference : aller-retour GPU -> CPU -> GPU
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
struct RoundtripResult {
|
||||
std::vector<double> samplesMs; // trie par ordre croissant
|
||||
double minMs = 0.0;
|
||||
double medianMs = 0.0;
|
||||
bool contentIdentical = false;
|
||||
};
|
||||
|
||||
/// Mesure un aller-retour complet de l'image (deja en TRANSFER_SRC_OPTIMAL) via
|
||||
/// le staging fourni. Inclut les soumissions, les attentes de file et les copies hote.
|
||||
RoundtripResult measure_cpu_roundtrip(DeviceContext& context, VkImage image, const BufferResource& staging,
|
||||
const std::vector<uint8_t>& pattern, int iterations);
|
||||
|
||||
} // namespace gpu_interop
|
||||
@@ -0,0 +1,353 @@
|
||||
// SPDX-License-Identifier: MPL-2.0
|
||||
//
|
||||
// PROTOTYPE JETABLE - pas du code produit.
|
||||
//
|
||||
// Etape 2 du risque n°1 de SEQUA (docs/SEQUA_Engine_resume_complet_v2.md §5.6) :
|
||||
// importer dans CUDA la memoire d'une texture Vulkan exportee en HANDLE Win32,
|
||||
// sans aucune copie CPU sur le chemin mesure.
|
||||
//
|
||||
// Ce programme :
|
||||
// 1. cree le device Vulkan headless et l'image 1920x1080 R8G8B8A8_UNORM
|
||||
// exportable remplie du motif deterministe (vk_common) ;
|
||||
// 2. exporte son HANDLE OPAQUE_WIN32 ;
|
||||
// 3. choisit le device CUDA dont l'UUID correspond au deviceUUID Vulkan ;
|
||||
// 4. importe la memoire (cudaImportExternalMemory) puis mappe une mipmapped
|
||||
// array sur l'image (cudaExternalMemoryGetMappedMipmappedArray) ;
|
||||
// 5. copie le niveau 0 vers un tampon device (cudaMemcpy2DFromArray), mesure
|
||||
// ce chemin GPU->GPU, puis copie vers l'hote pour verification seulement ;
|
||||
// 6. compare octet a octet au motif et rappelle l'aller-retour CPU de l'etape 1.
|
||||
//
|
||||
// Synchronisation : vkQueueWaitIdle avant l'export suffit pour cette etape.
|
||||
// Les semaphores externes sont l'etape 4 du prototype.
|
||||
//
|
||||
// C++ pur : aucun .cu, aucun noyau, uniquement l'API runtime CUDA.
|
||||
// Windows uniquement. Code non produit.
|
||||
|
||||
#include "vk_common.hpp"
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <array>
|
||||
#include <cstdio>
|
||||
#include <cstring>
|
||||
#include <stdexcept>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
namespace {
|
||||
|
||||
using namespace gpu_interop;
|
||||
|
||||
constexpr int kRoundtripIterations = 5;
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Erreurs CUDA
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
[[noreturn]] void throw_cuda(cudaError_t result, const char* what)
|
||||
{
|
||||
char buffer[512];
|
||||
std::snprintf(buffer, sizeof(buffer), "%s : %s (%d)", what, cudaGetErrorString(result), static_cast<int>(result));
|
||||
throw std::runtime_error(buffer);
|
||||
}
|
||||
|
||||
void check_cuda(cudaError_t result, const char* what)
|
||||
{
|
||||
if (result != cudaSuccess) {
|
||||
throw_cuda(result, what);
|
||||
}
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// RAII CUDA (ordre de destruction : buffer, mipmapped array, external memory)
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
struct CudaExternalMemory {
|
||||
cudaExternalMemory_t value = nullptr;
|
||||
|
||||
CudaExternalMemory() = default;
|
||||
CudaExternalMemory(const CudaExternalMemory&) = delete;
|
||||
CudaExternalMemory& operator=(const CudaExternalMemory&) = delete;
|
||||
~CudaExternalMemory()
|
||||
{
|
||||
if (value != nullptr) {
|
||||
(void)cudaDestroyExternalMemory(value);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct CudaMipmappedArray {
|
||||
cudaMipmappedArray_t value = nullptr;
|
||||
|
||||
CudaMipmappedArray() = default;
|
||||
CudaMipmappedArray(const CudaMipmappedArray&) = delete;
|
||||
CudaMipmappedArray& operator=(const CudaMipmappedArray&) = delete;
|
||||
~CudaMipmappedArray()
|
||||
{
|
||||
if (value != nullptr) {
|
||||
(void)cudaFreeMipmappedArray(value);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct CudaBuffer {
|
||||
void* value = nullptr;
|
||||
|
||||
CudaBuffer() = default;
|
||||
CudaBuffer(const CudaBuffer&) = delete;
|
||||
CudaBuffer& operator=(const CudaBuffer&) = delete;
|
||||
~CudaBuffer()
|
||||
{
|
||||
if (value != nullptr) {
|
||||
(void)cudaFree(value);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
struct CudaEventPair {
|
||||
cudaEvent_t start = nullptr;
|
||||
cudaEvent_t stop = nullptr;
|
||||
|
||||
CudaEventPair() = default;
|
||||
CudaEventPair(const CudaEventPair&) = delete;
|
||||
CudaEventPair& operator=(const CudaEventPair&) = delete;
|
||||
~CudaEventPair()
|
||||
{
|
||||
if (start != nullptr) {
|
||||
(void)cudaEventDestroy(start);
|
||||
}
|
||||
if (stop != nullptr) {
|
||||
(void)cudaEventDestroy(stop);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Selection du device CUDA par UUID
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
int select_cuda_device(const std::array<uint8_t, VK_UUID_SIZE>& vulkanUuid)
|
||||
{
|
||||
static_assert(sizeof(cudaUUID_t) == VK_UUID_SIZE, "cudaUUID_t doit contenir 16 octets");
|
||||
|
||||
int count = 0;
|
||||
check_cuda(cudaGetDeviceCount(&count), "cudaGetDeviceCount");
|
||||
if (count <= 0) {
|
||||
throw std::runtime_error("Aucun device CUDA detecte (pilote NVIDIA absent ou CUDA_VISIBLE_DEVICES vide).");
|
||||
}
|
||||
|
||||
std::string listing;
|
||||
for (int index = 0; index < count; ++index) {
|
||||
cudaDeviceProp properties{};
|
||||
check_cuda(cudaGetDeviceProperties(&properties, index), "cudaGetDeviceProperties");
|
||||
|
||||
std::array<uint8_t, VK_UUID_SIZE> uuid{};
|
||||
std::memcpy(uuid.data(), properties.uuid.bytes, VK_UUID_SIZE);
|
||||
std::printf(" CUDA device %d : %s, UUID %s\n", index, properties.name, uuid_to_string(uuid).c_str());
|
||||
|
||||
if (uuid == vulkanUuid) {
|
||||
return index;
|
||||
}
|
||||
listing += " - ";
|
||||
listing += properties.name;
|
||||
listing += " (UUID ";
|
||||
listing += uuid_to_string(uuid);
|
||||
listing += ")\n";
|
||||
}
|
||||
|
||||
throw std::runtime_error("Aucun device CUDA ne correspond au deviceUUID Vulkan " + uuid_to_string(vulkanUuid) +
|
||||
".\n" + listing);
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------------------
|
||||
// Programme
|
||||
// ---------------------------------------------------------------------------
|
||||
|
||||
int run()
|
||||
{
|
||||
std::printf("=== SEQUA - prototype gpu_interop_vk_cuda (etape 2) ===\n");
|
||||
std::printf("Image %ux%u R8G8B8A8_UNORM, %llu octets, DEVICE_LOCAL exportable\n\n", kImageWidth, kImageHeight,
|
||||
static_cast<unsigned long long>(kImageBytes));
|
||||
|
||||
// --- Vulkan : instance, device, image exportable -----------------------
|
||||
InstanceContext instanceContext;
|
||||
create_instance(instanceContext);
|
||||
|
||||
std::printf("[1] Selection du GPU Vulkan\n");
|
||||
const DeviceInfo deviceInfo = pick_physical_device(instanceContext.instance);
|
||||
std::printf(" retenu : %s, UUID %s, file %u\n", deviceInfo.properties.deviceName,
|
||||
uuid_to_string(deviceInfo.uuid).c_str(), deviceInfo.queueFamily);
|
||||
std::printf(" VK_KHR_external_semaphore_win32 : %s (verifie, non utilise a l'etape 2)\n",
|
||||
deviceInfo.externalSemaphoreWin32 ? "disponible" : "indisponible");
|
||||
print_external_format_support(deviceInfo.physicalDevice);
|
||||
|
||||
DeviceContext device;
|
||||
create_device(deviceInfo, false, device, "vkCreateDevice");
|
||||
|
||||
std::printf("[2] Image Vulkan exportable + motif deterministe\n");
|
||||
const std::vector<uint8_t> pattern = make_pattern();
|
||||
|
||||
ExportableImage exportable;
|
||||
create_filled_exportable_image(exportable, device, pattern);
|
||||
|
||||
std::printf(" memoire DEVICE_LOCAL type %u, %llu octets, allocation dediee=%s\n", exportable.memoryTypeIndex,
|
||||
static_cast<unsigned long long>(exportable.allocationSize), exportable.dedicated ? "oui" : "non");
|
||||
std::printf(" HANDLE Win32 0x%p exporte en %.3f ms\n", static_cast<void*>(exportable.handle.value),
|
||||
exportable.exportMs);
|
||||
|
||||
// --- CUDA : device correspondant ---------------------------------------
|
||||
std::printf("[3] Selection du device CUDA par deviceUUID\n");
|
||||
const int cudaDevice = select_cuda_device(deviceInfo.uuid);
|
||||
check_cuda(cudaSetDevice(cudaDevice), "cudaSetDevice");
|
||||
// Force l'initialisation du contexte avant de mesurer l'import.
|
||||
check_cuda(cudaFree(nullptr), "cudaFree(0)");
|
||||
std::printf(" device CUDA %d retenu, contexte initialise\n", cudaDevice);
|
||||
|
||||
// Mesures remplies dans la portee CUDA ci-dessous.
|
||||
double importMs = 0.0;
|
||||
double mapMs = 0.0;
|
||||
double copyMs = 0.0;
|
||||
double copyEventMs = 0.0;
|
||||
double copyGigabytesPerSecond = 0.0;
|
||||
size_t mismatch = 0;
|
||||
bool identical = false;
|
||||
|
||||
// Portee CUDA : le tampon, la mipmapped array et l'external memory sont
|
||||
// liberes a la sortie du bloc, donc avant que Vulkan ne retouche l'image
|
||||
// pour l'aller-retour de reference.
|
||||
{
|
||||
// --- Import de la memoire Vulkan -----------------------------------
|
||||
std::printf("[4] cudaImportExternalMemory\n");
|
||||
cudaExternalMemoryHandleDesc handleDesc{};
|
||||
handleDesc.type = cudaExternalMemoryHandleTypeOpaqueWin32;
|
||||
handleDesc.handle.win32.handle = exportable.handle.value;
|
||||
handleDesc.handle.win32.name = nullptr;
|
||||
handleDesc.size = exportable.allocationSize;
|
||||
// L'allocation Vulkan est dediee (VkMemoryDedicatedAllocateInfo) :
|
||||
// le drapeau cudaExternalMemoryDedicated decrit cette propriete.
|
||||
handleDesc.flags = exportable.dedicated ? static_cast<unsigned int>(cudaExternalMemoryDedicated) : 0u;
|
||||
|
||||
CudaExternalMemory externalMemory;
|
||||
const Clock::time_point importStart = Clock::now();
|
||||
check_cuda(cudaImportExternalMemory(&externalMemory.value, &handleDesc), "cudaImportExternalMemory");
|
||||
importMs = elapsed_ms(importStart);
|
||||
|
||||
std::printf(" taille %llu octets, flags=0x%x (dedicated=%s), %.3f ms\n",
|
||||
static_cast<unsigned long long>(handleDesc.size), handleDesc.flags,
|
||||
exportable.dedicated ? "oui" : "non", importMs);
|
||||
|
||||
// --- Mappage de l'image en mipmapped array CUDA --------------------
|
||||
std::printf("[5] cudaExternalMemoryGetMappedMipmappedArray\n");
|
||||
cudaExternalMemoryMipmappedArrayDesc arrayDesc{};
|
||||
arrayDesc.offset = 0;
|
||||
arrayDesc.formatDesc = cudaCreateChannelDesc<uchar4>(); // 8 bits x 4, non signe
|
||||
arrayDesc.extent = make_cudaExtent(kImageWidth, kImageHeight, 0); // depth 0 = tableau 2D
|
||||
arrayDesc.flags = 0; // l'image n'est pas liee comme cible de rendu cote CUDA
|
||||
arrayDesc.numLevels = 1;
|
||||
|
||||
CudaMipmappedArray mipmappedArray;
|
||||
const Clock::time_point mapStart = Clock::now();
|
||||
check_cuda(cudaExternalMemoryGetMappedMipmappedArray(&mipmappedArray.value, externalMemory.value, &arrayDesc),
|
||||
"cudaExternalMemoryGetMappedMipmappedArray");
|
||||
cudaArray_t level0 = nullptr;
|
||||
check_cuda(cudaGetMipmappedArrayLevel(&level0, mipmappedArray.value, 0), "cudaGetMipmappedArrayLevel");
|
||||
mapMs = elapsed_ms(mapStart);
|
||||
|
||||
std::printf(" format %d/%d/%d/%d bits, extent %ux%u, %u niveau(x), %.3f ms\n", arrayDesc.formatDesc.x,
|
||||
arrayDesc.formatDesc.y, arrayDesc.formatDesc.z, arrayDesc.formatDesc.w, kImageWidth, kImageHeight,
|
||||
arrayDesc.numLevels, mapMs);
|
||||
|
||||
// --- Copie GPU -> GPU (chemin mesure) -----------------------------
|
||||
std::printf("[6] cudaMemcpy2DFromArray vers un tampon device (GPU->GPU)\n");
|
||||
CudaBuffer deviceBuffer;
|
||||
check_cuda(cudaMalloc(&deviceBuffer.value, static_cast<size_t>(kImageBytes)), "cudaMalloc");
|
||||
|
||||
CudaEventPair copyEvents;
|
||||
check_cuda(cudaEventCreate(©Events.start), "cudaEventCreate (start)");
|
||||
check_cuda(cudaEventCreate(©Events.stop), "cudaEventCreate (stop)");
|
||||
|
||||
const size_t rowBytes = static_cast<size_t>(kImageWidth) * static_cast<size_t>(kBytesPerPixel);
|
||||
const Clock::time_point copyStart = Clock::now();
|
||||
check_cuda(cudaEventRecord(copyEvents.start, 0), "cudaEventRecord (start)");
|
||||
check_cuda(cudaMemcpy2DFromArray(deviceBuffer.value, rowBytes, level0, 0, 0, rowBytes,
|
||||
static_cast<size_t>(kImageHeight), cudaMemcpyDeviceToDevice),
|
||||
"cudaMemcpy2DFromArray");
|
||||
check_cuda(cudaEventRecord(copyEvents.stop, 0), "cudaEventRecord (stop)");
|
||||
check_cuda(cudaEventSynchronize(copyEvents.stop), "cudaEventSynchronize (stop)");
|
||||
copyMs = elapsed_ms(copyStart);
|
||||
|
||||
float copyEventMsRaw = 0.0f;
|
||||
check_cuda(cudaEventElapsedTime(©EventMsRaw, copyEvents.start, copyEvents.stop), "cudaEventElapsedTime");
|
||||
copyEventMs = static_cast<double>(copyEventMsRaw);
|
||||
const float copyEventSeconds = copyEventMsRaw / 1000.0f;
|
||||
copyGigabytesPerSecond =
|
||||
copyEventSeconds > 0.0f ? static_cast<double>(kImageBytes) / (static_cast<double>(copyEventSeconds) * 1.0e9)
|
||||
: 0.0;
|
||||
|
||||
std::printf(" %zu octets, std::chrono %.3f ms, cudaEvent %.3f ms (%.1f Go/s)\n",
|
||||
static_cast<size_t>(kImageBytes), copyMs, copyEventMs, copyGigabytesPerSecond);
|
||||
|
||||
// --- Verification (copie vers l'hote, hors chemin mesure) ----------
|
||||
std::printf("[7] Verification (copie device->host, hors chemin mesure)\n");
|
||||
std::vector<uint8_t> hostPixels(static_cast<size_t>(kImageBytes));
|
||||
check_cuda(cudaMemcpy(hostPixels.data(), deviceBuffer.value, hostPixels.size(), cudaMemcpyDeviceToHost),
|
||||
"cudaMemcpy (device->host)");
|
||||
|
||||
mismatch = first_mismatch(pattern, hostPixels.data(), hostPixels.size());
|
||||
identical = mismatch == hostPixels.size();
|
||||
std::printf(" comparaison octet a octet : %s\n", identical ? "IDENTIQUE" : "DIFFERENT");
|
||||
if (!identical) {
|
||||
print_pattern_mismatch(mismatch, hostPixels.size());
|
||||
}
|
||||
}
|
||||
|
||||
std::printf("[8] Ressources CUDA liberees (buffer, mipmapped array, external memory)\n");
|
||||
|
||||
// --- Reference : aller-retour Vulkan GPU -> CPU -> GPU ----------------
|
||||
const RoundtripResult roundtrip =
|
||||
measure_cpu_roundtrip(device, exportable.image.image, exportable.staging, pattern, kRoundtripIterations);
|
||||
|
||||
std::printf("[9] Reference : aller-retour GPU->CPU->GPU de l'etape 1 (%d essais)\n", kRoundtripIterations);
|
||||
for (size_t index = 0; index < roundtrip.samplesMs.size(); ++index) {
|
||||
std::printf(" echantillon %llu (trie) : %.3f ms\n", static_cast<unsigned long long>(index + 1),
|
||||
roundtrip.samplesMs[index]);
|
||||
}
|
||||
std::printf(" min %.3f ms, mediane %.3f ms, contenu %s\n", roundtrip.minMs, roundtrip.medianMs,
|
||||
roundtrip.contentIdentical ? "identique" : "DIFFERENT");
|
||||
|
||||
// Toutes les operations CUDA sont terminees et les objets CUDA liberes :
|
||||
// le HANDLE peut etre ferme. cudaImportExternalMemory ne prend pas
|
||||
// possession du handle (CUDA en garde sa propre reference).
|
||||
const bool handleClosed = exportable.handle.close();
|
||||
std::printf("[10] CloseHandle du HANDLE exporte : %s\n", handleClosed ? "OK" : "ECHEC");
|
||||
|
||||
// --- Resume -------------------------------------------------------------
|
||||
std::printf("\n=== RESULTATS ===\n");
|
||||
std::printf(" Export HANDLE Win32 : %8.3f ms\n", exportable.exportMs);
|
||||
std::printf(" Import CUDA : %8.3f ms\n", importMs);
|
||||
std::printf(" Mappage mipmapped array : %8.3f ms\n", mapMs);
|
||||
std::printf(" Copie GPU->GPU array->buffer : %8.3f ms (cudaEvent), %.3f ms (chrono), %.1f Go/s\n", copyEventMs,
|
||||
copyMs, copyGigabytesPerSecond);
|
||||
std::printf(" Aller-retour GPU->CPU->GPU : %8.3f ms (mediane de %d essais)\n", roundtrip.medianMs,
|
||||
kRoundtripIterations);
|
||||
std::printf(" Verification octet a octet : %s\n", identical ? "OK" : "ECHEC");
|
||||
|
||||
if (!identical) {
|
||||
std::printf("\nECHEC : le contenu importe differe du motif.\n");
|
||||
return 2;
|
||||
}
|
||||
std::printf("\nOK : texture Vulkan lue par CUDA sans copie CPU sur le chemin mesure.\n");
|
||||
return 0;
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
int main()
|
||||
{
|
||||
try {
|
||||
return run();
|
||||
} catch (const std::exception& error) {
|
||||
std::printf("\nERREUR : %s\n", error.what());
|
||||
return 1;
|
||||
}
|
||||
}
|
||||
Reference in new issue
Block a user