diff --git a/prototypes/gpu_interop/CMakeLists.txt b/prototypes/gpu_interop/CMakeLists.txt index e7ff709..d5cb2d8 100644 --- a/prototypes/gpu_interop/CMakeLists.txt +++ b/prototypes/gpu_interop/CMakeLists.txt @@ -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() diff --git a/prototypes/gpu_interop/README.md b/prototypes/gpu_interop/README.md index f85c6fd..3a20564 100644 --- a/prototypes/gpu_interop/README.md +++ b/prototypes/gpu_interop/README.md @@ -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. diff --git a/prototypes/gpu_interop/main.cpp b/prototypes/gpu_interop/main.cpp index 07663dc..36b7e3d 100644 --- a/prototypes/gpu_interop/main.cpp +++ b/prototypes/gpu_interop/main.cpp @@ -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 - -#include -#include - -#include #include -#include -#include #include -#include -#include #include -#include #include 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(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(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(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& 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(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 enumerate_instance_extensions() -{ - uint32_t count = 0; - check(vkEnumerateInstanceExtensionProperties(nullptr, &count, nullptr), "vkEnumerateInstanceExtensionProperties"); - std::vector extensions(count); - if (count != 0) { - check(vkEnumerateInstanceExtensionProperties(nullptr, &count, extensions.data()), - "vkEnumerateInstanceExtensionProperties"); - } - return extensions; -} - -std::vector enumerate_device_extensions(VkPhysicalDevice physicalDevice) -{ - uint32_t count = 0; - check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, nullptr), - "vkEnumerateDeviceExtensionProperties"); - std::vector extensions(count); - if (count != 0) { - check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, extensions.data()), - "vkEnumerateDeviceExtensionProperties"); - } - return extensions; -} - -bool has_extension(const std::vector& 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(VK_API_VERSION_MAJOR(loaderVersion)), - static_cast(VK_API_VERSION_MINOR(loaderVersion)), - static_cast(VK_API_VERSION_PATCH(loaderVersion))); - throw std::runtime_error(buffer); - } - - const std::vector available = enumerate_instance_extensions(); - std::vector 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(enabled.size()); - info.ppEnabledExtensionNames = enabled.data(); - - check(vkCreateInstance(&info, nullptr, &out.instance), "vkCreateInstance"); -} - -std::array 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 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 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 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 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 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(VK_API_VERSION_MAJOR(info.properties.apiVersion)), - static_cast(VK_API_VERSION_MINOR(info.properties.apiVersion)), - static_cast(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 available = enumerate_device_extensions(info.physicalDevice); - std::vector 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(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 make_pattern() -{ - std::vector pixels(static_cast(kImageBytes)); - for (uint32_t y = 0; y < kHeight; ++y) { - for (uint32_t x = 0; x < kWidth; ++x) { - const size_t offset = (static_cast(y) * kWidth + x) * 4; - pixels[offset + 0] = static_cast(x & 0xFFu); // R - pixels[offset + 1] = static_cast(y & 0xFFu); // G - pixels[offset + 2] = static_cast((x ^ y) & 0xFFu); // B - pixels[offset + 3] = 0xFFu; // A - } - } - return pixels; -} - -size_t first_mismatch(const std::vector& 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(offset), - static_cast(total)); - std::printf(" pixel (%llu, %llu), composante %llu\n", - static_cast(offset / (4 * kWidth)), - static_cast((offset % (4 * kWidth)) / 4), - static_cast(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(kImageBytes)); // --- Instance ----------------------------------------------------------- @@ -757,86 +65,17 @@ int run() } std::printf(" correspondance deviceUUID : OK\n"); - // --- Motif + staging ---------------------------------------------------- + // --- Motif + image exportable remplie ----------------------------------- const std::vector 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(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(exportable.allocationSize)); std::printf("[4] Motif deterministe uploade (staging buffer + vkCmdCopyBufferToImage)\n"); - - // --- Export ------------------------------------------------------------- - const auto getWin32Handle = reinterpret_cast( - 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(exported.value), exportMs); + std::printf("[5] Export : HANDLE Win32 0x%p en %.3f ms\n", + static_cast(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(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(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(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 hostPixels(static_cast(kImageBytes)); - std::vector 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(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"); diff --git a/prototypes/gpu_interop/vk_common.cpp b/prototypes/gpu_interop/vk_common.cpp new file mode 100644 index 0000000..72c4799 --- /dev/null +++ b/prototypes/gpu_interop/vk_common.cpp @@ -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 +#include +#include +#include +#include + +namespace gpu_interop { + +namespace { + +std::vector enumerate_instance_extensions() +{ + uint32_t count = 0; + check(vkEnumerateInstanceExtensionProperties(nullptr, &count, nullptr), "vkEnumerateInstanceExtensionProperties"); + std::vector extensions(count); + if (count != 0) { + check(vkEnumerateInstanceExtensionProperties(nullptr, &count, extensions.data()), + "vkEnumerateInstanceExtensionProperties"); + } + return extensions; +} + +std::vector enumerate_device_extensions(VkPhysicalDevice physicalDevice) +{ + uint32_t count = 0; + check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, nullptr), + "vkEnumerateDeviceExtensionProperties"); + std::vector extensions(count); + if (count != 0) { + check(vkEnumerateDeviceExtensionProperties(physicalDevice, nullptr, &count, extensions.data()), + "vkEnumerateDeviceExtensionProperties"); + } + return extensions; +} + +bool has_extension(const std::vector& 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(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(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& 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(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(VK_API_VERSION_MAJOR(loaderVersion)), + static_cast(VK_API_VERSION_MINOR(loaderVersion)), + static_cast(VK_API_VERSION_PATCH(loaderVersion))); + throw std::runtime_error(buffer); + } + + const std::vector available = enumerate_instance_extensions(); + std::vector 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(enabled.size()); + info.ppEnabledExtensionNames = enabled.data(); + + check(vkCreateInstance(&info, nullptr, &out.instance), "vkCreateInstance"); +} + +std::array 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 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 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 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 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(VK_API_VERSION_MAJOR(info.properties.apiVersion)), + static_cast(VK_API_VERSION_MINOR(info.properties.apiVersion)), + static_cast(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 available = enumerate_device_extensions(info.physicalDevice); + std::vector 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(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 make_pattern() +{ + std::vector pixels(static_cast(kImageBytes)); + for (uint32_t y = 0; y < kImageHeight; ++y) { + for (uint32_t x = 0; x < kImageWidth; ++x) { + const size_t offset = (static_cast(y) * kImageWidth + x) * 4; + pixels[offset + 0] = static_cast(x & 0xFFu); // R + pixels[offset + 1] = static_cast(y & 0xFFu); // G + pixels[offset + 2] = static_cast((x ^ y) & 0xFFu); // B + pixels[offset + 3] = 0xFFu; // A + } + } + return pixels; +} + +size_t first_mismatch(const std::vector& 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(offset), + static_cast(total)); + std::printf(" pixel (%llu, %llu), composante %llu\n", + static_cast(offset / (4 * kImageWidth)), + static_cast((offset % (4 * kImageWidth)) / 4), + static_cast(offset % 4)); +} + +// --------------------------------------------------------------------------- +// Image exportable prete a partager +// --------------------------------------------------------------------------- + +void create_filled_exportable_image(ExportableImage& out, DeviceContext& context, const std::vector& pattern) +{ + if (pattern.size() != static_cast(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( + 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& pattern, int iterations) +{ + RoundtripResult result; + std::vector hostPixels(static_cast(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 diff --git a/prototypes/gpu_interop/vk_common.hpp b/prototypes/gpu_interop/vk_common.hpp new file mode 100644 index 0000000..b2852e1 --- /dev/null +++ b/prototypes/gpu_interop/vk_common.hpp @@ -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 + +#include +#include + +#include +#include +#include +#include +#include +#include + +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(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& 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 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 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 make_pattern(); +size_t first_mismatch(const std::vector& 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& pattern); + +// --------------------------------------------------------------------------- +// Reference : aller-retour GPU -> CPU -> GPU +// --------------------------------------------------------------------------- + +struct RoundtripResult { + std::vector 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& pattern, int iterations); + +} // namespace gpu_interop diff --git a/prototypes/gpu_interop/vk_cuda.cpp b/prototypes/gpu_interop/vk_cuda.cpp new file mode 100644 index 0000000..2888688 --- /dev/null +++ b/prototypes/gpu_interop/vk_cuda.cpp @@ -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 + +#include +#include +#include +#include +#include +#include + +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(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& 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 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(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 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(exportable.allocationSize), exportable.dedicated ? "oui" : "non"); + std::printf(" HANDLE Win32 0x%p exporte en %.3f ms\n", static_cast(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(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(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(); // 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(kImageBytes)), "cudaMalloc"); + + CudaEventPair copyEvents; + check_cuda(cudaEventCreate(©Events.start), "cudaEventCreate (start)"); + check_cuda(cudaEventCreate(©Events.stop), "cudaEventCreate (stop)"); + + const size_t rowBytes = static_cast(kImageWidth) * static_cast(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(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(copyEventMsRaw); + const float copyEventSeconds = copyEventMsRaw / 1000.0f; + copyGigabytesPerSecond = + copyEventSeconds > 0.0f ? static_cast(kImageBytes) / (static_cast(copyEventSeconds) * 1.0e9) + : 0.0; + + std::printf(" %zu octets, std::chrono %.3f ms, cudaEvent %.3f ms (%.1f Go/s)\n", + static_cast(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 hostPixels(static_cast(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(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; + } +}