diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index ddb82cfb77..5964a7aad9 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -262,6 +262,12 @@ "Note: Following two tests disabled due to defect - EXSWHTEC-153", "Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String", "Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String", + "Note: Test disabled due to defect - EXSWHTEC-163", + "Unit_hipModuleGetGlobal_Negative_Hmod_Is_Nullptr", + "Note: Test disabled due to defect - EXSWHTEC-164", + "Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String", + "Note: Test disabled due to defect - EXSWHTEC-165", + "Unit_hipModuleGetGlobal_Negative_Dptr_And_Bytes_Are_Nullptr", #endif #if defined VEGA20 "=== SWDEV-419112 Below tests fail in stress test on 29/08/23 ===", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 25436e508f..c26e5cdd67 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -363,6 +363,12 @@ "Note: Following two tests disabled due to defect - EXSWHTEC-153", "Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String", "Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String", + "Note: Test disabled due to defect - EXSWHTEC-163", + "Unit_hipModuleGetGlobal_Negative_Hmod_Is_Nullptr", + "Note: Test disabled due to defect - EXSWHTEC-164", + "Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String", + "Note: Test disabled due to defect - EXSWHTEC-165", + "Unit_hipModuleGetGlobal_Negative_Dptr_And_Bytes_Are_Nullptr", #endif "End of json" ] diff --git a/projects/hip-tests/catch/unit/module/CMakeLists.txt b/projects/hip-tests/catch/unit/module/CMakeLists.txt index 27f368f5ed..3d2611753c 100644 --- a/projects/hip-tests/catch/unit/module/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/module/CMakeLists.txt @@ -27,6 +27,7 @@ set(TEST_SRC hipModuleUnload.cc hipModuleGetFunction.cc hipModuleLaunchKernel.cc + hipModuleGetGlobal.cc ) add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code @@ -39,6 +40,12 @@ add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc) add_custom_target(launch_kernel_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code) +add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code + COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_global_test_module.cc -o get_global_test_module.code + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_global_test_module.cc) +add_custom_target(get_global_test_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code) + + add_custom_target(empty_module.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc @@ -143,6 +150,7 @@ hip_add_exe_to_target(NAME ModuleTest add_dependencies(build_tests empty_module.code) add_dependencies(ModuleTest get_function_module) add_dependencies(ModuleTest launch_kernel_module) +add_dependencies(ModuleTest get_global_test_module) if(HIP_PLATFORM MATCHES "amd") add_dependencies(build_tests copyKernel.code copyKernel.s) diff --git a/projects/hip-tests/catch/unit/module/get_global_test_module.cc b/projects/hip-tests/catch/unit/module/get_global_test_module.cc new file mode 100644 index 0000000000..98e58c7c54 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/get_global_test_module.cc @@ -0,0 +1,42 @@ +/* +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 "hip/hip_runtime.h" +#include "hip/hip_runtime_api.h" + +#include "hipModuleGetGlobal.hh" + +#define HIP_MODULE_GET_GLOBAL_TEST_DEFINE_DEVICE_GLOBALS(type) \ + __device__ type type##_var = 0; \ + __device__ type type##_arr[kArraySize] = {}; \ + extern "C" { \ + __global__ void type##_var_address_validation_kernel(void* ptr, bool* out) { \ + *out = static_cast(&type##_var) == ptr; \ + } \ + __global__ void type##_arr_address_validation_kernel(void* ptr, bool* out) { \ + *out = static_cast(type##_arr) == ptr; \ + } \ + } + +HIP_MODULE_GET_GLOBAL_TEST_DEFINE_DEVICE_GLOBALS(int) +HIP_MODULE_GET_GLOBAL_TEST_DEFINE_DEVICE_GLOBALS(float) +HIP_MODULE_GET_GLOBAL_TEST_DEFINE_DEVICE_GLOBALS(char) +HIP_MODULE_GET_GLOBAL_TEST_DEFINE_DEVICE_GLOBALS(double) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc b/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc new file mode 100644 index 0000000000..32f46f0caf --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc @@ -0,0 +1,145 @@ +/* +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 +#include +#include +#include + +#include +#include +#include +#include + +#include "hip_module_common.hh" +#include "hipModuleGetGlobal.hh" + +template +static void HipModuleGetGlobalTest(hipModule_t module, const std::string global_name) { + constexpr auto size = N * sizeof(T); + + hipDeviceptr_t global; + size_t global_size = 0; + HIP_CHECK(hipModuleGetGlobal(&global, &global_size, module, global_name.c_str())); + REQUIRE(global != 0); + REQUIRE(size == global_size); + + hipFunction_t kernel = nullptr; + const auto kernel_name = global_name + "_address_validation_kernel"; + HIP_CHECK(hipModuleGetFunction(&kernel, module, kernel_name.c_str())); + LinearAllocGuard equal_addresses(LinearAllocs::hipMalloc, sizeof(bool)); + HIP_CHECK(hipMemset(equal_addresses.ptr(), false, sizeof(*equal_addresses.ptr()))); + bool* equal_addresses_ptr = equal_addresses.ptr(); + void* kernel_args[2] = {&global, &equal_addresses_ptr}; + HIP_CHECK(hipModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_args, nullptr)); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipDeviceSynchronize()); + bool ok; + HIP_CHECK(hipMemcpy(&ok, equal_addresses_ptr, sizeof(ok), hipMemcpyDeviceToHost)); + REQUIRE(ok); + + constexpr T expected_value = 42; + std::array fill_buffer; + std::fill_n(fill_buffer.begin(), N, expected_value); + HIP_CHECK(hipMemcpyHtoD(global, fill_buffer.data(), size)); + + + std::array read_buffer; + HIP_CHECK(hipMemcpyDtoH(read_buffer.data(), global, size)); + ArrayFindIfNot(read_buffer.data(), expected_value, read_buffer.size()); +} + +#define HIP_MODULE_GET_GLOBAL_S(expr) #expr +#define HIP_MODULE_GET_GLOBAL_TEST(type, module) \ + SECTION("array") { \ + HipModuleGetGlobalTest(module, HIP_MODULE_GET_GLOBAL_S(type##_arr)); \ + } \ + SECTION("scalar") { \ + HipModuleGetGlobalTest(module, HIP_MODULE_GET_GLOBAL_S(type##_var)); \ + } + +static inline hipModule_t GetModule() { + HIP_CHECK(hipFree(nullptr)); + const static auto mg = ModuleGuard::LoadModule("get_global_test_module.code"); + return mg.module(); +} + +TEST_CASE("Unit_hipModuleGetGlobal_Positive_Basic") { + hipModule_t module = GetModule(); + + SECTION("int") { HIP_MODULE_GET_GLOBAL_TEST(int, module); } + + SECTION("float") { HIP_MODULE_GET_GLOBAL_TEST(float, module); } + + SECTION("char") { HIP_MODULE_GET_GLOBAL_TEST(char, module); } + + SECTION("double") { HIP_MODULE_GET_GLOBAL_TEST(double, module); } +} + +TEST_CASE("Unit_hipModuleGetGlobal_Positive_Parameters") { + hipModule_t module = GetModule(); + hipDeviceptr_t global = 0; + size_t global_size = 0; + + SECTION("dptr == nullptr") { + HIP_CHECK(hipModuleGetGlobal(nullptr, &global_size, module, "int_var")); + } + + SECTION("bytes == nullptr") { + HIP_CHECK(hipModuleGetGlobal(&global, nullptr, module, "int_var")); + } +} + +TEST_CASE("Unit_hipModuleGetGlobal_Negative_Parameters") { + hipModule_t module = GetModule(); + hipDeviceptr_t global = 0; + size_t global_size = 0; + + SECTION("name == nullptr") { + HIP_CHECK_ERROR(hipModuleGetGlobal(&global, &global_size, module, nullptr), + hipErrorInvalidValue); + } + + SECTION("name == invalid name") { + HIP_CHECK_ERROR(hipModuleGetGlobal(&global, &global_size, module, "dummy"), hipErrorNotFound); + } +} + +TEST_CASE("Unit_hipModuleGetGlobal_Negative_Hmod_Is_Nullptr") { + hipDeviceptr_t global = 0; + size_t global_size = 0; + + HIP_CHECK_ERROR(hipModuleGetGlobal(&global, &global_size, nullptr, "int_var"), + hipErrorInvalidResourceHandle); +} + +TEST_CASE("Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String") { + hipModule_t module = GetModule(); + hipDeviceptr_t global = 0; + size_t global_size = 0; + + HIP_CHECK_ERROR(hipModuleGetGlobal(&global, &global_size, module, ""), hipErrorInvalidValue); +} + +TEST_CASE("Unit_hipModuleGetGlobal_Negative_Dptr_And_Bytes_Are_Nullptr") { + hipModule_t module = GetModule(); + HIP_CHECK_ERROR(hipModuleGetGlobal(nullptr, nullptr, module, "int_var"), hipErrorInvalidValue); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.hh b/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.hh new file mode 100644 index 0000000000..8bd773f032 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipModuleGetGlobal.hh @@ -0,0 +1,28 @@ +/* +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. +*/ + +#pragma once + +#include + +namespace { +constexpr size_t kArraySize = 5; +} // anonymous namespace \ No newline at end of file