diff --git a/catch/unit/vulkan_interop/CMakeLists.txt b/catch/unit/vulkan_interop/CMakeLists.txt index a0af39c7b7..f3b5073743 100644 --- a/catch/unit/vulkan_interop/CMakeLists.txt +++ b/catch/unit/vulkan_interop/CMakeLists.txt @@ -2,6 +2,7 @@ set(COMMON_SHARED_SRC vulkan_test.cc) set(TEST_SRC hipExternalMemoryGetMappedBuffer.cc + hipExternalMemoryGetMappedMipmappedArray.cc hipImportExternalMemory.cc hipDestroyExternalMemory.cc hipWaitExternalSemaphoresAsync.cc diff --git a/catch/unit/vulkan_interop/hipExternalMemoryGetMappedMipmappedArray.cc b/catch/unit/vulkan_interop/hipExternalMemoryGetMappedMipmappedArray.cc new file mode 100644 index 0000000000..b954968206 --- /dev/null +++ b/catch/unit/vulkan_interop/hipExternalMemoryGetMappedMipmappedArray.cc @@ -0,0 +1,214 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "vulkan_test.hh" + +constexpr bool enable_validation = false; + +template bool WriteAndValidateData(hipArray_t& array, size_t array_size) { + T* write_data = new T[array_size]; + for (size_t i = 0; i < array_size; i++) { + write_data[i] = rand() % 10; + } + HIP_CHECK( + hipMemcpyToArray(array, 0, 0, write_data, array_size * sizeof(T), hipMemcpyHostToDevice)); + + T* read_data = new T[array_size]; + HIP_CHECK( + hipMemcpyFromArray(read_data, array, 0, 0, array_size * sizeof(T), hipMemcpyDeviceToHost)); + + bool is_valid = true; + for (size_t i = 0; i < array_size && is_valid; i++) { + if (write_data[i] != read_data[i]) { + is_valid = false; + } + } + + free(read_data); + free(write_data); + + return is_valid; +} + +TEST_CASE("Unit_hipExternalMemoryGetMappedMipmappedArray_Vulkan_Positive_Read_Write") { + VulkanTest vkt(enable_validation); + using type = uint8_t; + constexpr uint32_t count = 16384; + + const auto vk_storage = + vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + if (vk_storage.memory == nullptr) { + return; + } + + const auto ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size); + hipExternalMemory_t ext_memory; + HIP_CHECK(hipImportExternalMemory(&ext_memory, &ext_mem_desc)); + + hipExternalMemoryMipmappedArrayDesc mipmapped_arr_desc = {}; + mipmapped_arr_desc.extent = {}; + mipmapped_arr_desc.extent.width = GENERATE(32, 128, 256); + mipmapped_arr_desc.extent.height = GENERATE(1, 4, 16); + mipmapped_arr_desc.extent.depth = GENERATE(0, 4); + mipmapped_arr_desc.flags = hipArrayDefault; + mipmapped_arr_desc.formatDesc = hipCreateChannelDesc(); + mipmapped_arr_desc.numLevels = GENERATE(1, 2, 4); + mipmapped_arr_desc.offset = 0; + + hipMipmappedArray_t mipmapped_arr = nullptr; + HIP_CHECK( + hipExternalMemoryGetMappedMipmappedArray(&mipmapped_arr, ext_memory, &mipmapped_arr_desc)); + + hipArray_t level_arr = nullptr; + HIP_CHECK(hipGetMipmappedArrayLevel(&level_arr, mipmapped_arr, 1)); + + size_t level_arr_size = mipmapped_arr_desc.extent.width * mipmapped_arr_desc.extent.height * + mipmapped_arr_desc.extent.depth; + + REQUIRE(WriteAndValidateData(level_arr, level_arr_size) == true); + + HIP_CHECK(hipFreeArray(level_arr)); + HIP_CHECK(hipFreeMipmappedArray(mipmapped_arr)); + HIP_CHECK(hipDestroyExternalMemory(ext_memory)); +} + +TEST_CASE("Unit_hipExternalMemoryGetMappedMipmappedArray_Vulkan_Array_Layered") { + VulkanTest vkt(enable_validation); + using type = uint8_t; + constexpr uint32_t count = 16384; + + const auto vk_storage = + vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + if (vk_storage.memory == nullptr) { + return; + } + + const auto ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size); + hipExternalMemory_t ext_memory; + HIP_CHECK(hipImportExternalMemory(&ext_memory, &ext_mem_desc)); + + hipExternalMemoryMipmappedArrayDesc mipmapped_arr_desc = {}; + mipmapped_arr_desc.extent = {}; + mipmapped_arr_desc.extent.width = GENERATE(32, 128, 256); + mipmapped_arr_desc.extent.height = GENERATE(0, 4); + mipmapped_arr_desc.extent.depth = 16; + mipmapped_arr_desc.flags = hipArrayLayered; + mipmapped_arr_desc.formatDesc = hipCreateChannelDesc(); + mipmapped_arr_desc.numLevels = GENERATE(1, 2, 4); + mipmapped_arr_desc.offset = 0; + + hipMipmappedArray_t mipmapped_arr = nullptr; + HIP_CHECK( + hipExternalMemoryGetMappedMipmappedArray(&mipmapped_arr, ext_memory, &mipmapped_arr_desc)); + + HIP_CHECK(hipFreeMipmappedArray(mipmapped_arr)); + HIP_CHECK(hipDestroyExternalMemory(ext_memory)); +} + +TEST_CASE("Unit_hipExternalMemoryGetMappedMipmappedArray_Vulkan_Array_Cubemap") { + VulkanTest vkt(enable_validation); + using type = uint8_t; + // cubemap HIP array is allocated if all three extents are non-zero and the hipArrayCubemap + // flag is set. Width must be equal to height, and depth must be six + constexpr uint32_t cube_size = 32; + constexpr uint32_t depth = 6; + constexpr uint32_t ext_mem_size = cube_size * cube_size * depth; + + const auto vk_storage = + vkt.CreateMappedStorage(ext_mem_size, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + if (vk_storage.memory == nullptr) { + return; + } + + const auto ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size); + hipExternalMemory_t ext_memory; + HIP_CHECK(hipImportExternalMemory(&ext_memory, &ext_mem_desc)); + + hipExternalMemoryMipmappedArrayDesc mipmapped_arr_desc = {}; + mipmapped_arr_desc.extent = {}; + mipmapped_arr_desc.extent.width = cube_size; + mipmapped_arr_desc.extent.height = cube_size; + mipmapped_arr_desc.extent.depth = depth; + mipmapped_arr_desc.flags = hipArrayCubemap; + mipmapped_arr_desc.formatDesc = hipCreateChannelDesc(); + mipmapped_arr_desc.numLevels = GENERATE(1, 2, 4); + mipmapped_arr_desc.offset = 0; + + hipMipmappedArray_t mipmapped_arr = nullptr; + HIP_CHECK( + hipExternalMemoryGetMappedMipmappedArray(&mipmapped_arr, ext_memory, &mipmapped_arr_desc)); + + HIP_CHECK(hipFreeMipmappedArray(mipmapped_arr)); + HIP_CHECK(hipDestroyExternalMemory(ext_memory)); +} + +TEST_CASE("Unit_hipExternalMemoryGetMappedMipmappedArray_Vulkan_Negative_Parameters") { + CHECK_IMAGE_SUPPORT + + VulkanTest vkt(enable_validation); + using type = uint8_t; + constexpr uint32_t count = 256; + + const auto vk_storage = + vkt.CreateMappedStorage(count, VK_BUFFER_USAGE_TRANSFER_DST_BIT, true); + if (vk_storage.memory == nullptr) { + return; + } + + const auto ext_mem_desc = vkt.BuildMemoryDescriptor(vk_storage.memory, vk_storage.size); + hipExternalMemory_t ext_memory; + HIP_CHECK(hipImportExternalMemory(&ext_memory, &ext_mem_desc)); + + hipExternalMemoryMipmappedArrayDesc mipmapped_arr_desc = {}; + mipmapped_arr_desc.extent = {.width = count, .height = 0, .depth = 0}; + mipmapped_arr_desc.flags = hipArrayDefault; + mipmapped_arr_desc.formatDesc = hipCreateChannelDesc(); + mipmapped_arr_desc.numLevels = 2; + mipmapped_arr_desc.offset = 0; + + hipMipmappedArray_t mipmapped_arr = nullptr; + + SECTION("Nullptr_Array") { + HIP_CHECK_ERROR( + hipExternalMemoryGetMappedMipmappedArray(nullptr, ext_memory, &mipmapped_arr_desc), + hipErrorInvalidValue); + } + + SECTION("Nullptr_ExternalMemory") { + HIP_CHECK_ERROR( + hipExternalMemoryGetMappedMipmappedArray(&mipmapped_arr, nullptr, &mipmapped_arr_desc), + hipErrorInvalidValue); + } + + SECTION("Nullptr_ArrayDescription") { + HIP_CHECK_ERROR(hipExternalMemoryGetMappedMipmappedArray(&mipmapped_arr, ext_memory, nullptr), + hipErrorInvalidValue); + } + + SECTION("Invalid_Offset") { + mipmapped_arr_desc.offset = 1 * sizeof(type); + HIP_CHECK_ERROR( + hipExternalMemoryGetMappedMipmappedArray(&mipmapped_arr, ext_memory, &mipmapped_arr_desc), + hipErrorInvalidValue); + } + + HIP_CHECK(hipDestroyExternalMemory(ext_memory)); +}