diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index f479b39e15..2f89c545c7 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -48,6 +48,10 @@ "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout", "Unit_hipFuncSetAttribute_Positive_Parameters", "Unit_hipFuncSetAttribute_Negative_Parameters", + "NOTE: The following test is disabled due to defect - EXSWHTEC-241", + "Unit_hipFuncGetAttributes_Negative_Parameters", + "NOTE: The following test is disabled due to defect - EXSWHTEC-242", + "Unit_hipFuncGetAttributes_Positive_Basic", "NOTE: The following test is disabled due to defect - EXSWHTEC-243", "Unit_hipExtLaunchKernel_Negative_Parameters", "NOTE: The following test is disabled due to defect - EXSWHTEC-244", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 097f165983..df0f98d53c 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -213,10 +213,15 @@ "Unit_hipVectorTypes_test_on_device", "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", + "NOTE: The following test is disabled due to defect - EXSWHTEC-241", + "Unit_hipFuncGetAttributes_Negative_Parameters", + "NOTE: The following test is disabled due to defect - EXSWHTEC-242", + "Unit_hipFuncGetAttributes_Positive_Basic", "NOTE: The following test is disabled due to defect - EXSWHTEC-243", "Unit_hipExtLaunchKernel_Negative_Parameters", "NOTE: The following test is disabled due to defect - EXSWHTEC-244", "Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters", + "Unit_hipMemAddressFree_negative", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipGraphAddMemcpyNode_Negative_Parameters", "Unit_hipMemCreate_ChkWithKerLaunch", diff --git a/catch/include/utils.hh b/catch/include/utils.hh index 5efd6a5125..f025768c14 100644 --- a/catch/include/utils.hh +++ b/catch/include/utils.hh @@ -169,3 +169,9 @@ inline bool DeviceAttributesSupport(const int device, Attributes... attributes) }; return (... && DeviceAttributeSupport(device, attributes)); } + +inline int GetDeviceAttribute(int device, const hipDeviceAttribute_t attr) { + int value = 0; + HIP_CHECK(hipDeviceGetAttribute(&value, attr, device)); + return value; +} diff --git a/catch/unit/executionControl/hipFuncGetAttributes.cc b/catch/unit/executionControl/hipFuncGetAttributes.cc new file mode 100644 index 0000000000..e97f44300e --- /dev/null +++ b/catch/unit/executionControl/hipFuncGetAttributes.cc @@ -0,0 +1,73 @@ +/* +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 + +constexpr size_t kConstSizeBytes = 128; +__constant__ char const_data[kConstSizeBytes]; + +__global__ void attribute_test_kernel() {} + +TEST_CASE("Unit_hipFuncGetAttributes_Positive_Basic") { + hipFuncAttributes attr; + HIP_CHECK(hipFuncGetAttributes(&attr, reinterpret_cast(attribute_test_kernel))); + + SECTION("binaryVersion") { +#if HT_NVIDIA + const auto major = GetDeviceAttribute(0, hipDeviceAttributeComputeCapabilityMajor); + const auto minor = GetDeviceAttribute(0, hipDeviceAttributeComputeCapabilityMinor); + REQUIRE(attr.binaryVersion == major * 10 + minor); +#elif HT_AMD + REQUIRE(attr.binaryVersion > 0); +#endif + } + + SECTION("cacheModeCA") { REQUIRE((attr.cacheModeCA == 0 || attr.cacheModeCA == 1)); } + + SECTION("constSizeBytes") { REQUIRE(attr.constSizeBytes == kConstSizeBytes); } + + SECTION("maxThreadsPerBlock") { + REQUIRE(attr.maxThreadsPerBlock == GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock)); + } + + SECTION("numRegs") { REQUIRE(attr.numRegs >= 0); } + + SECTION("ptxVersion") { REQUIRE(attr.ptxVersion > 0); } + + SECTION("sharedSizeBytes") { + REQUIRE(attr.sharedSizeBytes <= + GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock)); + } +} + +TEST_CASE("Unit_hipFuncGetAttributes_Negative_Parameters") { + SECTION("attr == nullptr") { + HIP_CHECK_ERROR(hipFuncGetAttributes(nullptr, reinterpret_cast(attribute_test_kernel)), + hipErrorInvalidValue); + } + SECTION("func == nullptr") { + hipFuncAttributes attr; + HIP_CHECK_ERROR(hipFuncGetAttributes(&attr, nullptr), hipErrorInvalidDeviceFunction); + } +} \ No newline at end of file