From 5dfc8b40c04855de8be49ff5d9ca5151d8ad7021 Mon Sep 17 00:00:00 2001 From: pghafari <61991377+pghafari@users.noreply.github.com> Date: Wed, 6 Oct 2021 07:57:22 -0400 Subject: [PATCH] SWDEV-245532 - HIP - Vulkan interop example (#2375) Change-Id: I713ff8fd9da49aa521aee278d1476fd8b4739fc5 [ROCm/hip-tests commit: fa8e4f1f8b96aee01b41ed5e791e71e093411efe] --- .../20_hip_vulkan/SineWaveSimulation.h | 101 + .../20_hip_vulkan/SineWaveSimulation.hip | 147 ++ .../20_hip_vulkan/VulkanBaseApp.cpp | 1724 +++++++++++++++++ .../2_Cookbook/20_hip_vulkan/VulkanBaseApp.h | 146 ++ .../2_Cookbook/20_hip_vulkan/buildcmd.txt | 11 + .../2_Cookbook/20_hip_vulkan/linmath.h | 502 +++++ .../samples/2_Cookbook/20_hip_vulkan/main.cpp | 454 +++++ .../2_Cookbook/20_hip_vulkan/sinewave.frag | 38 + .../2_Cookbook/20_hip_vulkan/sinewave.vert | 43 + 9 files changed, 3166 insertions(+) create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.h create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.hip create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.cpp create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.h create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/linmath.h create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/main.cpp create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.frag create mode 100644 projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.vert diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.h b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.h new file mode 100644 index 0000000000..9126cb3e14 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.h @@ -0,0 +1,101 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/* + * Modifications Copyright (C)2021 Advanced + * Micro Devices, Inc. All rights reserved. + */ + +#pragma once +#ifndef __SINESIM_H__ +#define __SINESIM_H__ + +#include +#include +#include +#include "linmath.h" + +class SineWaveSimulation +{ + float *m_heightMap; + size_t m_width, m_height; + int m_blocks, m_threads; +public: + SineWaveSimulation(size_t width, size_t height); + ~SineWaveSimulation(); + void initSimulation(float *heightMap); + void stepSimulation(float time, hipStream_t stream = 0); + void initCudaLaunchConfig(int device); + int initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE); + + size_t getWidth() const { + return m_width; + } + size_t getHeight() const { + return m_height; + } +}; + +template +void check(T result, char const* const func, const char* const file, + int const line) { + if (result) { + fprintf(stderr, "HIP error at %s:%d code=%d \"%s\" \n", file, line, static_cast(result), func); + // static_cast(result), _cudaGetErrorEnum(result), func); + exit(EXIT_FAILURE); + } +} + +// This will output the proper CUDA error strings in the event +// that a CUDA host call returns an error +#define checkHIPErrors(val) check((val), #val, __FILE__, __LINE__) + +// This will output the proper error string when calling cudaGetLastError +#define getLastHIPError(msg) __getLastHIPError(msg, __FILE__, __LINE__) + +inline void __getLastHIPError(const char* errorMessage, const char* file, + const int line) { + hipError_t err = hipGetLastError(); + + if (hipSuccess != err) { + fprintf(stderr, + "%s(%i) : getLastHIPError() HIP error :" + " %s : (%d) %s.\n", + file, line, errorMessage, static_cast(err), + hipGetErrorString(err)); + exit(EXIT_FAILURE); + } + +} + + + +#ifndef MAX +#define MAX(a, b) (a > b ? a : b) +#endif + +#endif // __SINESIM_H__ diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.hip b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.hip new file mode 100644 index 0000000000..41adae8065 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/SineWaveSimulation.hip @@ -0,0 +1,147 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/* + * Modifications Copyright (C)2021 Advanced + * Micro Devices, Inc. All rights reserved. + */ + +#include "SineWaveSimulation.h" +#include +//#include +#include "hip/hip_runtime.h" + + +__global__ void sinewave(float *heightMap, unsigned int width, unsigned int height, float time) +{ + const float freq = 4.0f; + const size_t stride = gridDim.x * blockDim.x; + + // Iterate through the entire array in a way that is + // independent of the grid configuration + for (size_t tid = blockIdx.x * blockDim.x + threadIdx.x; tid < width * height; tid += stride) { + // Calculate the x, y coordinates + const size_t y = tid / width; + const size_t x = tid - y * width; + // Normalize x, y to [0,1] + const float u = ((2.0f * x) / width) - 1.0f; + const float v = ((2.0f * y) / height) - 1.0f; + // Calculate the new height value + const float w = 0.5f * sinf(u * freq + time) * cosf(v * freq + time); + // Store this new height value + heightMap[tid] = w; + } +} + +SineWaveSimulation::SineWaveSimulation(size_t width, size_t height) + : m_heightMap(nullptr), m_width(width), m_height(height) +{ +} + +void SineWaveSimulation::initCudaLaunchConfig(int device) +{ + hipDeviceProp_t prop = {}; + checkHIPErrors(hipSetDevice(device)); + checkHIPErrors(hipGetDeviceProperties(&prop, device)); + + // We don't need large block sizes, since there's not much inter-thread communication + m_threads = prop.warpSize; + + // Use the occupancy calculator and fill the gpu as best as we can + checkHIPErrors(hipOccupancyMaxActiveBlocksPerMultiprocessor(&m_blocks, sinewave, prop.warpSize, 0)); + m_blocks *= prop.multiProcessorCount; + + // Go ahead and the clamp the blocks to the minimum needed for this height/width + m_blocks = std::min(m_blocks, (int)((m_width * m_height + m_threads - 1) / m_threads)); +} + +int SineWaveSimulation::initCuda(uint8_t *vkDeviceUUID, size_t UUID_SIZE) +{ + int current_device = 0; + int device_count = 0; + int devices_prohibited = 0; + + hipDeviceProp_t deviceProp; + checkHIPErrors(hipGetDeviceCount(&device_count)); + + if (device_count == 0) { + fprintf(stderr, "CUDA error: no devices supporting CUDA.\n"); + exit(EXIT_FAILURE); + } + + // Find the GPU which is selected by Vulkan + while (current_device < device_count) { + hipGetDeviceProperties(&deviceProp, current_device); + + if ((deviceProp.computeMode != hipComputeModeProhibited)) { + // Compare the cuda device UUID with vulkan UUID + // FIXME + int ret = 0; // memcmp((void*)&deviceProp.uuid, vkDeviceUUID, UUID_SIZE); + if (ret == 0) + { + checkHIPErrors(hipSetDevice(current_device)); + checkHIPErrors(hipGetDeviceProperties(&deviceProp, current_device)); + printf("GPU Device %d: \"%s\" with compute capability %d.%d\n\n", + current_device, deviceProp.name, deviceProp.major, + deviceProp.minor); + + return current_device; + } + + } else { + devices_prohibited++; + } + + current_device++; + } + + if (devices_prohibited == device_count) { + fprintf(stderr, + "HIP error:" + " No Vulkan-HIP Interop capable GPU found.\n"); + exit(EXIT_FAILURE); + } + + return -1; +} + +SineWaveSimulation::~SineWaveSimulation() +{ + m_heightMap = NULL; +} + +void SineWaveSimulation::initSimulation(float *heights) +{ + m_heightMap = heights; +} + +void SineWaveSimulation::stepSimulation(float time, hipStream_t stream) +{ + hipLaunchKernelGGL(sinewave, dim3(m_blocks), dim3(m_threads), 0, stream , m_heightMap, m_width, m_height, time); + getLastHIPError("Failed to launch CUDA simulation"); + //hipStreamSynchronize(stream); +} diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.cpp b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.cpp new file mode 100644 index 0000000000..deb86adeb8 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.cpp @@ -0,0 +1,1724 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/* + * This file contains basic cross-platform setup paths in working with Vulkan + * and rendering window. It is largely based off of tutorials provided here: + * https://vulkan-tutorial.com/ +*/ + +/* + * Modifications Copyright (C)2021 Advanced + * Micro Devices, Inc. All rights reserved. + */ + +#include +#include +#include +#include +#include +#include +#include + +#include "VulkanBaseApp.h" + +#define GLFW_INCLUDE_VULKAN +#define GLM_FORCE_DEPTH_ZERO_TO_ONE +#include + +#ifdef _WIN64 +#include +#include +#include +#endif /* _WIN64 */ + +#ifndef countof +#define countof(x) (sizeof(x) / sizeof(*(x))) +#endif + +static const char *validationLayers[] = { "VK_LAYER_KHRONOS_validation" }; +static const size_t MAX_FRAMES_IN_FLIGHT = 5; + +void VulkanBaseApp::resizeCallback(GLFWwindow *window, int width, int height) +{ + VulkanBaseApp *app = reinterpret_cast(glfwGetWindowUserPointer(window)); + app->m_framebufferResized = true; +} + +static VKAPI_ATTR VkBool32 VKAPI_CALL debugCallback(VkDebugUtilsMessageSeverityFlagBitsEXT messageSeverity, VkDebugUtilsMessageTypeFlagsEXT messageType, const VkDebugUtilsMessengerCallbackDataEXT *pCallbackData, void *pUserData) +{ + std::cerr << "validation layer: " << pCallbackData->pMessage << std::endl; + + return VK_FALSE; +} + +VulkanBaseApp::VulkanBaseApp(const std::string& appName, bool enableValidation) : + m_appName(appName), + m_enableValidation(enableValidation), + m_instance(VK_NULL_HANDLE), + m_window(nullptr), + m_debugMessenger(VK_NULL_HANDLE), + m_surface(VK_NULL_HANDLE), + m_physicalDevice(VK_NULL_HANDLE), + m_device(VK_NULL_HANDLE), + m_graphicsQueue(VK_NULL_HANDLE), + m_presentQueue(VK_NULL_HANDLE), + m_swapChain(VK_NULL_HANDLE), + m_vkDeviceUUID(), + m_swapChainImages(), + m_swapChainFormat(), + m_swapChainExtent(), + m_swapChainImageViews(), + m_shaderFiles(), + m_renderPass(), + m_pipelineLayout(VK_NULL_HANDLE), + m_graphicsPipeline(VK_NULL_HANDLE), + m_swapChainFramebuffers(), + m_commandPool(VK_NULL_HANDLE), + m_commandBuffers(), + m_imageAvailableSemaphores(), + m_renderFinishedSemaphores(), + m_inFlightFences(), + m_uniformBuffers(), + m_uniformMemory(), + m_descriptorSetLayout(VK_NULL_HANDLE), + m_descriptorPool(VK_NULL_HANDLE), + m_descriptorSets(), + m_depthImage(VK_NULL_HANDLE), + m_depthImageMemory(VK_NULL_HANDLE), + m_depthImageView(VK_NULL_HANDLE), + m_currentFrame(0), + m_framebufferResized(false) +{ +} + +VkExternalSemaphoreHandleTypeFlagBits VulkanBaseApp::getDefaultSemaphoreHandleType() +{ +#ifdef _WIN64 + return IsWindows8OrGreater() ? + VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT : + VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT; +#else + return VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT; +#endif +} + +VkExternalMemoryHandleTypeFlagBits VulkanBaseApp::getDefaultMemHandleType() +{ +#ifdef _WIN64 + return IsWindows8Point1OrGreater() ? + VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT : + VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT; +#else + return VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; +#endif +} + +VulkanBaseApp::~VulkanBaseApp() +{ + cleanupSwapChain(); + + if (m_descriptorSetLayout != VK_NULL_HANDLE) { + vkDestroyDescriptorSetLayout(m_device, m_descriptorSetLayout, nullptr); + } + + for (size_t i = 0; i < m_renderFinishedSemaphores.size(); i++) { + vkDestroySemaphore(m_device, m_renderFinishedSemaphores[i], nullptr); + vkDestroySemaphore(m_device, m_imageAvailableSemaphores[i], nullptr); + vkDestroyFence(m_device, m_inFlightFences[i], nullptr); + } + if (m_commandPool != VK_NULL_HANDLE) { + vkDestroyCommandPool(m_device, m_commandPool, nullptr); + } + + if (m_device != VK_NULL_HANDLE) { + vkDestroyDevice(m_device, nullptr); + } + + if (m_enableValidation) { + PFN_vkDestroyDebugUtilsMessengerEXT func = (PFN_vkDestroyDebugUtilsMessengerEXT)vkGetInstanceProcAddr(m_instance, "vkDestroyDebugUtilsMessengerEXT"); + if (func != nullptr) { + func(m_instance, m_debugMessenger, nullptr); + } + } + + if (m_surface != VK_NULL_HANDLE) { + vkDestroySurfaceKHR(m_instance, m_surface, nullptr); + } + + if (m_instance != VK_NULL_HANDLE) { + vkDestroyInstance(m_instance, nullptr); + } + + if (m_window) { + glfwDestroyWindow(m_window); + } + + glfwTerminate(); +} + +void VulkanBaseApp::init() +{ + initWindow(); + initVulkan(); +} + +VkCommandBuffer VulkanBaseApp::beginSingleTimeCommands() +{ + VkCommandBufferAllocateInfo allocInfo = {}; + allocInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO; + allocInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + allocInfo.commandPool = m_commandPool; + allocInfo.commandBufferCount = 1; + + VkCommandBuffer commandBuffer; + vkAllocateCommandBuffers(m_device, &allocInfo, &commandBuffer); + + VkCommandBufferBeginInfo beginInfo = {}; + beginInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + beginInfo.flags = VK_COMMAND_BUFFER_USAGE_ONE_TIME_SUBMIT_BIT; + + vkBeginCommandBuffer(commandBuffer, &beginInfo); + + return commandBuffer; +} + +void VulkanBaseApp::endSingleTimeCommands(VkCommandBuffer commandBuffer) +{ + vkEndCommandBuffer(commandBuffer); + + VkSubmitInfo submitInfo = {}; + submitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + submitInfo.commandBufferCount = 1; + submitInfo.pCommandBuffers = &commandBuffer; + + vkQueueSubmit(m_graphicsQueue, 1, &submitInfo, VK_NULL_HANDLE); + vkQueueWaitIdle(m_graphicsQueue); + + vkFreeCommandBuffers(m_device, m_commandPool, 1, &commandBuffer); +} + +void VulkanBaseApp::initWindow() +{ + glfwInit(); + + glfwWindowHint(GLFW_CLIENT_API, GLFW_NO_API); + glfwWindowHint(GLFW_RESIZABLE, GLFW_FALSE); + + m_window = glfwCreateWindow(1280, 800, m_appName.c_str(), nullptr, nullptr); + glfwSetWindowUserPointer(m_window, this); + glfwSetFramebufferSizeCallback(m_window, resizeCallback); +} + + +std::vector VulkanBaseApp::getRequiredExtensions() const +{ + return std::vector(); +} + +std::vector VulkanBaseApp::getRequiredDeviceExtensions() const +{ + return std::vector(); +} + +void VulkanBaseApp::initVulkan() +{ + createInstance(); + createSurface(); + createDevice(); + createSwapChain(); + createImageViews(); + createRenderPass(); + createDescriptorSetLayout(); + createGraphicsPipeline(); + createCommandPool(); + createDepthResources(); + createFramebuffers(); + initVulkanApp(); + createUniformBuffers(); + createDescriptorPool(); + createDescriptorSets(); + createCommandBuffers(); + createSyncObjects(); +} + +#ifdef _WIN64 +class WindowsSecurityAttributes +{ +protected: + SECURITY_ATTRIBUTES m_winSecurityAttributes; + PSECURITY_DESCRIPTOR m_winPSecurityDescriptor; + +public: + WindowsSecurityAttributes(); + SECURITY_ATTRIBUTES *operator&(); + ~WindowsSecurityAttributes(); +}; + +WindowsSecurityAttributes::WindowsSecurityAttributes() +{ + m_winPSecurityDescriptor = (PSECURITY_DESCRIPTOR)calloc(1, SECURITY_DESCRIPTOR_MIN_LENGTH + 2 * sizeof(void **)); + if (!m_winPSecurityDescriptor) { + throw std::runtime_error("Failed to allocate memory for security descriptor"); + } + + PSID *ppSID = (PSID *)((PBYTE)m_winPSecurityDescriptor + SECURITY_DESCRIPTOR_MIN_LENGTH); + PACL *ppACL = (PACL *)((PBYTE)ppSID + sizeof(PSID *)); + + InitializeSecurityDescriptor(m_winPSecurityDescriptor, SECURITY_DESCRIPTOR_REVISION); + + SID_IDENTIFIER_AUTHORITY sidIdentifierAuthority = SECURITY_WORLD_SID_AUTHORITY; + AllocateAndInitializeSid(&sidIdentifierAuthority, 1, SECURITY_WORLD_RID, 0, 0, 0, 0, 0, 0, 0, ppSID); + + EXPLICIT_ACCESS explicitAccess; + ZeroMemory(&explicitAccess, sizeof(EXPLICIT_ACCESS)); + explicitAccess.grfAccessPermissions = STANDARD_RIGHTS_ALL | SPECIFIC_RIGHTS_ALL; + explicitAccess.grfAccessMode = SET_ACCESS; + explicitAccess.grfInheritance = INHERIT_ONLY; + explicitAccess.Trustee.TrusteeForm = TRUSTEE_IS_SID; + explicitAccess.Trustee.TrusteeType = TRUSTEE_IS_WELL_KNOWN_GROUP; + explicitAccess.Trustee.ptstrName = (LPTSTR) * ppSID; + + SetEntriesInAcl(1, &explicitAccess, NULL, ppACL); + + SetSecurityDescriptorDacl(m_winPSecurityDescriptor, TRUE, *ppACL, FALSE); + + m_winSecurityAttributes.nLength = sizeof(m_winSecurityAttributes); + m_winSecurityAttributes.lpSecurityDescriptor = m_winPSecurityDescriptor; + m_winSecurityAttributes.bInheritHandle = TRUE; +} + +SECURITY_ATTRIBUTES * +WindowsSecurityAttributes::operator&() +{ + return &m_winSecurityAttributes; +} + +WindowsSecurityAttributes::~WindowsSecurityAttributes() +{ + PSID *ppSID = (PSID *)((PBYTE)m_winPSecurityDescriptor + SECURITY_DESCRIPTOR_MIN_LENGTH); + PACL *ppACL = (PACL *)((PBYTE)ppSID + sizeof(PSID *)); + + if (*ppSID) { + FreeSid(*ppSID); + } + if (*ppACL) { + LocalFree(*ppACL); + } + free(m_winPSecurityDescriptor); +} +#endif /* _WIN64 */ + + +static VkFormat findSupportedFormat(VkPhysicalDevice physicalDevice, const std::vector& candidates, VkImageTiling tiling, VkFormatFeatureFlags features) +{ + for (VkFormat format : candidates) { + VkFormatProperties props; + vkGetPhysicalDeviceFormatProperties(physicalDevice, format, &props); + if (tiling == VK_IMAGE_TILING_LINEAR && (props.linearTilingFeatures & features) == features) { + return format; + } + else if (tiling == VK_IMAGE_TILING_OPTIMAL && (props.optimalTilingFeatures & features) == features) { + return format; + } + } + throw std::runtime_error("Failed to find supported format!"); +} + +static uint32_t findMemoryType(VkPhysicalDevice physicalDevice, uint32_t typeFilter, VkMemoryPropertyFlags properties) +{ + VkPhysicalDeviceMemoryProperties memProperties; + vkGetPhysicalDeviceMemoryProperties(physicalDevice, &memProperties); + for (uint32_t i = 0; i < memProperties.memoryTypeCount; i++) { + if (typeFilter & (1 << i) && (memProperties.memoryTypes[i].propertyFlags & properties) == properties) { + return i; + } + } + return ~0; +} + +static bool supportsValidationLayers() +{ + std::vector availableLayers; + uint32_t layerCount; + + vkEnumerateInstanceLayerProperties(&layerCount, nullptr); + availableLayers.resize(layerCount); + vkEnumerateInstanceLayerProperties(&layerCount, availableLayers.data()); + + for (const char * layerName : validationLayers) { + bool layerFound = false; + + for (const auto & layerProperties : availableLayers) { + if (strcmp(layerName, layerProperties.layerName) == 0) { + layerFound = true; + break; + } + } + + if (!layerFound) { + return false; + } + } + + return true; +} + +void VulkanBaseApp::createInstance() +{ + if (m_enableValidation && !supportsValidationLayers()) { + throw std::runtime_error("Validation requested, but not supported!"); + } + + VkApplicationInfo appInfo = {}; + appInfo.sType = VK_STRUCTURE_TYPE_APPLICATION_INFO; + appInfo.pApplicationName = m_appName.c_str(); + appInfo.applicationVersion = VK_MAKE_VERSION(1, 0, 0); + appInfo.pEngineName = "No Engine"; + appInfo.engineVersion = VK_MAKE_VERSION(1, 0, 0); + appInfo.apiVersion = VK_API_VERSION_1_0; + + VkInstanceCreateInfo createInfo = {}; + createInfo.sType = VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO; + createInfo.pApplicationInfo = &appInfo; + + std::vector exts = getRequiredExtensions(); + + { + uint32_t glfwExtensionCount = 0; + const char **glfwExtensions; + + glfwExtensions = glfwGetRequiredInstanceExtensions(&glfwExtensionCount); + + exts.insert(exts.begin(), glfwExtensions, glfwExtensions + glfwExtensionCount); + + if (m_enableValidation) { + exts.push_back(VK_EXT_DEBUG_UTILS_EXTENSION_NAME); + } + } + + createInfo.enabledExtensionCount = static_cast(exts.size()); + createInfo.ppEnabledExtensionNames = exts.data(); + VkDebugUtilsMessengerCreateInfoEXT debugCreateInfo = {}; + if (m_enableValidation) { + createInfo.enabledLayerCount = static_cast(countof(validationLayers)); + createInfo.ppEnabledLayerNames = validationLayers; + + debugCreateInfo.sType = VK_STRUCTURE_TYPE_DEBUG_UTILS_MESSENGER_CREATE_INFO_EXT; + debugCreateInfo.messageSeverity = VK_DEBUG_UTILS_MESSAGE_SEVERITY_VERBOSE_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_SEVERITY_WARNING_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_SEVERITY_ERROR_BIT_EXT; + debugCreateInfo.messageType = VK_DEBUG_UTILS_MESSAGE_TYPE_GENERAL_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_TYPE_VALIDATION_BIT_EXT | VK_DEBUG_UTILS_MESSAGE_TYPE_PERFORMANCE_BIT_EXT; + debugCreateInfo.pfnUserCallback = debugCallback; + + createInfo.pNext = &debugCreateInfo; + } + else { + createInfo.enabledLayerCount = 0; + createInfo.pNext = nullptr; + } + + if (vkCreateInstance(&createInfo, nullptr, &m_instance) != VK_SUCCESS) { + throw std::runtime_error("Failed to create Vulkan instance!"); + } + + if (m_enableValidation) { + PFN_vkCreateDebugUtilsMessengerEXT func = (PFN_vkCreateDebugUtilsMessengerEXT)vkGetInstanceProcAddr(m_instance, "vkCreateDebugUtilsMessengerEXT"); + if (func == nullptr || func(m_instance, &debugCreateInfo, nullptr, &m_debugMessenger) != VK_SUCCESS) { + throw std::runtime_error("Failed to set up debug messenger!"); + } + } +} + +void VulkanBaseApp::createSurface() +{ + if (glfwCreateWindowSurface(m_instance, m_window, nullptr, &m_surface) != VK_SUCCESS) { + throw std::runtime_error("failed to create window surface!"); + } +} + +static bool findGraphicsQueueIndicies(VkPhysicalDevice device, VkSurfaceKHR surface, uint32_t& graphicsFamily, uint32_t& presentFamily) +{ + uint32_t queueFamilyCount = 0; + + vkGetPhysicalDeviceQueueFamilyProperties(device, &queueFamilyCount, nullptr); + + std::vector queueFamilies(queueFamilyCount); + vkGetPhysicalDeviceQueueFamilyProperties(device, &queueFamilyCount, queueFamilies.data()); + + graphicsFamily = presentFamily = ~0; + + for (uint32_t i = 0; i < queueFamilyCount; i++) { + + if (queueFamilies[i].queueCount > 0) { + if (graphicsFamily == ~0 && queueFamilies[i].queueFlags & VK_QUEUE_GRAPHICS_BIT) { + graphicsFamily = i; + } + uint32_t presentSupport = 0; + vkGetPhysicalDeviceSurfaceSupportKHR(device, i, surface, &presentSupport); + if (presentFamily == ~0 && presentSupport) { + presentFamily = i; + } + if (presentFamily != ~0 && graphicsFamily != ~0) { + break; + } + } + } + + return graphicsFamily != ~0 && presentFamily != ~0; +} + +static bool hasAllExtensions(VkPhysicalDevice device, const std::vector& deviceExtensions) +{ + uint32_t extensionCount; + vkEnumerateDeviceExtensionProperties(device, nullptr, &extensionCount, nullptr); + std::vector availableExtensions(extensionCount); + vkEnumerateDeviceExtensionProperties(device, nullptr, &extensionCount, availableExtensions.data()); + + std::set requiredExtensions(deviceExtensions.begin(), deviceExtensions.end()); + + for (const auto & extension : availableExtensions) { + requiredExtensions.erase(extension.extensionName); + } + + return requiredExtensions.empty(); +} + +static void getSwapChainProperties(VkPhysicalDevice device, VkSurfaceKHR surface, VkSurfaceCapabilitiesKHR& capabilities, std::vector& formats, std::vector& presentModes) +{ + vkGetPhysicalDeviceSurfaceCapabilitiesKHR(device, surface, &capabilities); + uint32_t formatCount; + vkGetPhysicalDeviceSurfaceFormatsKHR(device, surface, &formatCount, nullptr); + if (formatCount != 0) { + formats.resize(formatCount); + vkGetPhysicalDeviceSurfaceFormatsKHR(device, surface, &formatCount, formats.data()); + } + uint32_t presentModeCount; + vkGetPhysicalDeviceSurfacePresentModesKHR(device, surface, &presentModeCount, nullptr); + if (presentModeCount != 0) { + presentModes.resize(presentModeCount); + vkGetPhysicalDeviceSurfacePresentModesKHR(device, surface, &presentModeCount, presentModes.data()); + } +} + +bool VulkanBaseApp::isSuitableDevice(VkPhysicalDevice dev) const +{ + uint32_t graphicsQueueIndex, presentQueueIndex; + std::vector deviceExtensions = getRequiredDeviceExtensions(); + VkSurfaceCapabilitiesKHR caps; + std::vector formats; + std::vector presentModes; + deviceExtensions.push_back(VK_KHR_SWAPCHAIN_EXTENSION_NAME); + getSwapChainProperties(dev, m_surface, caps, formats, presentModes); + return hasAllExtensions(dev, deviceExtensions) + && !formats.empty() && !presentModes.empty() + && findGraphicsQueueIndicies(dev, m_surface, graphicsQueueIndex, presentQueueIndex); +} + +void VulkanBaseApp::createDevice() +{ + { + uint32_t deviceCount = 0; + vkEnumeratePhysicalDevices(m_instance, &deviceCount, nullptr); + if (deviceCount == 0) { + throw std::runtime_error("Failed to find Vulkan capable GPUs!"); + } + std::vector phyDevs(deviceCount); + vkEnumeratePhysicalDevices(m_instance, &deviceCount, phyDevs.data()); + std::vector::iterator it = std::find_if(phyDevs.begin(), phyDevs.end(), + std::bind(&VulkanBaseApp::isSuitableDevice, this, std::placeholders::_1)); + if (it == phyDevs.end()) { + throw std::runtime_error("No suitable device found!"); + } + m_physicalDevice = *it; + } + + uint32_t graphicsQueueIndex, presentQueueIndex; + findGraphicsQueueIndicies(m_physicalDevice, m_surface, graphicsQueueIndex, presentQueueIndex); + + std::vector queueCreateInfos; + std::set uniqueFamilyIndices = { graphicsQueueIndex, presentQueueIndex }; + + float queuePriority = 1.0f; + + for (uint32_t queueFamily : uniqueFamilyIndices) { + VkDeviceQueueCreateInfo queueCreateInfo = {}; + queueCreateInfo.sType = VK_STRUCTURE_TYPE_DEVICE_QUEUE_CREATE_INFO; + queueCreateInfo.queueFamilyIndex = graphicsQueueIndex; + queueCreateInfo.queueCount = 1; + queueCreateInfo.pQueuePriorities = &queuePriority; + queueCreateInfos.push_back(queueCreateInfo); + } + + VkPhysicalDeviceFeatures deviceFeatures = {}; + deviceFeatures.fillModeNonSolid = true; + + VkDeviceCreateInfo createInfo = {}; + createInfo.sType = VK_STRUCTURE_TYPE_DEVICE_CREATE_INFO; + + createInfo.pQueueCreateInfos = queueCreateInfos.data(); + createInfo.queueCreateInfoCount = static_cast(queueCreateInfos.size()); + + createInfo.pEnabledFeatures = &deviceFeatures; + + std::vector deviceExtensions = getRequiredDeviceExtensions(); + deviceExtensions.push_back(VK_KHR_SWAPCHAIN_EXTENSION_NAME); + + createInfo.enabledExtensionCount = static_cast(deviceExtensions.size()); + createInfo.ppEnabledExtensionNames = deviceExtensions.data(); + + if (m_enableValidation) { + createInfo.enabledLayerCount = static_cast(countof(validationLayers)); + createInfo.ppEnabledLayerNames = validationLayers; + } + else { + createInfo.enabledLayerCount = 0; + } + VkResult asdf = vkCreateDevice(m_physicalDevice, &createInfo, nullptr, &m_device); + if (asdf != VK_SUCCESS) { + throw std::runtime_error("failed to create logical device!"); + } + + vkGetDeviceQueue(m_device, graphicsQueueIndex, 0, &m_graphicsQueue); + vkGetDeviceQueue(m_device, presentQueueIndex, 0, &m_presentQueue); + + VkPhysicalDeviceIDProperties vkPhysicalDeviceIDProperties = {}; + vkPhysicalDeviceIDProperties.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_ID_PROPERTIES; + vkPhysicalDeviceIDProperties.pNext = NULL; + + VkPhysicalDeviceProperties2 vkPhysicalDeviceProperties2 = {}; + vkPhysicalDeviceProperties2.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PROPERTIES_2; + vkPhysicalDeviceProperties2.pNext = &vkPhysicalDeviceIDProperties; + + PFN_vkGetPhysicalDeviceProperties2 fpGetPhysicalDeviceProperties2; + fpGetPhysicalDeviceProperties2 = (PFN_vkGetPhysicalDeviceProperties2)vkGetInstanceProcAddr(m_instance, "vkGetPhysicalDeviceProperties2"); + if (fpGetPhysicalDeviceProperties2 == NULL) { + throw std::runtime_error("Vulkan: Proc address for \"vkGetPhysicalDeviceProperties2KHR\" not found.\n"); + } + + //fpGetPhysicalDeviceProperties2(m_physicalDevice, &vkPhysicalDeviceProperties2); + + memcpy(m_vkDeviceUUID, vkPhysicalDeviceIDProperties.deviceUUID, VK_UUID_SIZE); +} + +static VkSurfaceFormatKHR chooseSwapSurfaceFormat(const std::vector& availableFormats) +{ + if (availableFormats.size() == 1 && availableFormats[0].format == VK_FORMAT_UNDEFINED) { + return { VK_FORMAT_B8G8R8A8_UNORM, VK_COLOR_SPACE_SRGB_NONLINEAR_KHR }; + } + + for (const auto & availableFormat : availableFormats) { + if (availableFormat.format == VK_FORMAT_B8G8R8A8_UNORM && availableFormat.colorSpace == VK_COLOR_SPACE_SRGB_NONLINEAR_KHR) { + return availableFormat; + } + } + + return availableFormats[0]; +} + +static VkPresentModeKHR chooseSwapPresentMode(const std::vector& availablePresentModes) +{ + VkPresentModeKHR bestMode = VK_PRESENT_MODE_FIFO_KHR; + + for (const auto & availablePresentMode : availablePresentModes) { + if (availablePresentMode == VK_PRESENT_MODE_MAILBOX_KHR) { + return availablePresentMode; + } + else if (availablePresentMode == VK_PRESENT_MODE_IMMEDIATE_KHR) { + bestMode = availablePresentMode; + } + } + + return bestMode; +} + +static VkExtent2D chooseSwapExtent(GLFWwindow *window, const VkSurfaceCapabilitiesKHR& capabilities) +{ + if (capabilities.currentExtent.width != std::numeric_limits::max()) { + return capabilities.currentExtent; + } + else { + int width, height; + glfwGetFramebufferSize(window, &width, &height); + VkExtent2D actualExtent = { static_cast(width), static_cast(height) }; + + actualExtent.width = std::max(capabilities.minImageExtent.width, std::min(capabilities.maxImageExtent.width, actualExtent.width)); + actualExtent.height = std::max(capabilities.minImageExtent.height, std::min(capabilities.maxImageExtent.height, actualExtent.height)); + + return actualExtent; + } +} + +void VulkanBaseApp::createSwapChain() +{ + VkSurfaceCapabilitiesKHR capabilities; + VkSurfaceFormatKHR format; + VkPresentModeKHR presentMode; + VkExtent2D extent; + uint32_t imageCount; + + { + std::vector formats; + std::vector presentModes; + + getSwapChainProperties(m_physicalDevice, m_surface, capabilities, formats, presentModes); + format = chooseSwapSurfaceFormat(formats); + presentMode = chooseSwapPresentMode(presentModes); + extent = chooseSwapExtent(m_window, capabilities); + imageCount = capabilities.minImageCount + 1; + if (capabilities.maxImageCount > 0 && imageCount > capabilities.maxImageCount) { + imageCount = capabilities.maxImageCount; + } + } + + VkSwapchainCreateInfoKHR createInfo = {}; + createInfo.sType = VK_STRUCTURE_TYPE_SWAPCHAIN_CREATE_INFO_KHR; + createInfo.surface = m_surface; + + createInfo.minImageCount = imageCount; + createInfo.imageFormat = format.format; + createInfo.imageColorSpace = format.colorSpace; + createInfo.imageExtent = extent; + createInfo.imageArrayLayers = 1; + createInfo.imageUsage = VK_IMAGE_USAGE_COLOR_ATTACHMENT_BIT; + + uint32_t queueFamilyIndices[2]; + findGraphicsQueueIndicies(m_physicalDevice, m_surface, queueFamilyIndices[0], queueFamilyIndices[1]); + + if (queueFamilyIndices[0] != queueFamilyIndices[1]) { + createInfo.imageSharingMode = VK_SHARING_MODE_CONCURRENT; + createInfo.queueFamilyIndexCount = countof(queueFamilyIndices); + createInfo.pQueueFamilyIndices = queueFamilyIndices; + } + else { + createInfo.imageSharingMode = VK_SHARING_MODE_EXCLUSIVE; + } + + createInfo.preTransform = capabilities.currentTransform; + createInfo.compositeAlpha = VK_COMPOSITE_ALPHA_OPAQUE_BIT_KHR; + createInfo.presentMode = presentMode; + createInfo.clipped = VK_TRUE; + + createInfo.oldSwapchain = VK_NULL_HANDLE; + + if (vkCreateSwapchainKHR(m_device, &createInfo, nullptr, &m_swapChain) != VK_SUCCESS) { + throw std::runtime_error("failed to create swap chain!"); + } + + vkGetSwapchainImagesKHR(m_device, m_swapChain, &imageCount, nullptr); + m_swapChainImages.resize(imageCount); + vkGetSwapchainImagesKHR(m_device, m_swapChain, &imageCount, m_swapChainImages.data()); + + m_swapChainFormat = format.format; + m_swapChainExtent = extent; +} + +static VkImageView createImageView(VkDevice dev, VkImage image, VkFormat format, VkImageAspectFlags aspectFlags) +{ + VkImageView imageView; + VkImageViewCreateInfo createInfo = {}; + createInfo.sType = VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO; + createInfo.image = image; + createInfo.viewType = VK_IMAGE_VIEW_TYPE_2D; + createInfo.format = format; + createInfo.components.r = VK_COMPONENT_SWIZZLE_IDENTITY; + createInfo.components.g = VK_COMPONENT_SWIZZLE_IDENTITY; + createInfo.components.b = VK_COMPONENT_SWIZZLE_IDENTITY; + createInfo.components.a = VK_COMPONENT_SWIZZLE_IDENTITY; + createInfo.subresourceRange.aspectMask = aspectFlags; + createInfo.subresourceRange.baseMipLevel = 0; + createInfo.subresourceRange.levelCount = 1; + createInfo.subresourceRange.baseArrayLayer = 0; + createInfo.subresourceRange.layerCount = 1; + if (vkCreateImageView(dev, &createInfo, nullptr, &imageView) != VK_SUCCESS) { + throw std::runtime_error("Failed to create image views!"); + } + + return imageView; +} + +static void createImage(VkPhysicalDevice physicalDevice, VkDevice device, uint32_t width, uint32_t height, VkFormat format, VkImageTiling tiling, VkImageUsageFlags usage, VkMemoryPropertyFlags properties, VkImage& image, VkDeviceMemory& imageMemory) +{ + VkImageCreateInfo imageInfo = {}; + imageInfo.sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO; + imageInfo.imageType = VK_IMAGE_TYPE_2D; + imageInfo.extent.width = width; + imageInfo.extent.height = height; + imageInfo.extent.depth = 1; + imageInfo.mipLevels = 1; + imageInfo.arrayLayers = 1; + imageInfo.format = format; + imageInfo.tiling = tiling; + imageInfo.initialLayout = VK_IMAGE_LAYOUT_UNDEFINED; + imageInfo.usage = usage; + imageInfo.samples = VK_SAMPLE_COUNT_1_BIT; + imageInfo.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + + if (vkCreateImage(device, &imageInfo, nullptr, &image) != VK_SUCCESS) { + throw std::runtime_error("failed to create image!"); + } + + VkMemoryRequirements memRequirements; + vkGetImageMemoryRequirements(device, image, &memRequirements); + + VkMemoryAllocateInfo allocInfo = {}; + allocInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + allocInfo.allocationSize = memRequirements.size; + allocInfo.memoryTypeIndex = findMemoryType(physicalDevice, memRequirements.memoryTypeBits, properties); + + if (vkAllocateMemory(device, &allocInfo, nullptr, &imageMemory) != VK_SUCCESS) { + throw std::runtime_error("failed to allocate image memory!"); + } + + vkBindImageMemory(device, image, imageMemory, 0); +} + +void VulkanBaseApp::createImageViews() +{ + m_swapChainImageViews.resize(m_swapChainImages.size()); + + for (uint32_t i = 0; i < m_swapChainImages.size(); i++) { + m_swapChainImageViews[i] = createImageView(m_device, m_swapChainImages[i], m_swapChainFormat, VK_IMAGE_ASPECT_COLOR_BIT); + } +} + +void VulkanBaseApp::createRenderPass() +{ + VkAttachmentDescription colorAttachment = {}; + colorAttachment.format = m_swapChainFormat; + colorAttachment.samples = VK_SAMPLE_COUNT_1_BIT; + colorAttachment.loadOp = VK_ATTACHMENT_LOAD_OP_CLEAR; + colorAttachment.storeOp = VK_ATTACHMENT_STORE_OP_STORE; + colorAttachment.stencilLoadOp = VK_ATTACHMENT_LOAD_OP_DONT_CARE; + colorAttachment.stencilStoreOp = VK_ATTACHMENT_STORE_OP_DONT_CARE; + colorAttachment.initialLayout = VK_IMAGE_LAYOUT_UNDEFINED; + colorAttachment.finalLayout = VK_IMAGE_LAYOUT_PRESENT_SRC_KHR; + + VkAttachmentReference colorAttachmentRef = {}; + colorAttachmentRef.attachment = 0; + colorAttachmentRef.layout = VK_IMAGE_LAYOUT_COLOR_ATTACHMENT_OPTIMAL; + + VkAttachmentDescription depthAttachment = {}; + depthAttachment.format = findSupportedFormat(m_physicalDevice, + { VK_FORMAT_D32_SFLOAT, VK_FORMAT_D32_SFLOAT_S8_UINT, VK_FORMAT_D24_UNORM_S8_UINT }, + VK_IMAGE_TILING_OPTIMAL, VK_FORMAT_FEATURE_DEPTH_STENCIL_ATTACHMENT_BIT); + depthAttachment.samples = VK_SAMPLE_COUNT_1_BIT; + depthAttachment.loadOp = VK_ATTACHMENT_LOAD_OP_CLEAR; + depthAttachment.storeOp = VK_ATTACHMENT_STORE_OP_DONT_CARE; + depthAttachment.stencilLoadOp = VK_ATTACHMENT_LOAD_OP_DONT_CARE; + depthAttachment.stencilStoreOp = VK_ATTACHMENT_STORE_OP_DONT_CARE; + depthAttachment.initialLayout = VK_IMAGE_LAYOUT_UNDEFINED; + depthAttachment.finalLayout = VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL; + + VkAttachmentReference depthAttachmentRef = {}; + depthAttachmentRef.attachment = 1; + depthAttachmentRef.layout = VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL; + + VkSubpassDescription subpass = {}; + subpass.pipelineBindPoint = VK_PIPELINE_BIND_POINT_GRAPHICS; + subpass.colorAttachmentCount = 1; + subpass.pColorAttachments = &colorAttachmentRef; + subpass.pDepthStencilAttachment = &depthAttachmentRef; + + + VkSubpassDependency dependency = {}; + dependency.srcSubpass = VK_SUBPASS_EXTERNAL; + dependency.dstSubpass = 0; + dependency.srcStageMask = VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT; + dependency.srcAccessMask = 0; + dependency.dstStageMask = VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT; + dependency.dstAccessMask = VK_ACCESS_COLOR_ATTACHMENT_READ_BIT | VK_ACCESS_COLOR_ATTACHMENT_WRITE_BIT; + + VkAttachmentDescription attachments[] = {colorAttachment, depthAttachment}; + VkRenderPassCreateInfo renderPassInfo = {}; + renderPassInfo.sType = VK_STRUCTURE_TYPE_RENDER_PASS_CREATE_INFO; + renderPassInfo.attachmentCount = countof(attachments); + renderPassInfo.pAttachments = attachments; + renderPassInfo.subpassCount = 1; + renderPassInfo.pSubpasses = &subpass; + renderPassInfo.dependencyCount = 1; + renderPassInfo.pDependencies = &dependency; + + if (vkCreateRenderPass(m_device, &renderPassInfo, nullptr, &m_renderPass) != VK_SUCCESS) { + throw std::runtime_error("failed to create render pass!"); + } +} + +void VulkanBaseApp::createDescriptorSetLayout() +{ + VkDescriptorSetLayoutBinding uboLayoutBinding = {}; + uboLayoutBinding.binding = 0; + uboLayoutBinding.descriptorCount = 1; + uboLayoutBinding.descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER; + uboLayoutBinding.pImmutableSamplers = nullptr; + uboLayoutBinding.stageFlags = VK_SHADER_STAGE_VERTEX_BIT; + + VkDescriptorSetLayoutCreateInfo layoutInfo = {}; + layoutInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO; + layoutInfo.bindingCount = 1; + layoutInfo.pBindings = &uboLayoutBinding; + + if (vkCreateDescriptorSetLayout(m_device, &layoutInfo, nullptr, &m_descriptorSetLayout) != VK_SUCCESS) { + throw std::runtime_error("failed to create descriptor set layout!"); + } +} + +VkShaderModule createShaderModule(VkDevice device, const char *filename) +{ + std::vector shaderContents; + std::ifstream shaderFile(filename, std::ios_base::in | std::ios_base::binary); + VkShaderModuleCreateInfo createInfo = {}; + VkShaderModule shaderModule; + + if (!shaderFile.good()) { + throw std::runtime_error("Failed to load shader contents"); + } + readFile(shaderFile, shaderContents); + + createInfo.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO; + createInfo.codeSize = shaderContents.size(); + createInfo.pCode = reinterpret_cast(shaderContents.data()); + VkResult asdf = vkCreateShaderModule(device, &createInfo, nullptr, &shaderModule); + if ( asdf != VK_SUCCESS) { + throw std::runtime_error("Failed to create shader module!"); + } + + return shaderModule; +} + +void VulkanBaseApp::getVertexDescriptions(std::vector& bindingDesc, std::vector& attribDesc) +{ +} + +void VulkanBaseApp::getAssemblyStateInfo(VkPipelineInputAssemblyStateCreateInfo& info) +{ + +} + +void VulkanBaseApp::createGraphicsPipeline() +{ + std::vector shaderStageInfos(m_shaderFiles.size()); + for (size_t i = 0; i < m_shaderFiles.size(); i++) { + shaderStageInfos[i] = {}; + shaderStageInfos[i].sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO; + shaderStageInfos[i].stage = m_shaderFiles[i].first; + shaderStageInfos[i].module = createShaderModule(m_device, m_shaderFiles[i].second.c_str()); + shaderStageInfos[i].pName = "main"; + } + + VkPipelineVertexInputStateCreateInfo vertexInputInfo = {}; + + std::vector vertexBindingDescriptions; + std::vector vertexAttributeDescriptions; + + getVertexDescriptions(vertexBindingDescriptions, vertexAttributeDescriptions); + + vertexInputInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_VERTEX_INPUT_STATE_CREATE_INFO; + vertexInputInfo.vertexBindingDescriptionCount = static_cast(vertexBindingDescriptions.size()); + vertexInputInfo.pVertexBindingDescriptions = vertexBindingDescriptions.data(); + vertexInputInfo.vertexAttributeDescriptionCount = static_cast(vertexAttributeDescriptions.size()); + vertexInputInfo.pVertexAttributeDescriptions = vertexAttributeDescriptions.data(); + + VkPipelineInputAssemblyStateCreateInfo inputAssembly = {}; + getAssemblyStateInfo(inputAssembly); + + VkViewport viewport = {}; + viewport.x = 0.0f; + viewport.y = 0.0f; + viewport.width = (float)m_swapChainExtent.width; + viewport.height = (float)m_swapChainExtent.height; + viewport.minDepth = 0.0f; + viewport.maxDepth = 1.0f; + + VkRect2D scissor = {}; + scissor.offset = { 0, 0 }; + scissor.extent = m_swapChainExtent; + + VkPipelineViewportStateCreateInfo viewportState = {}; + viewportState.sType = VK_STRUCTURE_TYPE_PIPELINE_VIEWPORT_STATE_CREATE_INFO; + viewportState.viewportCount = 1; + viewportState.pViewports = &viewport; + viewportState.scissorCount = 1; + viewportState.pScissors = &scissor; + + VkPipelineRasterizationStateCreateInfo rasterizer = {}; + rasterizer.sType = VK_STRUCTURE_TYPE_PIPELINE_RASTERIZATION_STATE_CREATE_INFO; + rasterizer.depthClampEnable = VK_FALSE; + rasterizer.rasterizerDiscardEnable = VK_FALSE; + rasterizer.polygonMode = VK_POLYGON_MODE_LINE; + rasterizer.lineWidth = 1.0f; + rasterizer.cullMode = VK_CULL_MODE_NONE; + rasterizer.frontFace = VK_FRONT_FACE_CLOCKWISE; + rasterizer.depthBiasEnable = VK_FALSE; + + VkPipelineMultisampleStateCreateInfo multisampling = {}; + multisampling.sType = VK_STRUCTURE_TYPE_PIPELINE_MULTISAMPLE_STATE_CREATE_INFO; + multisampling.sampleShadingEnable = VK_FALSE; + multisampling.rasterizationSamples = VK_SAMPLE_COUNT_1_BIT; + multisampling.minSampleShading = 1.0f; // Optional + multisampling.pSampleMask = nullptr; // Optional + multisampling.alphaToCoverageEnable = VK_FALSE; // Optional + multisampling.alphaToOneEnable = VK_FALSE; // Optional + + VkPipelineDepthStencilStateCreateInfo depthStencil = {}; + depthStencil.sType = VK_STRUCTURE_TYPE_PIPELINE_DEPTH_STENCIL_STATE_CREATE_INFO; + depthStencil.depthTestEnable = VK_TRUE; + depthStencil.depthWriteEnable = VK_TRUE; + depthStencil.depthCompareOp = VK_COMPARE_OP_LESS; + depthStencil.depthBoundsTestEnable = VK_FALSE; + depthStencil.stencilTestEnable = VK_FALSE; + + VkPipelineColorBlendAttachmentState colorBlendAttachment = {}; + colorBlendAttachment.colorWriteMask = VK_COLOR_COMPONENT_R_BIT | VK_COLOR_COMPONENT_G_BIT | VK_COLOR_COMPONENT_B_BIT | VK_COLOR_COMPONENT_A_BIT; + colorBlendAttachment.blendEnable = VK_FALSE; + + VkPipelineColorBlendStateCreateInfo colorBlending = {}; + colorBlending.sType = VK_STRUCTURE_TYPE_PIPELINE_COLOR_BLEND_STATE_CREATE_INFO; + colorBlending.logicOpEnable = VK_FALSE; + colorBlending.logicOp = VK_LOGIC_OP_COPY; + colorBlending.attachmentCount = 1; + colorBlending.pAttachments = &colorBlendAttachment; + colorBlending.blendConstants[0] = 0.0f; + colorBlending.blendConstants[1] = 0.0f; + colorBlending.blendConstants[2] = 0.0f; + colorBlending.blendConstants[3] = 0.0f; + + VkPipelineLayoutCreateInfo pipelineLayoutInfo = {}; + pipelineLayoutInfo.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO; + pipelineLayoutInfo.setLayoutCount = 1; // Optional + pipelineLayoutInfo.pSetLayouts = &m_descriptorSetLayout; // Optional + pipelineLayoutInfo.pushConstantRangeCount = 0; // Optional + pipelineLayoutInfo.pPushConstantRanges = nullptr; // Optional + + if (vkCreatePipelineLayout(m_device, &pipelineLayoutInfo, nullptr, &m_pipelineLayout) != VK_SUCCESS) { + throw std::runtime_error("failed to create pipeline layout!"); + } + + VkGraphicsPipelineCreateInfo pipelineInfo = {}; + pipelineInfo.sType = VK_STRUCTURE_TYPE_GRAPHICS_PIPELINE_CREATE_INFO; + pipelineInfo.stageCount = static_cast(shaderStageInfos.size()); + pipelineInfo.pStages = shaderStageInfos.data(); + + pipelineInfo.pVertexInputState = &vertexInputInfo; + pipelineInfo.pInputAssemblyState = &inputAssembly; + pipelineInfo.pViewportState = &viewportState; + pipelineInfo.pRasterizationState = &rasterizer; + pipelineInfo.pMultisampleState = &multisampling; + pipelineInfo.pDepthStencilState = &depthStencil; // Optional + pipelineInfo.pColorBlendState = &colorBlending; + pipelineInfo.pDynamicState = nullptr; // Optional + + pipelineInfo.layout = m_pipelineLayout; + + pipelineInfo.renderPass = m_renderPass; + pipelineInfo.subpass = 0; + + pipelineInfo.basePipelineHandle = VK_NULL_HANDLE; // Optional + pipelineInfo.basePipelineIndex = -1; // Optional + + if (vkCreateGraphicsPipelines(m_device, VK_NULL_HANDLE, 1, &pipelineInfo, nullptr, &m_graphicsPipeline) != VK_SUCCESS) { + throw std::runtime_error("failed to create graphics pipeline!"); + } + + for (size_t i = 0; i < shaderStageInfos.size(); i++) { + vkDestroyShaderModule(m_device, shaderStageInfos[i].module, nullptr); + } +} + +void VulkanBaseApp::createFramebuffers() +{ + m_swapChainFramebuffers.resize(m_swapChainImageViews.size()); + for (size_t i = 0; i < m_swapChainImageViews.size(); i++) { + VkImageView attachments[] = { + m_swapChainImageViews[i], + m_depthImageView + }; + + VkFramebufferCreateInfo framebufferInfo = {}; + framebufferInfo.sType = VK_STRUCTURE_TYPE_FRAMEBUFFER_CREATE_INFO; + framebufferInfo.renderPass = m_renderPass; + framebufferInfo.attachmentCount = countof(attachments); + framebufferInfo.pAttachments = attachments; + framebufferInfo.width = m_swapChainExtent.width; + framebufferInfo.height = m_swapChainExtent.height; + framebufferInfo.layers = 1; + + if (vkCreateFramebuffer(m_device, &framebufferInfo, nullptr, &m_swapChainFramebuffers[i]) != VK_SUCCESS) { + throw std::runtime_error("failed to create framebuffer!"); + } + } +} + +void VulkanBaseApp::createCommandPool() +{ + VkCommandPoolCreateInfo poolInfo = {}; + uint32_t graphicsIndex, presentIndex; + + findGraphicsQueueIndicies(m_physicalDevice, m_surface, graphicsIndex, presentIndex); + + poolInfo.sType = VK_STRUCTURE_TYPE_COMMAND_POOL_CREATE_INFO; + poolInfo.queueFamilyIndex = graphicsIndex; + poolInfo.flags = 0; // Optional + + if (vkCreateCommandPool(m_device, &poolInfo, nullptr, &m_commandPool) != VK_SUCCESS) { + throw std::runtime_error("Failed to create command pool!"); + } +} + +static void transitionImageLayout(VulkanBaseApp *app, VkImage image, VkFormat format, VkImageLayout oldLayout, VkImageLayout newLayout) +{ + VkCommandBuffer commandBuffer = app->beginSingleTimeCommands(); + + VkImageMemoryBarrier barrier = {}; + barrier.sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER; + barrier.oldLayout = oldLayout; + barrier.newLayout = newLayout; + barrier.srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + barrier.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + barrier.image = image; + + if (newLayout == VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL) { + barrier.subresourceRange.aspectMask = VK_IMAGE_ASPECT_DEPTH_BIT; + + if (format == VK_FORMAT_D32_SFLOAT_S8_UINT || format == VK_FORMAT_D24_UNORM_S8_UINT) { + barrier.subresourceRange.aspectMask |= VK_IMAGE_ASPECT_STENCIL_BIT; + } + } + else { + barrier.subresourceRange.aspectMask = VK_IMAGE_ASPECT_COLOR_BIT; + } + + barrier.subresourceRange.baseMipLevel = 0; + barrier.subresourceRange.levelCount = 1; + barrier.subresourceRange.baseArrayLayer = 0; + barrier.subresourceRange.layerCount = 1; + + VkPipelineStageFlags sourceStage; + VkPipelineStageFlags destinationStage; + + if (oldLayout == VK_IMAGE_LAYOUT_UNDEFINED && newLayout == VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL) { + barrier.srcAccessMask = 0; + barrier.dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT; + + sourceStage = VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT; + destinationStage = VK_PIPELINE_STAGE_TRANSFER_BIT; + } + else if (oldLayout == VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL && newLayout == VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL) { + barrier.srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT; + barrier.dstAccessMask = VK_ACCESS_SHADER_READ_BIT; + + sourceStage = VK_PIPELINE_STAGE_TRANSFER_BIT; + destinationStage = VK_PIPELINE_STAGE_FRAGMENT_SHADER_BIT; + } + else if (oldLayout == VK_IMAGE_LAYOUT_UNDEFINED && newLayout == VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL) { + barrier.srcAccessMask = 0; + barrier.dstAccessMask = VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_READ_BIT | VK_ACCESS_DEPTH_STENCIL_ATTACHMENT_WRITE_BIT; + + sourceStage = VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT; + destinationStage = VK_PIPELINE_STAGE_EARLY_FRAGMENT_TESTS_BIT; + } + else { + throw std::invalid_argument("unsupported layout transition!"); + } + + vkCmdPipelineBarrier( + commandBuffer, + sourceStage, destinationStage, + 0, + 0, nullptr, + 0, nullptr, + 1, &barrier + ); + + app->endSingleTimeCommands(commandBuffer); +} + +void VulkanBaseApp::createDepthResources() +{ + VkFormat depthFormat = findSupportedFormat(m_physicalDevice, + { VK_FORMAT_D32_SFLOAT, VK_FORMAT_D32_SFLOAT_S8_UINT, VK_FORMAT_D24_UNORM_S8_UINT }, + VK_IMAGE_TILING_OPTIMAL, VK_FORMAT_FEATURE_DEPTH_STENCIL_ATTACHMENT_BIT); + createImage(m_physicalDevice, m_device, m_swapChainExtent.width, m_swapChainExtent.height, depthFormat, VK_IMAGE_TILING_OPTIMAL, VK_IMAGE_USAGE_DEPTH_STENCIL_ATTACHMENT_BIT, VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, m_depthImage, m_depthImageMemory); + m_depthImageView = createImageView(m_device, m_depthImage, depthFormat, VK_IMAGE_ASPECT_DEPTH_BIT); + transitionImageLayout(this, m_depthImage, depthFormat, VK_IMAGE_LAYOUT_UNDEFINED, VK_IMAGE_LAYOUT_DEPTH_STENCIL_ATTACHMENT_OPTIMAL); +} + +void VulkanBaseApp::createUniformBuffers() +{ + VkDeviceSize size = getUniformSize(); + if (size > 0) { + m_uniformBuffers.resize(m_swapChainImages.size()); + m_uniformMemory.resize(m_swapChainImages.size()); + for (size_t i = 0; i < m_uniformBuffers.size(); i++) { + createBuffer(getUniformSize(), + VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT, + m_uniformBuffers[i], m_uniformMemory[i]); + } + } +} + +void VulkanBaseApp::createDescriptorPool() +{ + VkDescriptorPoolSize poolSize = {}; + poolSize.type = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER; + poolSize.descriptorCount = static_cast(m_swapChainImages.size()); + VkDescriptorPoolCreateInfo poolInfo = {}; + poolInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO; + poolInfo.poolSizeCount = 1; + poolInfo.pPoolSizes = &poolSize; + poolInfo.maxSets = static_cast(m_swapChainImages.size());; + if (vkCreateDescriptorPool(m_device, &poolInfo, nullptr, &m_descriptorPool) != VK_SUCCESS) { + throw std::runtime_error("failed to create descriptor pool!"); + } +} + +void VulkanBaseApp::createDescriptorSets() +{ + std::vector layouts(m_swapChainImages.size(), m_descriptorSetLayout); + VkDescriptorSetAllocateInfo allocInfo = {}; + allocInfo.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO; + allocInfo.descriptorPool = m_descriptorPool; + allocInfo.descriptorSetCount = static_cast(m_swapChainImages.size()); + allocInfo.pSetLayouts = layouts.data(); + m_descriptorSets.resize(m_swapChainImages.size()); + + if (vkAllocateDescriptorSets(m_device, &allocInfo, m_descriptorSets.data()) != VK_SUCCESS) { + throw std::runtime_error("failed to allocate descriptor sets!"); + } + + VkDescriptorBufferInfo bufferInfo = {}; + bufferInfo.offset = 0; + bufferInfo.range = VK_WHOLE_SIZE; + VkWriteDescriptorSet descriptorWrite = {}; + descriptorWrite.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET; + descriptorWrite.dstBinding = 0; + descriptorWrite.dstArrayElement = 0; + descriptorWrite.descriptorType = VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER; + descriptorWrite.descriptorCount = 1; + descriptorWrite.pBufferInfo = &bufferInfo; + descriptorWrite.pImageInfo = nullptr; // Optional + descriptorWrite.pTexelBufferView = nullptr; // Optional + + for (size_t i = 0; i < m_swapChainImages.size(); i++) { + bufferInfo.buffer = m_uniformBuffers[i]; + descriptorWrite.dstSet = m_descriptorSets[i]; + vkUpdateDescriptorSets(m_device, 1, &descriptorWrite, 0, nullptr); + } +} + +void VulkanBaseApp::createCommandBuffers() +{ + m_commandBuffers.resize(m_swapChainFramebuffers.size()); + VkCommandBufferAllocateInfo allocInfo = {}; + allocInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_ALLOCATE_INFO; + allocInfo.commandPool = m_commandPool; + allocInfo.level = VK_COMMAND_BUFFER_LEVEL_PRIMARY; + allocInfo.commandBufferCount = (uint32_t)m_commandBuffers.size(); + + if (vkAllocateCommandBuffers(m_device, &allocInfo, m_commandBuffers.data()) != VK_SUCCESS) { + throw std::runtime_error("failed to allocate command buffers!"); + } + + for (size_t i = 0; i < m_commandBuffers.size(); i++) { + VkCommandBufferBeginInfo beginInfo = {}; + beginInfo.sType = VK_STRUCTURE_TYPE_COMMAND_BUFFER_BEGIN_INFO; + beginInfo.flags = VK_COMMAND_BUFFER_USAGE_SIMULTANEOUS_USE_BIT; + beginInfo.pInheritanceInfo = nullptr; // Optional + + if (vkBeginCommandBuffer(m_commandBuffers[i], &beginInfo) != VK_SUCCESS) { + throw std::runtime_error("failed to begin recording command buffer!"); + } + + VkRenderPassBeginInfo renderPassInfo = {}; + renderPassInfo.sType = VK_STRUCTURE_TYPE_RENDER_PASS_BEGIN_INFO; + renderPassInfo.renderPass = m_renderPass; + renderPassInfo.framebuffer = m_swapChainFramebuffers[i]; + + renderPassInfo.renderArea.offset = { 0, 0 }; + renderPassInfo.renderArea.extent = m_swapChainExtent; + + VkClearValue clearColors[2]; + clearColors[0].color = { 0.0f, 0.0f, 0.0f, 1.0f }; + clearColors[1].depthStencil = { 1.0f, 0 }; + renderPassInfo.clearValueCount = countof(clearColors); + renderPassInfo.pClearValues = clearColors; + + vkCmdBeginRenderPass(m_commandBuffers[i], &renderPassInfo, VK_SUBPASS_CONTENTS_INLINE); + + vkCmdBindPipeline(m_commandBuffers[i], VK_PIPELINE_BIND_POINT_GRAPHICS, m_graphicsPipeline); + + vkCmdBindDescriptorSets(m_commandBuffers[i], VK_PIPELINE_BIND_POINT_GRAPHICS, m_pipelineLayout, 0, 1, &m_descriptorSets[i], 0, nullptr); + + fillRenderingCommandBuffer(m_commandBuffers[i]); + + vkCmdEndRenderPass(m_commandBuffers[i]); + + if (vkEndCommandBuffer(m_commandBuffers[i]) != VK_SUCCESS) { + throw std::runtime_error("failed to record command buffer!"); + } + } +} + +void VulkanBaseApp::createSyncObjects() +{ + VkSemaphoreCreateInfo semaphoreInfo = {}; + semaphoreInfo.sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO; + VkFenceCreateInfo fenceInfo = {}; + fenceInfo.sType = VK_STRUCTURE_TYPE_FENCE_CREATE_INFO; + fenceInfo.flags = VK_FENCE_CREATE_SIGNALED_BIT; + + m_inFlightFences.resize(MAX_FRAMES_IN_FLIGHT); + m_imageAvailableSemaphores.resize(MAX_FRAMES_IN_FLIGHT); + m_renderFinishedSemaphores.resize(MAX_FRAMES_IN_FLIGHT); + + for (size_t i = 0; i < MAX_FRAMES_IN_FLIGHT; i++) { + if (vkCreateSemaphore(m_device, &semaphoreInfo, nullptr, &m_imageAvailableSemaphores[i]) != VK_SUCCESS) { + throw std::runtime_error("Failed to create image available semaphore!"); + } + if (vkCreateSemaphore(m_device, &semaphoreInfo, nullptr, &m_renderFinishedSemaphores[i]) != VK_SUCCESS) { + throw std::runtime_error("Failed to create image available semaphore!"); + } + if (vkCreateFence(m_device, &fenceInfo, nullptr, &m_inFlightFences[i]) != VK_SUCCESS) { + throw std::runtime_error("Failed to create image available semaphore!"); + } + } +} + +void VulkanBaseApp::getWaitFrameSemaphores(std::vector& wait, std::vector& waitStages) const +{ +} + +void VulkanBaseApp::getSignalFrameSemaphores(std::vector& signal) const +{ +} + +VkDeviceSize VulkanBaseApp::getUniformSize() const +{ + return VkDeviceSize(0); +} + +void VulkanBaseApp::updateUniformBuffer(uint32_t imageIndex) +{ +} + +void VulkanBaseApp::createBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer& buffer, VkDeviceMemory& bufferMemory) +{ + VkBufferCreateInfo bufferInfo = {}; + bufferInfo.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; + bufferInfo.size = size; + bufferInfo.usage = usage; + bufferInfo.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + + if (vkCreateBuffer(m_device, &bufferInfo, nullptr, &buffer) != VK_SUCCESS) { + throw std::runtime_error("failed to create buffer!"); + } + + VkMemoryRequirements memRequirements; + vkGetBufferMemoryRequirements(m_device, buffer, &memRequirements); + + VkMemoryAllocateInfo allocInfo = {}; + allocInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + allocInfo.allocationSize = memRequirements.size; + allocInfo.memoryTypeIndex = findMemoryType(m_physicalDevice, memRequirements.memoryTypeBits, properties); + + if (vkAllocateMemory(m_device, &allocInfo, nullptr, &bufferMemory) != VK_SUCCESS) { + throw std::runtime_error("failed to allocate buffer memory!"); + } + + vkBindBufferMemory(m_device, buffer, bufferMemory, 0); +} + +void VulkanBaseApp::createExternalBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkExternalMemoryHandleTypeFlagsKHR extMemHandleType, VkBuffer& buffer, VkDeviceMemory& bufferMemory) +{ + VkBufferCreateInfo bufferInfo = {}; + bufferInfo.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; + bufferInfo.size = size; + bufferInfo.usage = usage; + bufferInfo.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + + if (vkCreateBuffer(m_device, &bufferInfo, nullptr, &buffer) != VK_SUCCESS) { + throw std::runtime_error("failed to create buffer!"); + } + + VkMemoryRequirements memRequirements; + vkGetBufferMemoryRequirements(m_device, buffer, &memRequirements); + +#ifdef _WIN64 + WindowsSecurityAttributes winSecurityAttributes; + + VkExportMemoryWin32HandleInfoKHR vulkanExportMemoryWin32HandleInfoKHR = {}; + vulkanExportMemoryWin32HandleInfoKHR.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_WIN32_HANDLE_INFO_KHR; + vulkanExportMemoryWin32HandleInfoKHR.pNext = NULL; + vulkanExportMemoryWin32HandleInfoKHR.pAttributes = &winSecurityAttributes; + vulkanExportMemoryWin32HandleInfoKHR.dwAccess = DXGI_SHARED_RESOURCE_READ | DXGI_SHARED_RESOURCE_WRITE; + vulkanExportMemoryWin32HandleInfoKHR.name = (LPCWSTR)NULL; +#endif + VkExportMemoryAllocateInfoKHR vulkanExportMemoryAllocateInfoKHR = {}; + vulkanExportMemoryAllocateInfoKHR.sType = VK_STRUCTURE_TYPE_EXPORT_MEMORY_ALLOCATE_INFO_KHR; +#ifdef _WIN64 + vulkanExportMemoryAllocateInfoKHR.pNext = extMemHandleType & VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT_KHR ? &vulkanExportMemoryWin32HandleInfoKHR : NULL; + vulkanExportMemoryAllocateInfoKHR.handleTypes = extMemHandleType; +#else + vulkanExportMemoryAllocateInfoKHR.pNext = NULL; + vulkanExportMemoryAllocateInfoKHR.handleTypes = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; +#endif + VkMemoryAllocateInfo allocInfo = {}; + allocInfo.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + allocInfo.pNext = &vulkanExportMemoryAllocateInfoKHR; + allocInfo.allocationSize = memRequirements.size; + allocInfo.memoryTypeIndex = findMemoryType(m_physicalDevice, memRequirements.memoryTypeBits, properties); + + if (vkAllocateMemory(m_device, &allocInfo, nullptr, &bufferMemory) != VK_SUCCESS) { + throw std::runtime_error("failed to allocate external buffer memory!"); + } + + vkBindBufferMemory(m_device, buffer, bufferMemory, 0); +} + +void *VulkanBaseApp::getMemHandle(VkDeviceMemory memory, VkExternalMemoryHandleTypeFlagBits handleType) +{ +#ifdef _WIN64 + HANDLE handle = 0; + + VkMemoryGetWin32HandleInfoKHR vkMemoryGetWin32HandleInfoKHR = {}; + vkMemoryGetWin32HandleInfoKHR.sType = VK_STRUCTURE_TYPE_MEMORY_GET_WIN32_HANDLE_INFO_KHR; + vkMemoryGetWin32HandleInfoKHR.pNext = NULL; + vkMemoryGetWin32HandleInfoKHR.memory = memory; + vkMemoryGetWin32HandleInfoKHR.handleType = handleType; + + PFN_vkGetMemoryWin32HandleKHR fpGetMemoryWin32HandleKHR; + fpGetMemoryWin32HandleKHR = (PFN_vkGetMemoryWin32HandleKHR)vkGetDeviceProcAddr(m_device, "vkGetMemoryWin32HandleKHR"); + if (!fpGetMemoryWin32HandleKHR) { + throw std::runtime_error("Failed to retrieve vkGetMemoryWin32HandleKHR!"); + } + if (fpGetMemoryWin32HandleKHR(m_device, &vkMemoryGetWin32HandleInfoKHR, &handle) != VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve handle for buffer!"); + } + return (void *)handle; +#else + int fd = -1; + + VkMemoryGetFdInfoKHR vkMemoryGetFdInfoKHR = {}; + vkMemoryGetFdInfoKHR.sType = VK_STRUCTURE_TYPE_MEMORY_GET_FD_INFO_KHR; + vkMemoryGetFdInfoKHR.pNext = NULL; + vkMemoryGetFdInfoKHR.memory = memory; + vkMemoryGetFdInfoKHR.handleType = handleType; + + PFN_vkGetMemoryFdKHR fpGetMemoryFdKHR; + fpGetMemoryFdKHR = (PFN_vkGetMemoryFdKHR)vkGetDeviceProcAddr(m_device, "vkGetMemoryFdKHR"); + if (!fpGetMemoryFdKHR) { + throw std::runtime_error("Failed to retrieve vkGetMemoryWin32HandleKHR!"); + } + if (fpGetMemoryFdKHR(m_device, &vkMemoryGetFdInfoKHR, &fd) != VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve handle for buffer!"); + } + return (void *)(uintptr_t)fd; +#endif /* _WIN64 */ +} + +void *VulkanBaseApp::getSemaphoreHandle(VkSemaphore semaphore, VkExternalSemaphoreHandleTypeFlagBits handleType) +{ +#ifdef _WIN64 + HANDLE handle; + + VkSemaphoreGetWin32HandleInfoKHR semaphoreGetWin32HandleInfoKHR = {}; + semaphoreGetWin32HandleInfoKHR.sType = VK_STRUCTURE_TYPE_SEMAPHORE_GET_WIN32_HANDLE_INFO_KHR; + semaphoreGetWin32HandleInfoKHR.pNext = NULL; + semaphoreGetWin32HandleInfoKHR.semaphore = semaphore; + semaphoreGetWin32HandleInfoKHR.handleType = handleType; + + PFN_vkGetSemaphoreWin32HandleKHR fpGetSemaphoreWin32HandleKHR; + fpGetSemaphoreWin32HandleKHR = (PFN_vkGetSemaphoreWin32HandleKHR)vkGetDeviceProcAddr(m_device, "vkGetSemaphoreWin32HandleKHR"); + if (!fpGetSemaphoreWin32HandleKHR) { + throw std::runtime_error("Failed to retrieve vkGetMemoryWin32HandleKHR!"); + } + if (fpGetSemaphoreWin32HandleKHR(m_device, &semaphoreGetWin32HandleInfoKHR, &handle) != VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve handle for buffer!"); + } + + return (void *)handle; +#else + int fd; + + VkSemaphoreGetFdInfoKHR semaphoreGetFdInfoKHR = {}; + semaphoreGetFdInfoKHR.sType =VK_STRUCTURE_TYPE_SEMAPHORE_GET_FD_INFO_KHR; + semaphoreGetFdInfoKHR.pNext = NULL; + semaphoreGetFdInfoKHR.semaphore = semaphore; + semaphoreGetFdInfoKHR.handleType = handleType; + + PFN_vkGetSemaphoreFdKHR fpGetSemaphoreFdKHR; + fpGetSemaphoreFdKHR = (PFN_vkGetSemaphoreFdKHR)vkGetDeviceProcAddr(m_device, "vkGetSemaphoreFdKHR"); + if (!fpGetSemaphoreFdKHR) { + throw std::runtime_error("Failed to retrieve vkGetMemoryWin32HandleKHR!"); + } + if (fpGetSemaphoreFdKHR(m_device, &semaphoreGetFdInfoKHR, &fd) != VK_SUCCESS) { + throw std::runtime_error("Failed to retrieve handle for buffer!"); + } + + return (void *)(uintptr_t)fd; +#endif +} + +void VulkanBaseApp::createExternalSemaphore(VkSemaphore& semaphore, VkExternalSemaphoreHandleTypeFlagBits handleType) +{ + VkSemaphoreCreateInfo semaphoreInfo = {}; + semaphoreInfo.sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO; + VkExportSemaphoreCreateInfoKHR exportSemaphoreCreateInfo = {}; + exportSemaphoreCreateInfo.sType = VK_STRUCTURE_TYPE_EXPORT_SEMAPHORE_CREATE_INFO_KHR; + +#ifdef _WIN64 + WindowsSecurityAttributes winSecurityAttributes; + + VkExportSemaphoreWin32HandleInfoKHR exportSemaphoreWin32HandleInfoKHR = {}; + exportSemaphoreWin32HandleInfoKHR.sType = VK_STRUCTURE_TYPE_EXPORT_SEMAPHORE_WIN32_HANDLE_INFO_KHR; + exportSemaphoreWin32HandleInfoKHR.pNext = NULL; + exportSemaphoreWin32HandleInfoKHR.pAttributes = &winSecurityAttributes; + exportSemaphoreWin32HandleInfoKHR.dwAccess = DXGI_SHARED_RESOURCE_READ | DXGI_SHARED_RESOURCE_WRITE; + exportSemaphoreWin32HandleInfoKHR.name = (LPCWSTR)NULL; + exportSemaphoreCreateInfo.pNext = (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) ? &exportSemaphoreWin32HandleInfoKHR : NULL; +#else + exportSemaphoreCreateInfo.pNext = NULL; +#endif + exportSemaphoreCreateInfo.handleTypes = handleType; + semaphoreInfo.pNext = &exportSemaphoreCreateInfo; + + if (vkCreateSemaphore(m_device, &semaphoreInfo, nullptr, &semaphore) != VK_SUCCESS) { + throw std::runtime_error("failed to create synchronization objects for a CUDA-Vulkan!"); + } +} + +void VulkanBaseApp::importExternalBuffer(void *handle, VkExternalMemoryHandleTypeFlagBits handleType, size_t size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer& buffer, VkDeviceMemory& memory) +{ + VkBufferCreateInfo bufferInfo = {}; + bufferInfo.sType = VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO; + bufferInfo.size = size; + bufferInfo.usage = usage; + bufferInfo.sharingMode = VK_SHARING_MODE_EXCLUSIVE; + + if (vkCreateBuffer(m_device, &bufferInfo, nullptr, &buffer) != VK_SUCCESS) { + throw std::runtime_error("failed to create buffer!"); + } + + VkMemoryRequirements memRequirements; + vkGetBufferMemoryRequirements(m_device, buffer, &memRequirements); + +#ifdef _WIN64 + VkImportMemoryWin32HandleInfoKHR handleInfo = {}; + handleInfo.sType = VK_STRUCTURE_TYPE_IMPORT_MEMORY_WIN32_HANDLE_INFO_KHR; + handleInfo.pNext = NULL; + handleInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_WIN32_BIT; + handleInfo.handle = handle; + handleInfo.name = NULL; +#else + VkImportMemoryFdInfoKHR handleInfo = {}; + handleInfo.sType = VK_STRUCTURE_TYPE_IMPORT_MEMORY_FD_INFO_KHR; + handleInfo.pNext = NULL; + handleInfo.fd = (int)(uintptr_t)handle; + handleInfo.handleType = VK_EXTERNAL_MEMORY_HANDLE_TYPE_OPAQUE_FD_BIT; +#endif /* _WIN64 */ + + VkMemoryAllocateInfo memAllocation = {}; + memAllocation.sType = VK_STRUCTURE_TYPE_MEMORY_ALLOCATE_INFO; + memAllocation.pNext = (void *)&handleInfo; + memAllocation.allocationSize = size; + memAllocation.memoryTypeIndex = findMemoryType(m_physicalDevice, memRequirements.memoryTypeBits, properties); + + if (vkAllocateMemory(m_device, &memAllocation, nullptr, &memory) != VK_SUCCESS) { + throw std::runtime_error("Failed to import allocation!"); + } + + vkBindBufferMemory(m_device, buffer, memory, 0); +} + +void VulkanBaseApp::copyBuffer(VkBuffer dst, VkBuffer src, VkDeviceSize size) +{ + + VkCommandBuffer commandBuffer = beginSingleTimeCommands(); + + VkBufferCopy copyRegion = {}; + copyRegion.size = size; + vkCmdCopyBuffer(commandBuffer, src, dst, 1, ©Region); + + endSingleTimeCommands(commandBuffer); +} + +void VulkanBaseApp::drawFrame() +{ + size_t currentFrameIdx = m_currentFrame % MAX_FRAMES_IN_FLIGHT; + vkWaitForFences(m_device, 1, &m_inFlightFences[currentFrameIdx], VK_TRUE, std::numeric_limits::max()); + + uint32_t imageIndex; + VkResult result = vkAcquireNextImageKHR(m_device, m_swapChain, std::numeric_limits::max(), m_imageAvailableSemaphores[currentFrameIdx], VK_NULL_HANDLE, &imageIndex); + if (result == VK_ERROR_OUT_OF_DATE_KHR) { + recreateSwapChain(); + } + else if (result != VK_SUCCESS && result != VK_SUBOPTIMAL_KHR) { + throw std::runtime_error("Failed to acquire swap chain image!"); + } + + updateUniformBuffer(imageIndex); + + VkSubmitInfo submitInfo = {}; + submitInfo.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO; + + std::vector waitSemaphores; + std::vector waitStages; + + waitSemaphores.push_back(m_imageAvailableSemaphores[currentFrameIdx]); + waitStages.push_back(VK_PIPELINE_STAGE_COLOR_ATTACHMENT_OUTPUT_BIT); + //getWaitFrameSemaphores(waitSemaphores, waitStages); + + submitInfo.waitSemaphoreCount = (uint32_t)waitSemaphores.size(); + submitInfo.pWaitSemaphores = waitSemaphores.data(); + submitInfo.pWaitDstStageMask = waitStages.data(); + + submitInfo.commandBufferCount = 1; + submitInfo.pCommandBuffers = &m_commandBuffers[imageIndex]; + + std::vector signalSemaphores; + //getSignalFrameSemaphores(signalSemaphores); + signalSemaphores.push_back(m_renderFinishedSemaphores[currentFrameIdx]); + submitInfo.signalSemaphoreCount = (uint32_t)signalSemaphores.size(); + submitInfo.pSignalSemaphores = signalSemaphores.data(); + + vkResetFences(m_device, 1, &m_inFlightFences[currentFrameIdx]); + + if (vkQueueSubmit(m_graphicsQueue, 1, &submitInfo, m_inFlightFences[currentFrameIdx]) != VK_SUCCESS) { + throw std::runtime_error("failed to submit draw command buffer!"); + } + + VkPresentInfoKHR presentInfo = {}; + presentInfo.sType = VK_STRUCTURE_TYPE_PRESENT_INFO_KHR; + presentInfo.waitSemaphoreCount = 1; + presentInfo.pWaitSemaphores = &m_renderFinishedSemaphores[currentFrameIdx]; + + VkSwapchainKHR swapChains[] = { m_swapChain }; + presentInfo.swapchainCount = 1; + presentInfo.pSwapchains = swapChains; + presentInfo.pImageIndices = &imageIndex; + + result = vkQueuePresentKHR(m_presentQueue, &presentInfo); + if (result == VK_ERROR_OUT_OF_DATE_KHR || result == VK_SUBOPTIMAL_KHR || m_framebufferResized) { + recreateSwapChain(); + m_framebufferResized = false; + } + else if (result != VK_SUCCESS) { + throw std::runtime_error("Failed to acquire swap chain image!"); + } + + m_currentFrame++; +} + +void VulkanBaseApp::cleanupSwapChain() +{ + + if (m_depthImageView != VK_NULL_HANDLE) { + vkDestroyImageView(m_device, m_depthImageView, nullptr); + } + if (m_depthImage != VK_NULL_HANDLE) { + vkDestroyImage(m_device, m_depthImage, nullptr); + } + if (m_depthImageMemory != VK_NULL_HANDLE) { + vkFreeMemory(m_device, m_depthImageMemory, nullptr); + } + + for (size_t i = 0; i < m_uniformBuffers.size(); i++) { + vkDestroyBuffer(m_device, m_uniformBuffers[i], nullptr); + vkFreeMemory(m_device, m_uniformMemory[i], nullptr); + } + + if (m_descriptorPool != VK_NULL_HANDLE) { + vkDestroyDescriptorPool(m_device, m_descriptorPool, nullptr); + } + + for (size_t i = 0; i < m_swapChainFramebuffers.size(); i++) { + vkDestroyFramebuffer(m_device, m_swapChainFramebuffers[i], nullptr); + } + + if (m_graphicsPipeline != VK_NULL_HANDLE) { + vkDestroyPipeline(m_device, m_graphicsPipeline, nullptr); + } + + if (m_pipelineLayout != VK_NULL_HANDLE) { + vkDestroyPipelineLayout(m_device, m_pipelineLayout, nullptr); + } + + if (m_renderPass != VK_NULL_HANDLE) { + vkDestroyRenderPass(m_device, m_renderPass, nullptr); + } + + for (size_t i = 0; i < m_swapChainImageViews.size(); i++) { + vkDestroyImageView(m_device, m_swapChainImageViews[i], nullptr); + } + + if (m_swapChain != VK_NULL_HANDLE) { + vkDestroySwapchainKHR(m_device, m_swapChain, nullptr); + } +} + +void VulkanBaseApp::recreateSwapChain() +{ + int width, height; + + glfwGetFramebufferSize(m_window, &width, &height); + while (width == 0 || height == 0) { + glfwWaitEvents(); + glfwGetFramebufferSize(m_window, &width, &height); + } + + vkDeviceWaitIdle(m_device); + + cleanupSwapChain(); + + createSwapChain(); + createImageViews(); + createRenderPass(); + createGraphicsPipeline(); + createDepthResources(); + createFramebuffers(); + createUniformBuffers(); + createDescriptorPool(); + createDescriptorSets(); + createCommandBuffers(); +} + +void VulkanBaseApp::mainLoop() +{ + while (!glfwWindowShouldClose(m_window)) { + glfwPollEvents(); + drawFrame(); + } + vkDeviceWaitIdle(m_device); +} + +void readFile(std::istream& s, std::vector& data) +{ + s.seekg(0, std::ios_base::end); + data.resize(s.tellg()); + s.clear(); + s.seekg(0, std::ios_base::beg); + s.read(data.data(), data.size()); +} diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.h b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.h new file mode 100644 index 0000000000..03007db309 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/VulkanBaseApp.h @@ -0,0 +1,146 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + + +/* + * Modifications Copyright (C)2021 Advanced + * Micro Devices, Inc. All rights reserved. + */ + +#pragma once +#ifndef __VULKANBASEAPP_H__ +#define __VULKANBASEAPP_H__ + +#include +#include +#include +#ifdef _WIN64 +#define NOMINMAX +#include +#include +#endif /* _WIN64 */ + +struct GLFWwindow; + +class VulkanBaseApp +{ +public: + VulkanBaseApp(const std::string& appName, bool enableValidation = false); + static VkExternalSemaphoreHandleTypeFlagBits getDefaultSemaphoreHandleType(); + static VkExternalMemoryHandleTypeFlagBits getDefaultMemHandleType(); + virtual ~VulkanBaseApp(); + void init(); + void *getMemHandle(VkDeviceMemory memory, VkExternalMemoryHandleTypeFlagBits handleType); + void *getSemaphoreHandle(VkSemaphore semaphore, VkExternalSemaphoreHandleTypeFlagBits handleType); + void createExternalSemaphore(VkSemaphore& semaphore, VkExternalSemaphoreHandleTypeFlagBits handleType); + void createBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer& buffer, VkDeviceMemory& bufferMemory); + void createExternalBuffer(VkDeviceSize size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkExternalMemoryHandleTypeFlagsKHR extMemHandleType, VkBuffer& buffer, VkDeviceMemory& bufferMemory); + void importExternalBuffer(void *handle, VkExternalMemoryHandleTypeFlagBits handleType, size_t size, VkBufferUsageFlags usage, VkMemoryPropertyFlags properties, VkBuffer& buffer, VkDeviceMemory& memory); + void copyBuffer(VkBuffer dst, VkBuffer src, VkDeviceSize size); + VkCommandBuffer beginSingleTimeCommands(); + void endSingleTimeCommands(VkCommandBuffer commandBuffer); + void mainLoop(); +protected: + const std::string m_appName; + const bool m_enableValidation; + VkInstance m_instance; + VkDebugUtilsMessengerEXT m_debugMessenger; + VkSurfaceKHR m_surface; + VkPhysicalDevice m_physicalDevice; + VkDevice m_device; + VkQueue m_graphicsQueue; + VkQueue m_presentQueue; + VkSwapchainKHR m_swapChain; + std::vector m_swapChainImages; + VkFormat m_swapChainFormat; + VkExtent2D m_swapChainExtent; + std::vector m_swapChainImageViews; + std::vector > m_shaderFiles; + VkRenderPass m_renderPass; + VkPipelineLayout m_pipelineLayout; + VkPipeline m_graphicsPipeline; + std::vector m_swapChainFramebuffers; + VkCommandPool m_commandPool; + std::vector m_commandBuffers; + std::vector m_imageAvailableSemaphores; + std::vector m_renderFinishedSemaphores; + std::vector m_inFlightFences; + std::vector m_uniformBuffers; + std::vector m_uniformMemory; + VkDescriptorSetLayout m_descriptorSetLayout; + VkDescriptorPool m_descriptorPool; + std::vector m_descriptorSets; + VkImage m_depthImage; + VkDeviceMemory m_depthImageMemory; + VkImageView m_depthImageView; + size_t m_currentFrame; + bool m_framebufferResized; + uint8_t m_vkDeviceUUID[VK_UUID_SIZE]; + + virtual void initVulkanApp() {} + virtual void fillRenderingCommandBuffer(VkCommandBuffer& buffer) {} + virtual std::vector getRequiredExtensions() const; + virtual std::vector getRequiredDeviceExtensions() const; + virtual void getVertexDescriptions(std::vector& bindingDesc, std::vector& attribDesc); + virtual void getAssemblyStateInfo(VkPipelineInputAssemblyStateCreateInfo& info); + virtual void getWaitFrameSemaphores(std::vector& wait, std::vector< VkPipelineStageFlags>& waitStages) const; + virtual void getSignalFrameSemaphores(std::vector& signal) const; + virtual VkDeviceSize getUniformSize() const; + virtual void updateUniformBuffer(uint32_t imageIndex); + virtual void drawFrame(); +private: + GLFWwindow *m_window; + + void initWindow(); + void initVulkan(); + void createInstance(); + void createSurface(); + void createDevice(); + void createSwapChain(); + void createImageViews(); + void createRenderPass(); + void createDescriptorSetLayout(); + void createGraphicsPipeline(); + void createFramebuffers(); + void createCommandPool(); + void createDepthResources(); + void createUniformBuffers(); + void createDescriptorPool(); + void createDescriptorSets(); + void createCommandBuffers(); + void createSyncObjects(); + + void cleanupSwapChain(); + void recreateSwapChain(); + + bool isSuitableDevice(VkPhysicalDevice dev) const; + static void resizeCallback(GLFWwindow *window, int width, int height); +}; + +void readFile(std::istream& s, std::vector& data); + +#endif /* __VULKANBASEAPP_H__ */ diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt new file mode 100644 index 0000000000..b847596510 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/buildcmd.txt @@ -0,0 +1,11 @@ +• Install hip and visual studio +• Install vulkan sdk from vulkan.lunarg.com +• Download GLFW binaries from glfw.org +• Convert sinwave.farg and vert files to spv +o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.vert -V -o vert.spv +o c:\VulkanSDK\1.2.182.0\bin\glslangValidator.exe sinewave.frag -V -o frag.spv + +• set HCC_AMDGPU_TARGET=gfx906:sramecc-:xnack- (for your graphic card, you can get the name from hipinfo ) +$• hipcc -v *.cpp *.hip -Ic:\VulkanSDK\1.2.182.0\include -L c:\VulkanSDK\1.2.182.0\lib -Ic:\glfw-3.3.4.bin.WIN64\include -L c:\glfw-3.3.4.bin.WIN64\lib-vc2019 -Ic:\hip\include\hip -lglfw3dll -lvulkan-1 -ladvapi32 -std=c++14 +• run a.exe, you should see a 3D sinewave simulation + diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/linmath.h b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/linmath.h new file mode 100644 index 0000000000..dbedbc163a --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/linmath.h @@ -0,0 +1,502 @@ +/* + * Copyright (c) 2015-2016 The Khronos Group Inc. + * Copyright (c) 2015-2016 Valve Corporation + * Copyright (c) 2015-2016 LunarG, Inc. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + * + * Relicensed from the WTFPL (http://www.wtfpl.net/faq/). + */ + +#ifndef LINMATH_H +#define LINMATH_H + +#define _USE_MATH_DEFINES +#include + +// Converts degrees to radians. +#define degreesToRadians(angleDegrees) (angleDegrees * M_PI / 180.0) + +// Converts radians to degrees. +#define radiansToDegrees(angleRadians) (angleRadians * 180.0 / M_PI) + +typedef float vec3[3]; +static inline void vec3_add(vec3 r, vec3 const a, vec3 const b) { + int i; + for (i = 0; i < 3; ++i) r[i] = a[i] + b[i]; +} +static inline void vec3_sub(vec3 r, vec3 const a, vec3 const b) { + int i; + for (i = 0; i < 3; ++i) r[i] = a[i] - b[i]; +} +static inline void vec3_scale(vec3 r, vec3 const v, float const s) { + int i; + for (i = 0; i < 3; ++i) r[i] = v[i] * s; +} +static inline float vec3_mul_inner(vec3 const a, vec3 const b) { + float p = 0.f; + int i; + for (i = 0; i < 3; ++i) p += b[i] * a[i]; + return p; +} +static inline void vec3_mul_cross(vec3 r, vec3 const a, vec3 const b) { + r[0] = a[1] * b[2] - a[2] * b[1]; + r[1] = a[2] * b[0] - a[0] * b[2]; + r[2] = a[0] * b[1] - a[1] * b[0]; +} +static inline float vec3_len(vec3 const v) { return sqrtf(vec3_mul_inner(v, v)); } +static inline void vec3_norm(vec3 r, vec3 const v) { + float k = 1.f / vec3_len(v); + vec3_scale(r, v, k); +} +static inline void vec3_reflect(vec3 r, vec3 const v, vec3 const n) { + float p = 2.f * vec3_mul_inner(v, n); + int i; + for (i = 0; i < 3; ++i) r[i] = v[i] - p * n[i]; +} + +typedef float vec4[4]; +static inline void vec4_add(vec4 r, vec4 const a, vec4 const b) { + int i; + for (i = 0; i < 4; ++i) r[i] = a[i] + b[i]; +} +static inline void vec4_sub(vec4 r, vec4 const a, vec4 const b) { + int i; + for (i = 0; i < 4; ++i) r[i] = a[i] - b[i]; +} +static inline void vec4_scale(vec4 r, vec4 v, float s) { + int i; + for (i = 0; i < 4; ++i) r[i] = v[i] * s; +} +static inline float vec4_mul_inner(vec4 a, vec4 b) { + float p = 0.f; + int i; + for (i = 0; i < 4; ++i) p += b[i] * a[i]; + return p; +} +static inline void vec4_mul_cross(vec4 r, vec4 a, vec4 b) { + r[0] = a[1] * b[2] - a[2] * b[1]; + r[1] = a[2] * b[0] - a[0] * b[2]; + r[2] = a[0] * b[1] - a[1] * b[0]; + r[3] = 1.f; +} +static inline float vec4_len(vec4 v) { return sqrtf(vec4_mul_inner(v, v)); } +static inline void vec4_norm(vec4 r, vec4 v) { + float k = 1.f / vec4_len(v); + vec4_scale(r, v, k); +} +static inline void vec4_reflect(vec4 r, vec4 v, vec4 n) { + float p = 2.f * vec4_mul_inner(v, n); + int i; + for (i = 0; i < 4; ++i) r[i] = v[i] - p * n[i]; +} + +typedef vec4 mat4x4[4]; +static inline void mat4x4_identity(mat4x4 M) { + int i, j; + for (i = 0; i < 4; ++i) + for (j = 0; j < 4; ++j) M[i][j] = i == j ? 1.f : 0.f; +} +static inline void mat4x4_dup(mat4x4 M, mat4x4 N) { + int i, j; + for (i = 0; i < 4; ++i) + for (j = 0; j < 4; ++j) M[i][j] = N[i][j]; +} +static inline void mat4x4_row(vec4 r, mat4x4 M, int i) { + int k; + for (k = 0; k < 4; ++k) r[k] = M[k][i]; +} +static inline void mat4x4_col(vec4 r, mat4x4 M, int i) { + int k; + for (k = 0; k < 4; ++k) r[k] = M[i][k]; +} +static inline void mat4x4_transpose(mat4x4 M, mat4x4 N) { + int i, j; + for (j = 0; j < 4; ++j) + for (i = 0; i < 4; ++i) M[i][j] = N[j][i]; +} +static inline void mat4x4_add(mat4x4 M, mat4x4 a, mat4x4 b) { + int i; + for (i = 0; i < 4; ++i) vec4_add(M[i], a[i], b[i]); +} +static inline void mat4x4_sub(mat4x4 M, mat4x4 a, mat4x4 b) { + int i; + for (i = 0; i < 4; ++i) vec4_sub(M[i], a[i], b[i]); +} +static inline void mat4x4_scale(mat4x4 M, mat4x4 a, float k) { + int i; + for (i = 0; i < 4; ++i) vec4_scale(M[i], a[i], k); +} +static inline void mat4x4_scale_aniso(mat4x4 M, mat4x4 a, float x, float y, float z) { + int i; + vec4_scale(M[0], a[0], x); + vec4_scale(M[1], a[1], y); + vec4_scale(M[2], a[2], z); + for (i = 0; i < 4; ++i) { + M[3][i] = a[3][i]; + } +} +static inline void mat4x4_mul(mat4x4 M, mat4x4 a, mat4x4 b) { + int k, r, c; + for (c = 0; c < 4; ++c) + for (r = 0; r < 4; ++r) { + M[c][r] = 0.f; + for (k = 0; k < 4; ++k) M[c][r] += a[k][r] * b[c][k]; + } +} +static inline void mat4x4_mul_vec4(vec4 r, mat4x4 M, vec4 v) { + int i, j; + for (j = 0; j < 4; ++j) { + r[j] = 0.f; + for (i = 0; i < 4; ++i) r[j] += M[i][j] * v[i]; + } +} +static inline void mat4x4_translate(mat4x4 T, float x, float y, float z) { + mat4x4_identity(T); + T[3][0] = x; + T[3][1] = y; + T[3][2] = z; +} +static inline void mat4x4_translate_in_place(mat4x4 M, float x, float y, float z) { + vec4 t = {x, y, z, 0}; + vec4 r; + int i; + for (i = 0; i < 4; ++i) { + mat4x4_row(r, M, i); + M[3][i] += vec4_mul_inner(r, t); + } +} +static inline void mat4x4_from_vec3_mul_outer(mat4x4 M, vec3 a, vec3 b) { + int i, j; + for (i = 0; i < 4; ++i) + for (j = 0; j < 4; ++j) M[i][j] = i < 3 && j < 3 ? a[i] * b[j] : 0.f; +} +static inline void mat4x4_rotate(mat4x4 R, mat4x4 M, float x, float y, float z, float angle) { + float s = sinf(angle); + float c = cosf(angle); + vec3 u = {x, y, z}; + + if (vec3_len(u) > 1e-4) { + vec3_norm(u, u); + mat4x4 T; + mat4x4_from_vec3_mul_outer(T, u, u); + + mat4x4 S = {{0, u[2], -u[1], 0}, {-u[2], 0, u[0], 0}, {u[1], -u[0], 0, 0}, {0, 0, 0, 0}}; + mat4x4_scale(S, S, s); + + mat4x4 C; + mat4x4_identity(C); + mat4x4_sub(C, C, T); + + mat4x4_scale(C, C, c); + + mat4x4_add(T, T, C); + mat4x4_add(T, T, S); + + T[3][3] = 1.; + mat4x4_mul(R, M, T); + } else { + mat4x4_dup(R, M); + } +} +static inline void mat4x4_rotate_X(mat4x4 Q, mat4x4 M, float angle) { + float s = sinf(angle); + float c = cosf(angle); + mat4x4 R = {{1.f, 0.f, 0.f, 0.f}, {0.f, c, s, 0.f}, {0.f, -s, c, 0.f}, {0.f, 0.f, 0.f, 1.f}}; + mat4x4_mul(Q, M, R); +} +static inline void mat4x4_rotate_Y(mat4x4 Q, mat4x4 M, float angle) { + float s = sinf(angle); + float c = cosf(angle); + mat4x4 R = {{c, 0.f, s, 0.f}, {0.f, 1.f, 0.f, 0.f}, {-s, 0.f, c, 0.f}, {0.f, 0.f, 0.f, 1.f}}; + mat4x4_mul(Q, M, R); +} +static inline void mat4x4_rotate_Z(mat4x4 Q, mat4x4 M, float angle) { + float s = sinf(angle); + float c = cosf(angle); + mat4x4 R = {{c, s, 0.f, 0.f}, {-s, c, 0.f, 0.f}, {0.f, 0.f, 1.f, 0.f}, {0.f, 0.f, 0.f, 1.f}}; + mat4x4_mul(Q, M, R); +} +static inline void mat4x4_invert(mat4x4 T, mat4x4 M) { + float s[6]; + float c[6]; + s[0] = M[0][0] * M[1][1] - M[1][0] * M[0][1]; + s[1] = M[0][0] * M[1][2] - M[1][0] * M[0][2]; + s[2] = M[0][0] * M[1][3] - M[1][0] * M[0][3]; + s[3] = M[0][1] * M[1][2] - M[1][1] * M[0][2]; + s[4] = M[0][1] * M[1][3] - M[1][1] * M[0][3]; + s[5] = M[0][2] * M[1][3] - M[1][2] * M[0][3]; + + c[0] = M[2][0] * M[3][1] - M[3][0] * M[2][1]; + c[1] = M[2][0] * M[3][2] - M[3][0] * M[2][2]; + c[2] = M[2][0] * M[3][3] - M[3][0] * M[2][3]; + c[3] = M[2][1] * M[3][2] - M[3][1] * M[2][2]; + c[4] = M[2][1] * M[3][3] - M[3][1] * M[2][3]; + c[5] = M[2][2] * M[3][3] - M[3][2] * M[2][3]; + + /* Assumes it is invertible */ + float idet = 1.0f / (s[0] * c[5] - s[1] * c[4] + s[2] * c[3] + s[3] * c[2] - s[4] * c[1] + s[5] * c[0]); + + T[0][0] = (M[1][1] * c[5] - M[1][2] * c[4] + M[1][3] * c[3]) * idet; + T[0][1] = (-M[0][1] * c[5] + M[0][2] * c[4] - M[0][3] * c[3]) * idet; + T[0][2] = (M[3][1] * s[5] - M[3][2] * s[4] + M[3][3] * s[3]) * idet; + T[0][3] = (-M[2][1] * s[5] + M[2][2] * s[4] - M[2][3] * s[3]) * idet; + + T[1][0] = (-M[1][0] * c[5] + M[1][2] * c[2] - M[1][3] * c[1]) * idet; + T[1][1] = (M[0][0] * c[5] - M[0][2] * c[2] + M[0][3] * c[1]) * idet; + T[1][2] = (-M[3][0] * s[5] + M[3][2] * s[2] - M[3][3] * s[1]) * idet; + T[1][3] = (M[2][0] * s[5] - M[2][2] * s[2] + M[2][3] * s[1]) * idet; + + T[2][0] = (M[1][0] * c[4] - M[1][1] * c[2] + M[1][3] * c[0]) * idet; + T[2][1] = (-M[0][0] * c[4] + M[0][1] * c[2] - M[0][3] * c[0]) * idet; + T[2][2] = (M[3][0] * s[4] - M[3][1] * s[2] + M[3][3] * s[0]) * idet; + T[2][3] = (-M[2][0] * s[4] + M[2][1] * s[2] - M[2][3] * s[0]) * idet; + + T[3][0] = (-M[1][0] * c[3] + M[1][1] * c[1] - M[1][2] * c[0]) * idet; + T[3][1] = (M[0][0] * c[3] - M[0][1] * c[1] + M[0][2] * c[0]) * idet; + T[3][2] = (-M[3][0] * s[3] + M[3][1] * s[1] - M[3][2] * s[0]) * idet; + T[3][3] = (M[2][0] * s[3] - M[2][1] * s[1] + M[2][2] * s[0]) * idet; +} +static inline void mat4x4_orthonormalize(mat4x4 R, mat4x4 M) { + mat4x4_dup(R, M); + float s = 1.; + vec3 h; + + vec3_norm(R[2], R[2]); + + s = vec3_mul_inner(R[1], R[2]); + vec3_scale(h, R[2], s); + vec3_sub(R[1], R[1], h); + vec3_norm(R[2], R[2]); + + s = vec3_mul_inner(R[1], R[2]); + vec3_scale(h, R[2], s); + vec3_sub(R[1], R[1], h); + vec3_norm(R[1], R[1]); + + s = vec3_mul_inner(R[0], R[1]); + vec3_scale(h, R[1], s); + vec3_sub(R[0], R[0], h); + vec3_norm(R[0], R[0]); +} + +static inline void mat4x4_frustum(mat4x4 M, float l, float r, float b, float t, float n, float f) { + M[0][0] = 2.f * n / (r - l); + M[0][1] = M[0][2] = M[0][3] = 0.f; + + M[1][1] = 2.f * n / (t - b); + M[1][0] = M[1][2] = M[1][3] = 0.f; + + M[2][0] = (r + l) / (r - l); + M[2][1] = (t + b) / (t - b); + M[2][2] = -(f + n) / (f - n); + M[2][3] = -1.f; + + M[3][2] = -2.f * (f * n) / (f - n); + M[3][0] = M[3][1] = M[3][3] = 0.f; +} +static inline void mat4x4_ortho(mat4x4 M, float l, float r, float b, float t, float n, float f) { + M[0][0] = 2.f / (r - l); + M[0][1] = M[0][2] = M[0][3] = 0.f; + + M[1][1] = 2.f / (t - b); + M[1][0] = M[1][2] = M[1][3] = 0.f; + + M[2][2] = -2.f / (f - n); + M[2][0] = M[2][1] = M[2][3] = 0.f; + + M[3][0] = -(r + l) / (r - l); + M[3][1] = -(t + b) / (t - b); + M[3][2] = -(f + n) / (f - n); + M[3][3] = 1.f; +} +static inline void mat4x4_perspective(mat4x4 m, float y_fov, float aspect, float n, float f) { + /* NOTE: Degrees are an unhandy unit to work with. + * linmath.h uses radians for everything! */ + float const a = (float)(1.f / tan(y_fov / 2.f)); + + m[0][0] = a / aspect; + m[0][1] = 0.f; + m[0][2] = 0.f; + m[0][3] = 0.f; + + m[1][0] = 0.f; + m[1][1] = a; + m[1][2] = 0.f; + m[1][3] = 0.f; + + m[2][0] = 0.f; + m[2][1] = 0.f; + m[2][2] = -((f + n) / (f - n)); + m[2][3] = -1.f; + + m[3][0] = 0.f; + m[3][1] = 0.f; + m[3][2] = -((2.f * f * n) / (f - n)); + m[3][3] = 0.f; +} +static inline void mat4x4_look_at(mat4x4 m, vec3 eye, vec3 center, vec3 up) { + /* Adapted from Android's OpenGL Matrix.java. */ + /* See the OpenGL GLUT documentation for gluLookAt for a description */ + /* of the algorithm. We implement it in a straightforward way: */ + + /* TODO: The negation of of can be spared by swapping the order of + * operands in the following cross products in the right way. */ + vec3 f; + vec3_sub(f, center, eye); + vec3_norm(f, f); + + vec3 s; + vec3_mul_cross(s, f, up); + vec3_norm(s, s); + + vec3 t; + vec3_mul_cross(t, s, f); + + m[0][0] = s[0]; + m[0][1] = t[0]; + m[0][2] = -f[0]; + m[0][3] = 0.f; + + m[1][0] = s[1]; + m[1][1] = t[1]; + m[1][2] = -f[1]; + m[1][3] = 0.f; + + m[2][0] = s[2]; + m[2][1] = t[2]; + m[2][2] = -f[2]; + m[2][3] = 0.f; + + m[3][0] = 0.f; + m[3][1] = 0.f; + m[3][2] = 0.f; + m[3][3] = 1.f; + + mat4x4_translate_in_place(m, -eye[0], -eye[1], -eye[2]); +} + +typedef float quat[4]; +static inline void quat_identity(quat q) { + q[0] = q[1] = q[2] = 0.f; + q[3] = 1.f; +} +static inline void quat_add(quat r, quat a, quat b) { + int i; + for (i = 0; i < 4; ++i) r[i] = a[i] + b[i]; +} +static inline void quat_sub(quat r, quat a, quat b) { + int i; + for (i = 0; i < 4; ++i) r[i] = a[i] - b[i]; +} +static inline void quat_mul(quat r, quat p, quat q) { + vec3 w; + vec3_mul_cross(r, p, q); + vec3_scale(w, p, q[3]); + vec3_add(r, r, w); + vec3_scale(w, q, p[3]); + vec3_add(r, r, w); + r[3] = p[3] * q[3] - vec3_mul_inner(p, q); +} +static inline void quat_scale(quat r, quat v, float s) { + int i; + for (i = 0; i < 4; ++i) r[i] = v[i] * s; +} +static inline float quat_inner_product(quat a, quat b) { + float p = 0.f; + int i; + for (i = 0; i < 4; ++i) p += b[i] * a[i]; + return p; +} +static inline void quat_conj(quat r, quat q) { + int i; + for (i = 0; i < 3; ++i) r[i] = -q[i]; + r[3] = q[3]; +} +#define quat_norm vec4_norm +static inline void quat_mul_vec3(vec3 r, quat q, vec3 v) { + quat v_ = {v[0], v[1], v[2], 0.f}; + + quat_conj(r, q); + quat_norm(r, r); + quat_mul(r, v_, r); + quat_mul(r, q, r); +} +static inline void mat4x4_from_quat(mat4x4 M, quat q) { + float a = q[3]; + float b = q[0]; + float c = q[1]; + float d = q[2]; + float a2 = a * a; + float b2 = b * b; + float c2 = c * c; + float d2 = d * d; + + M[0][0] = a2 + b2 - c2 - d2; + M[0][1] = 2.f * (b * c + a * d); + M[0][2] = 2.f * (b * d - a * c); + M[0][3] = 0.f; + + M[1][0] = 2 * (b * c - a * d); + M[1][1] = a2 - b2 + c2 - d2; + M[1][2] = 2.f * (c * d + a * b); + M[1][3] = 0.f; + + M[2][0] = 2.f * (b * d + a * c); + M[2][1] = 2.f * (c * d - a * b); + M[2][2] = a2 - b2 - c2 + d2; + M[2][3] = 0.f; + + M[3][0] = M[3][1] = M[3][2] = 0.f; + M[3][3] = 1.f; +} + +static inline void mat4x4o_mul_quat(mat4x4 R, mat4x4 M, quat q) { + /* XXX: The way this is written only works for othogonal matrices. */ + /* TODO: Take care of non-orthogonal case. */ + quat_mul_vec3(R[0], q, M[0]); + quat_mul_vec3(R[1], q, M[1]); + quat_mul_vec3(R[2], q, M[2]); + + R[3][0] = R[3][1] = R[3][2] = 0.f; + R[3][3] = 1.f; +} +static inline void quat_from_mat4x4(quat q, mat4x4 M) { + float r = 0.f; + int i; + + int perm[] = {0, 1, 2, 0, 1}; + int *p = perm; + + for (i = 0; i < 3; i++) { + float m = M[i][i]; + if (m < r) continue; + m = r; + p = &perm[i]; + } + + r = sqrtf(1.f + M[p[0]][p[0]] - M[p[1]][p[1]] - M[p[2]][p[2]]); + + if (r < 1e-6) { + q[0] = 1.f; + q[1] = q[2] = q[3] = 0.f; + return; + } + + q[0] = r / 2.f; + q[1] = (M[p[0]][p[1]] - M[p[1]][p[0]]) / (2.f * r); + q[2] = (M[p[2]][p[0]] - M[p[0]][p[2]]) / (2.f * r); + q[3] = (M[p[2]][p[1]] - M[p[1]][p[2]]) / (2.f * r); +} + +#endif diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/main.cpp b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/main.cpp new file mode 100644 index 0000000000..0788243e52 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/main.cpp @@ -0,0 +1,454 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +/* + * Modifications Copyright (C)2021 Advanced + * Micro Devices, Inc. All rights reserved. + */ + +#include "VulkanBaseApp.h" + +#include +#include +#include +#include +#include "linmath.h" +#include "hip_runtime.h" +#include "SineWaveSimulation.h" + + +typedef float vec2[2]; +std::string execution_path; + +#ifndef NDEBUG +#define ENABLE_VALIDATION (false) +#else +#define ENABLE_VALIDATION (true) +#endif + +class VulkanCudaSineWave : public VulkanBaseApp +{ + + typedef struct UniformBufferObject_st { + mat4x4 modelViewProj; + } UniformBufferObject; + + VkBuffer m_heightBuffer, m_xyBuffer, m_indexBuffer; + VkDeviceMemory m_heightMemory, m_xyMemory, m_indexMemory; + UniformBufferObject m_ubo; + VkSemaphore m_vkWaitSemaphore, m_vkSignalSemaphore; + SineWaveSimulation m_sim; + hipStream_t m_stream; + hipExternalSemaphore_t m_cudaWaitSemaphore, m_cudaSignalSemaphore; + hipExternalMemory_t m_cudaVertMem; + float *m_cudaHeightMap; + using chrono_tp = std::chrono::time_point; + chrono_tp m_lastTime; + size_t m_lastFrame; +public: + VulkanCudaSineWave(size_t width, size_t height) : + VulkanBaseApp("vulkanCudaSineWave", ENABLE_VALIDATION), + m_heightBuffer(VK_NULL_HANDLE), + m_xyBuffer(VK_NULL_HANDLE), + m_indexBuffer(VK_NULL_HANDLE), + m_heightMemory(VK_NULL_HANDLE), + m_xyMemory(VK_NULL_HANDLE), + m_indexMemory(VK_NULL_HANDLE), + m_ubo(), + m_sim(width, height), + m_stream(0), + m_vkWaitSemaphore(VK_NULL_HANDLE), + m_vkSignalSemaphore(VK_NULL_HANDLE), + m_cudaWaitSemaphore(), + m_cudaSignalSemaphore(), + m_cudaVertMem(), + m_cudaHeightMap(nullptr), + m_lastFrame(0) { + // Our index buffer can only index 32-bits of the vertex buffer + if ((width * height) > (1ULL << 32ULL)) { + throw std::runtime_error("Requested height and width is too large for this sample!"); + } + // Add our compiled vulkan shader files + TCHAR buffer[MAX_PATH] = { 0 }; + GetModuleFileName(NULL, buffer, MAX_PATH); + std::string str3 = std::string(buffer); + std::string str1 = "vert.spv" ; //sdkFindFilePath("sinewave.vert", execution_path.c_str()); + char* vertex_shader_path = strdup(str1.c_str()); + + std::string str2 = "frag.spv" ; //sdkFindFilePath("sinewave.frag", execution_path.c_str()); + char* fragment_shader_path = strdup(str2.c_str()); + + m_shaderFiles.push_back(std::make_pair(VK_SHADER_STAGE_VERTEX_BIT, vertex_shader_path)); + m_shaderFiles.push_back(std::make_pair(VK_SHADER_STAGE_FRAGMENT_BIT, fragment_shader_path)); + + } + ~VulkanCudaSineWave() { + // Make sure there's no pending work before we start tearing down + checkHIPErrors(hipStreamSynchronize(m_stream)); + + if (m_vkSignalSemaphore != VK_NULL_HANDLE) { + checkHIPErrors(hipDestroyExternalSemaphore(m_cudaSignalSemaphore)); + vkDestroySemaphore(m_device, m_vkSignalSemaphore, nullptr); + } + if (m_vkWaitSemaphore != VK_NULL_HANDLE) { + checkHIPErrors(hipDestroyExternalSemaphore(m_cudaWaitSemaphore)); + vkDestroySemaphore(m_device, m_vkWaitSemaphore, nullptr); + } + + if (m_xyBuffer != VK_NULL_HANDLE) { + vkDestroyBuffer(m_device, m_xyBuffer, nullptr); + } + if (m_xyMemory != VK_NULL_HANDLE) { + vkFreeMemory(m_device, m_xyMemory, nullptr); + } + + if (m_heightBuffer != VK_NULL_HANDLE) { + vkDestroyBuffer(m_device, m_heightBuffer, nullptr); + } + if (m_heightMemory != VK_NULL_HANDLE) { + vkFreeMemory(m_device, m_heightMemory, nullptr); + } + if (m_cudaHeightMap) { + checkHIPErrors(hipDestroyExternalMemory(m_cudaVertMem)); + } + + if (m_indexBuffer != VK_NULL_HANDLE) { + vkDestroyBuffer(m_device, m_indexBuffer, nullptr); + } + if (m_indexMemory != VK_NULL_HANDLE) { + vkFreeMemory(m_device, m_indexMemory, nullptr); + } + } + + void fillRenderingCommandBuffer(VkCommandBuffer& commandBuffer) { + VkBuffer vertexBuffers[] = { m_heightBuffer, m_xyBuffer }; + VkDeviceSize offsets[] = { 0, 0 }; + vkCmdBindVertexBuffers(commandBuffer, 0, sizeof(vertexBuffers) / sizeof(vertexBuffers[0]), vertexBuffers, offsets); + vkCmdBindIndexBuffer(commandBuffer, m_indexBuffer, 0, VK_INDEX_TYPE_UINT32); + vkCmdDrawIndexed(commandBuffer, (uint32_t)((m_sim.getWidth() - 1) * (m_sim.getHeight() - 1) * 6), 1, 0, 0, 0); + } + + void getVertexDescriptions(std::vector& bindingDesc, std::vector& attribDesc) { + bindingDesc.resize(2); + attribDesc.resize(2); + + bindingDesc[0].binding = 0; + bindingDesc[0].stride = sizeof(float); + bindingDesc[0].inputRate = VK_VERTEX_INPUT_RATE_VERTEX; + + bindingDesc[1].binding = 1; + bindingDesc[1].stride = sizeof(vec2); + bindingDesc[1].inputRate = VK_VERTEX_INPUT_RATE_VERTEX; + + attribDesc[0].binding = 0; + attribDesc[0].location = 0; + attribDesc[0].format = VK_FORMAT_R32_SFLOAT; + attribDesc[0].offset = 0; + + attribDesc[1].binding = 1; + attribDesc[1].location = 1; + attribDesc[1].format = VK_FORMAT_R32G32_SFLOAT; + attribDesc[1].offset = 0; + } + + void getAssemblyStateInfo(VkPipelineInputAssemblyStateCreateInfo& info) { + info.sType = VK_STRUCTURE_TYPE_PIPELINE_INPUT_ASSEMBLY_STATE_CREATE_INFO; + info.topology = VK_PRIMITIVE_TOPOLOGY_TRIANGLE_LIST; + info.primitiveRestartEnable = VK_FALSE; + } + + void getWaitFrameSemaphores(std::vector& wait, std::vector< VkPipelineStageFlags>& waitStages) const { + if (m_currentFrame != 0) { + // Have vulkan wait until cuda is done with the vertex buffer before rendering + // We don't do this on the first frame, as the wait semaphore hasn't been initialized yet + wait.push_back(m_vkWaitSemaphore); + // We want to wait until all the pipeline commands are complete before letting cuda work + waitStages.push_back(VK_PIPELINE_STAGE_ALL_COMMANDS_BIT); + } + } + + void getSignalFrameSemaphores(std::vector& signal) const { + // Add this semaphore for vulkan to signal once the vertex buffer is ready for cuda to modify + signal.push_back(m_vkSignalSemaphore); + } + + void initVulkanApp() { + int cuda_device = -1; + + // Select cuda device where vulkan is running. + cuda_device = m_sim.initCuda(m_vkDeviceUUID, VK_UUID_SIZE); + if (cuda_device == -1) + { + printf("Error: No CUDA-Vulkan interop capable device found\n"); + exit(EXIT_FAILURE); + } + + m_sim.initCudaLaunchConfig(cuda_device); + + // Create the cuda stream we'll be using + checkHIPErrors(hipStreamCreateWithFlags(&m_stream, hipStreamNonBlocking)); + + const size_t nVerts = m_sim.getWidth() * m_sim.getHeight(); + const size_t nInds = (m_sim.getWidth() - 1) * (m_sim.getHeight() - 1) * 6; + + // Create the height map cuda will write to + createExternalBuffer(nVerts * sizeof(float), + VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT, + VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, + getDefaultMemHandleType(), + m_heightBuffer, m_heightMemory); + + // Create the vertex buffer that will hold the xy coordinates for the grid + createBuffer(nVerts * sizeof(vec2), + VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_VERTEX_BUFFER_BIT, + VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, + m_xyBuffer, m_xyMemory); + + // Create the index buffer that references from both buffers above + createBuffer(nInds * sizeof(uint32_t), + VK_BUFFER_USAGE_TRANSFER_DST_BIT | VK_BUFFER_USAGE_INDEX_BUFFER_BIT, + VK_MEMORY_PROPERTY_DEVICE_LOCAL_BIT, + m_indexBuffer, m_indexMemory); + + // Import the height map into cuda and retrieve a device pointer to use + importHipExternalMemory((void **)&m_cudaHeightMap, m_cudaVertMem, m_heightMemory, nVerts * sizeof(*m_cudaHeightMap), getDefaultMemHandleType()); + // Set the height map to use in the simulation + m_sim.initSimulation(m_cudaHeightMap); + + { + // Set up the initial values for the vertex buffers with Vulkan + void *stagingBase; + VkBuffer stagingBuffer; + VkDeviceMemory stagingMemory; + VkDeviceSize stagingSz = std::max(nVerts * sizeof(vec2), nInds * sizeof(uint32_t)); + createBuffer(stagingSz, VK_BUFFER_USAGE_TRANSFER_SRC_BIT, + VK_MEMORY_PROPERTY_HOST_VISIBLE_BIT | VK_MEMORY_PROPERTY_HOST_COHERENT_BIT, stagingBuffer, stagingMemory); + + vkMapMemory(m_device, stagingMemory, 0, stagingSz, 0, &stagingBase); + + memset(stagingBase, 0, nVerts * sizeof(float)); + copyBuffer(m_heightBuffer, stagingBuffer, nVerts * sizeof(float)); + + for (size_t y = 0; y < m_sim.getHeight(); y++) { + for (size_t x = 0; x < m_sim.getWidth(); x++) { + vec2 *stagedVert = (vec2 *)stagingBase; + stagedVert[y * m_sim.getWidth() + x][0] = (2.0f * x) / (m_sim.getWidth() - 1) - 1; + stagedVert[y * m_sim.getWidth() + x][1] = (2.0f * y) / (m_sim.getHeight() - 1) - 1; + } + } + copyBuffer(m_xyBuffer, stagingBuffer, nVerts * sizeof(vec2)); + + { + uint32_t *indices = (uint32_t *)stagingBase; + for (size_t y = 0; y < m_sim.getHeight() - 1; y++) { + for (size_t x = 0; x < m_sim.getWidth() - 1; x++) { + indices[0] = (uint32_t)((y + 0) * m_sim.getWidth() + (x + 0)); + indices[1] = (uint32_t)((y + 1) * m_sim.getWidth() + (x + 0)); + indices[2] = (uint32_t)((y + 0) * m_sim.getWidth() + (x + 1)); + indices[3] = (uint32_t)((y + 1) * m_sim.getWidth() + (x + 0)); + indices[4] = (uint32_t)((y + 1) * m_sim.getWidth() + (x + 1)); + indices[5] = (uint32_t)((y + 0) * m_sim.getWidth() + (x + 1)); + indices += 6; + } + } + } + copyBuffer(m_indexBuffer, stagingBuffer, nInds * sizeof(uint32_t)); + + vkUnmapMemory(m_device, stagingMemory); + vkDestroyBuffer(m_device, stagingBuffer, nullptr); + vkFreeMemory(m_device, stagingMemory, nullptr); + } + + // Create the semaphore vulkan will signal when it's done with the vertex buffer + createExternalSemaphore(m_vkSignalSemaphore, getDefaultSemaphoreHandleType()); + // Create the semaphore vulkan will wait for before using the vertex buffer + createExternalSemaphore(m_vkWaitSemaphore, getDefaultSemaphoreHandleType()); + // Import the semaphore cuda will use -- vulkan's signal will be cuda's wait + importCudaExternalSemaphore(m_cudaWaitSemaphore, m_vkSignalSemaphore, getDefaultSemaphoreHandleType()); + // Import the semaphore cuda will use -- cuda's signal will be vulkan's wait + importCudaExternalSemaphore(m_cudaSignalSemaphore, m_vkWaitSemaphore, getDefaultSemaphoreHandleType()); + + } + + void importHipExternalMemory(void **cudaPtr, hipExternalMemory_t& cudaMem, VkDeviceMemory& vkMem, VkDeviceSize size, VkExternalMemoryHandleTypeFlagBits handleType) { + hipExternalMemoryHandleDesc externalMemoryHandleDesc = {}; + + if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) { + externalMemoryHandleDesc.type = hipExternalMemoryHandleTypeOpaqueWin32; + } + else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) { + externalMemoryHandleDesc.type = hipExternalMemoryHandleTypeOpaqueWin32Kmt; + } + else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) { + externalMemoryHandleDesc.type = hipExternalMemoryHandleTypeOpaqueFd; + } + else { + throw std::runtime_error("Unknown handle type requested!"); + } + + externalMemoryHandleDesc.size = size; + +#ifdef _WIN64 + externalMemoryHandleDesc.handle.win32.handle = (HANDLE)getMemHandle(vkMem, handleType); +#else + externalMemoryHandleDesc.handle.fd = (int)(uintptr_t)getMemHandle(vkMem, handleType); +#endif + checkHIPErrors(hipImportExternalMemory(&cudaMem, &externalMemoryHandleDesc)); + + hipExternalMemoryBufferDesc externalMemBufferDesc = {}; + externalMemBufferDesc.offset = 0; + externalMemBufferDesc.size = size; + externalMemBufferDesc.flags = 0; + + checkHIPErrors(hipExternalMemoryGetMappedBuffer(cudaPtr, cudaMem, &externalMemBufferDesc)); + } + + void importCudaExternalSemaphore(hipExternalSemaphore_t& cudaSem, VkSemaphore& vkSem, VkExternalSemaphoreHandleTypeFlagBits handleType) { + hipExternalSemaphoreHandleDesc externalSemaphoreHandleDesc = {}; + + if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_BIT) { + externalSemaphoreHandleDesc.type = hipExternalSemaphoreHandleTypeOpaqueWin32; + } + else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_WIN32_KMT_BIT) { + externalSemaphoreHandleDesc.type = hipExternalSemaphoreHandleTypeOpaqueWin32Kmt; + } + else if (handleType & VK_EXTERNAL_SEMAPHORE_HANDLE_TYPE_OPAQUE_FD_BIT) { + externalSemaphoreHandleDesc.type = hipExternalSemaphoreHandleTypeOpaqueFd; + } + else { + throw std::runtime_error("Unknown handle type requested!"); + } + +#ifdef _WIN64 + externalSemaphoreHandleDesc.handle.win32.handle = (HANDLE)getSemaphoreHandle(vkSem, handleType); +#else + externalSemaphoreHandleDesc.handle.fd = (int)(uintptr_t)getSemaphoreHandle(vkSem, handleType); +#endif + + externalSemaphoreHandleDesc.flags = 0; + + checkHIPErrors(hipImportExternalSemaphore(&cudaSem, &externalSemaphoreHandleDesc)); + } + + VkDeviceSize getUniformSize() const { + return sizeof(UniformBufferObject); + } + + void updateUniformBuffer(uint32_t imageIndex) { + { + mat4x4 view, proj; + vec3 eye = { 1.75f, 1.75f, 1.25f }; + vec3 center = { 0.0f, 0.0f, -0.25f }; + vec3 up = { 0.0f, 0.0f, 1.0f }; + + mat4x4_perspective(proj, (float)degreesToRadians(45.0f), m_swapChainExtent.width / (float)m_swapChainExtent.height, 0.1f, 10.0f); + proj[1][1] *= -1.0f; // Flip y axis + + mat4x4_look_at(view, eye, center, up); + mat4x4_mul(m_ubo.modelViewProj, proj, view); + } + + void *data; + vkMapMemory(m_device, m_uniformMemory[imageIndex], 0, getUniformSize(), 0, &data); + memcpy(data, &m_ubo, sizeof(m_ubo)); + vkUnmapMemory(m_device, m_uniformMemory[imageIndex]); + } + + std::vector getRequiredExtensions() const { + std::vector extensions; + extensions.push_back(VK_KHR_EXTERNAL_MEMORY_CAPABILITIES_EXTENSION_NAME); + extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_CAPABILITIES_EXTENSION_NAME); + return extensions; + } + + std::vector getRequiredDeviceExtensions() const { + std::vector extensions; + extensions.push_back(VK_KHR_EXTERNAL_MEMORY_EXTENSION_NAME); + extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_EXTENSION_NAME); +#ifdef _WIN64 + extensions.push_back(VK_KHR_EXTERNAL_MEMORY_WIN32_EXTENSION_NAME); + extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_WIN32_EXTENSION_NAME); +#else + extensions.push_back(VK_KHR_EXTERNAL_MEMORY_FD_EXTENSION_NAME); + extensions.push_back(VK_KHR_EXTERNAL_SEMAPHORE_FD_EXTENSION_NAME); +#endif /* _WIN64 */ + return extensions; + } + + void drawFrame() { + static chrono_tp startTime = std::chrono::high_resolution_clock::now(); + + chrono_tp currentTime = std::chrono::high_resolution_clock::now(); + float time = std::chrono::duration(currentTime - startTime).count(); + + if (m_currentFrame == 0) { + m_lastTime = startTime; + } + + float frame_time = std::chrono::duration(currentTime - m_lastTime).count(); + + hipExternalSemaphoreWaitParams waitParams = {}; + waitParams.flags = 0; + waitParams.params.fence.value = 1; + + hipExternalSemaphoreSignalParams signalParams = {}; + signalParams.flags = 0; + signalParams.params.fence.value = 0; + + + // Have vulkan draw the current frame... + VulkanBaseApp::drawFrame(); + // Wait for vulkan to complete it's work + checkHIPErrors(hipWaitExternalSemaphoresAsync(&m_cudaWaitSemaphore, &waitParams, 1, m_stream)); + // Now step the simulation + m_sim.stepSimulation(time, m_stream); + + // Signal vulkan to continue with the updated buffers + checkHIPErrors(hipSignalExternalSemaphoresAsync(&m_cudaSignalSemaphore, &signalParams, 1, m_stream)); + + // Output a naive measurement of the frames per second every five seconds + if (frame_time > 5) { + std::cout << "Average FPS (over " + << std::fixed << std::setprecision(2) << frame_time + << " seconds): " + << std::fixed << std::setprecision(2) + << ((m_currentFrame - m_lastFrame) / frame_time) + << std::endl; + m_lastFrame = m_currentFrame; + m_lastTime = currentTime; + } + } +}; + +int main(int argc, char **argv) +{ + execution_path = argv[0]; + VulkanCudaSineWave app((1ULL << 8ULL), (1ULL << 8ULL)); + app.init(); + app.mainLoop(); + return 0; +} diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.frag b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.frag new file mode 100644 index 0000000000..c850c7a248 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.frag @@ -0,0 +1,38 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + + +#version 450 +#extension GL_ARB_separate_shader_objects : enable + +layout(location = 0) in vec3 fragColor; + +layout(location = 0) out vec4 outColor; + +void main() { + outColor = vec4(fragColor, 1.0); +} \ No newline at end of file diff --git a/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.vert b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.vert new file mode 100644 index 0000000000..9157430756 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/20_hip_vulkan/sinewave.vert @@ -0,0 +1,43 @@ +/* Copyright (c) 2020, NVIDIA CORPORATION. All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions + * are met: + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * * Neither the name of NVIDIA CORPORATION nor the names of its + * contributors may be used to endorse or promote products derived + * from this software without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY + * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE + * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY + * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT + * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE + * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +#version 450 +#extension GL_ARB_separate_shader_objects : enable + +layout(binding = 0) uniform UniformBufferObject { + mat4 modelViewProj; +} ubo; + +layout(location = 0) in float height; +layout(location = 1) in vec2 xyPos; + +layout(location = 0) out vec3 fragColor; + +void main() { + gl_Position = ubo.modelViewProj * vec4(xyPos.xy, height, 1.0f); + fragColor = vec3(0.0f, (height + 0.5f), 0.0f); +}