EXSWHTEC-108 - Implement tests for hipFuncGetAttributes #58

Change-Id: I7aa979434724aa3d39343a6c69f67de83e61de4c

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",

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",

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;
+}

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 <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+#include <utils.hh>
+
+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<void*>(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<void*>(attribute_test_kernel)),
+                    hipErrorInvalidValue);
+  }
+  SECTION("func == nullptr") {
+    hipFuncAttributes attr;
+    HIP_CHECK_ERROR(hipFuncGetAttributes(&attr, nullptr), hipErrorInvalidDeviceFunction);
+  }
+}