diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 66107df9e1..441b3dc9ed 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -136,6 +136,8 @@ "=== Below 2 tests are disable due to defect EXSWHTEC-369 ===", "Unit_Device_ilogbf_Accuracy_Positive", "Unit_Device_ilogb_Accuracy_Positive", + "NOTE: The following test is disabled due to defect - EXSWHTEC-245", + "Unit_hipFuncGetAttribute_Negative_Parameters", "Unit_hipMemAddressFree_negative", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipMemAddressReserve_Negative", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 5631444f2b..dacd9280c5 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -229,6 +229,9 @@ "=== Below 2 tests are disable due to defect EXSWHTEC-369 ===", "Unit_Device_ilogbf_Accuracy_Positive", "Unit_Device_ilogb_Accuracy_Positive", + "NOTE: The following test is disabled due to defect - EXSWHTEC-245", + "Unit_hipFuncGetAttribute_Negative_Parameters", + "Unit_hipMemAddressFree_negative", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipGraphAddMemcpyNode_Negative_Parameters", "Unit_hipMemCreate_ChkWithKerLaunch", diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index 76ca9e9ec6..6ae01bfbc5 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -31,6 +31,7 @@ set(TEST_SRC hipModuleGetTexRef.cc hipModuleLaunchCooperativeKernel.cc hipModuleLaunchCooperativeKernelMultiDevice.cc + hipFuncGetAttribute.cc ) add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code @@ -63,6 +64,9 @@ add_custom_target(get_tex_ref_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/get # Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900 --offload-arch=gfx906" # having space at the start/end of OFFLOAD_ARCH_STR can cause build failures +add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code + COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc -o get_function_module.code + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/get_function_module.cc) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC diff --git a/catch/unit/module/hipFuncGetAttribute.cc b/catch/unit/module/hipFuncGetAttribute.cc new file mode 100644 index 0000000000..c55b5179d3 --- /dev/null +++ b/catch/unit/module/hipFuncGetAttribute.cc @@ -0,0 +1,96 @@ +/* +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_module_common.hh" + +#include +#include +#include + +static hipModule_t GetModule() { + HIP_CHECK(hipFree(nullptr)); + static const auto mg = ModuleGuard::LoadModule("get_function_module.code"); + return mg.module(); +} + +TEST_CASE("Unit_hipFuncGetAttribute_Positive_Basic") { + hipFunction_t kernel = GetKernel(GetModule(), "GlobalKernel"); + + int value; + + SECTION("binaryVersion") { + HIP_CHECK(hipFuncGetAttribute(&value, HIP_FUNC_ATTRIBUTE_BINARY_VERSION, kernel)); +#if HT_NVIDIA + const auto major = GetDeviceAttribute(hipDeviceAttributeComputeCapabilityMajor, 0); + const auto minor = GetDeviceAttribute(hipDeviceAttributeComputeCapabilityMinor, 0); + REQUIRE(value == major * 10 + minor); +#elif HT_AMD + REQUIRE(value > 0); +#endif + } + + SECTION("cacheModeCA") { + HIP_CHECK(hipFuncGetAttribute(&value, HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA, kernel)); + REQUIRE((value == 0 || value == 1)); + } + + SECTION("maxThreadsPerBlock") { + HIP_CHECK(hipFuncGetAttribute(&value, HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, kernel)); + REQUIRE(value == GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0)); + } + + SECTION("numRegs") { + HIP_CHECK(hipFuncGetAttribute(&value, HIP_FUNC_ATTRIBUTE_NUM_REGS, kernel)); + REQUIRE(value >= 0); + } + + SECTION("ptxVersion") { + HIP_CHECK(hipFuncGetAttribute(&value, HIP_FUNC_ATTRIBUTE_PTX_VERSION, kernel)); + REQUIRE(value > 0); + } + + SECTION("sharedSizeBytes") { + HIP_CHECK(hipFuncGetAttribute(&value, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, kernel)); + REQUIRE(value <= GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0)); + } +} + +TEST_CASE("Unit_hipFuncGetAttribute_Negative_Parameters") { + hipFunction_t kernel = GetKernel(GetModule(), "GlobalKernel"); + + int value; + + SECTION("value == nullptr") { + HIP_CHECK_ERROR(hipFuncGetAttribute(nullptr, HIP_FUNC_ATTRIBUTE_BINARY_VERSION, kernel), + hipErrorInvalidValue); + } + + SECTION("invalid attribute") { + HIP_CHECK_ERROR(hipFuncGetAttribute(&value, static_cast(-1), kernel), + hipErrorInvalidValue); + } + + SECTION("hfunc == nullptr") { + HIP_CHECK_ERROR(hipFuncGetAttribute(&value, HIP_FUNC_ATTRIBUTE_BINARY_VERSION, nullptr), + hipErrorInvalidResourceHandle); + } +} \ No newline at end of file diff --git a/catch/unit/module/hipModuleLaunchCooperativeKernel.cc b/catch/unit/module/hipModuleLaunchCooperativeKernel.cc index 0ca6a31293..cf92152bce 100644 --- a/catch/unit/module/hipModuleLaunchCooperativeKernel.cc +++ b/catch/unit/module/hipModuleLaunchCooperativeKernel.cc @@ -97,17 +97,17 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Positive_Parameters") { hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); SECTION("blockDim.x == maxBlockDimX") { - const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX); + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0); HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, x, 1, 1, 0, nullptr, nullptr)); } SECTION("blockDim.y == maxBlockDimY") { - const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY); + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0); HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, y, 1, 1, 0, nullptr, nullptr)); } SECTION("blockDim.z == maxBlockDimZ") { - const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ); + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0); HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, z, 1, 1, 0, nullptr, nullptr)); } } @@ -168,25 +168,25 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Negative_Parameters") { } SECTION("blockDim.x > maxBlockDimX") { - const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u; + const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0) + 1u; HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, x, 1, 1, 0, nullptr, nullptr), hipErrorInvalidValue); } SECTION("blockDim.y > maxBlockDimY") { - const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u; + const unsigned int y = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimY, 0) + 1u; HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, y, 1, 0, nullptr, nullptr), hipErrorInvalidValue); } SECTION("blockDim.z > maxBlockDimZ") { - const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u; + const unsigned int z = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimZ, 0) + 1u; HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, z, 0, nullptr, nullptr), hipErrorInvalidValue); } SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { - const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock); + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxThreadsPerBlock, 0); const unsigned int dim = std::ceil(std::cbrt(max)); HIP_CHECK_ERROR( hipModuleLaunchCooperativeKernel(f, 1, 1, 1, dim, dim, dim, 0, nullptr, nullptr), @@ -195,7 +195,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Negative_Parameters") { #if HT_AMD // Disabled due to defect EXSWHTEC-351 SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { - const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u; + const unsigned int max = GetDeviceAttribute(hipDeviceAttributeMaxSharedMemoryPerBlock, 0) + 1u; HIP_CHECK_ERROR(hipModuleLaunchCooperativeKernel(f, 1, 1, 1, 1, 1, 1, max, nullptr, nullptr), hipErrorInvalidValue); }