EXSWHTEC-100 - Miscellaneous modifications to existing memory tests (#96)

- Miscellaneous modifications to existing memory tests - Update config_amd_linux_common.json - Update config_amd_windows_common.json - Disable tests which failed on external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/96 [ROCm/hip-tests commit: 2f64855322]
2023-07-08 17:21:34 +02:00
parent 1bb4348c3b
commit d7f8a95ce1
8 changed files with 67 additions and 59 deletions
@@ -98,6 +98,8 @@
 	   "=== Below tests fail in stress test on 30/06/23 ===",
 	   "Unit_hipStreamValue_Wait32_Blocking_NoMask_Nor",
 	   "Unit_hipStreamValue_Write - TestParams<uint32_t, PtrType::HostPtr>",
-	   "Unit_hipMemcpyParam2DAsync_multiDevice-StreamOnDiffDevice"
+	   "Unit_hipMemcpyParam2DAsync_multiDevice-StreamOnDiffDevice",
+	   "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/96 ===",
+	   "Unit_hipHostGetDevicePointer_Negative"
 	]
 }
@@ -185,6 +185,8 @@
 	   "SWDEV-398981 fails in stress test",
 	   "Unit_hipStreamCreateWithPriority_MulthreadDefaultflag",
 	   "Note: UUID returned empty on some windows nodes",
-	   "Unit_hipDeviceGetUuid_Positive"
+	   "Unit_hipDeviceGetUuid_Positive",
+     "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/96 ===",
+	   "Unit_hipHostGetDevicePointer_Negative"
 	]
 }
@@ -49,11 +49,10 @@ using namespace std::chrono_literals;
 const std::chrono::duration<uint64_t, std::milli> delay = 50ms;
 constexpr size_t numAllocs = 10;

-#if HT_AMD /* Disabled because frequency based wait is timing out on nvidia platforms */
-TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float4) {
-  TestType* devPtr{};
+TEST_CASE("Unit_hipFreeImplicitSyncDev") {
+  int* devPtr{};
  size_t size_mult = GENERATE(1, 32, 64, 128, 256);
-  HIP_CHECK(hipMalloc(&devPtr, sizeof(TestType) * size_mult));
+  HIP_CHECK(hipMalloc(&devPtr, sizeof(*devPtr) * size_mult));

  HipTest::runKernelForDuration(delay);
  // make sure device is busy
@@ -62,11 +61,11 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncDev", "", char, float, float2, float
  HIP_CHECK(hipStreamQuery(nullptr));
 }

-TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, float4) {
-  TestType* hostPtr{};
+TEST_CASE("Unit_hipFreeImplicitSyncHost") {
+  int* hostPtr{};
  size_t size_mult = GENERATE(1, 32, 64, 128, 256);

-  HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(TestType) * size_mult));
+  HIP_CHECK(hipHostMalloc(&hostPtr, sizeof(*hostPtr) * size_mult));

  HipTest::runKernelForDuration(delay);
  // make sure device is busy
@@ -75,7 +74,7 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncHost", "", char, float, float2, floa
  HIP_CHECK(hipStreamQuery(nullptr));
 }

-#if HT_NVIDIA  // Meaningless at the moment, since we are not running wait kernel on nvidia.
+#if HT_NVIDIA
 TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, float4) {
  using vec_info = vector_info<TestType>;
  DriverContext ctx;
@@ -135,7 +134,6 @@ TEMPLATE_TEST_CASE("Unit_hipFreeImplicitSyncArray", "", char, float, float2, flo
  }
 }

-#endif
 #endif

 // Freeing a invalid pointer with on device
@@ -166,8 +164,6 @@ TEST_CASE("Unit_hipFreeNegativeHost") {
 #if HT_NVIDIA
 TEST_CASE("Unit_hipFreeNegativeArray") {
  DriverContext ctx;
-  hipArray_t arrayPtr{};
-  hiparray cuArrayPtr{};

  SECTION("ArrayFree") { HIP_CHECK(hipFreeArray(nullptr)); }
  SECTION("ArrayDestroy") {
@@ -21,11 +21,19 @@ THE SOFTWARE.
 */

 #include <hip_test_common.hh>
+#include <utils.hh>

 TEST_CASE("Unit_hipHostGetDevicePointer_Negative") {
  int* hPtr{nullptr};
+  int* dPtr{nullptr};
  HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int)));

+  if (!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) {
+    HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), hPtr, 0),
+                    hipErrorNotSupported);
+    return;
+  }
+
  SECTION("Nullptr as device") {
    HIP_CHECK_ERROR(hipHostGetDevicePointer(nullptr, hPtr, 0), hipErrorInvalidValue);
  }
@@ -36,13 +44,29 @@ TEST_CASE("Unit_hipHostGetDevicePointer_Negative") {
                    hipErrorInvalidValue);
  }

-  // Not adding check for flags since CUDA spec states that there might be more values added to it
+  SECTION("Non pinned memory as host") {
+    int* hPtr = reinterpret_cast<int*>(malloc(sizeof(*hPtr)));
+    HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), hPtr, 0),
+                    hipErrorInvalidValue);
+    free(hPtr);
+  }
+
+  SECTION("Flags non-zero") {
+    HIP_CHECK_ERROR(hipHostGetDevicePointer(reinterpret_cast<void**>(&dPtr), hPtr, 1),
+                    hipErrorInvalidValue);
+  }
+
  HIP_CHECK(hipHostFree(hPtr));
 }

 template <typename T> __global__ void set(T* ptr, T val) { *ptr = val; }

 TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") {
+  if(!DeviceAttributesSupport(0, hipDeviceAttributeCanMapHostMemory)) {
+    HipTest::HIP_SKIP_TEST("Device does not support mapping host memory"); 
+    return;
+  }
+
  int* hPtr{nullptr};
  HIP_CHECK(hipHostMalloc(&hPtr, sizeof(int)));

@@ -71,8 +95,8 @@ TEST_CASE("Unit_hipHostGetDevicePointer_UseCase") {
    HIP_CHECK(hipDeviceSynchronize());
    HIP_CHECK(hipHostUnregister(&res));

-    REQUIRE(value == 10);
+    REQUIRE(res == value);
  }

  HIP_CHECK(hipHostFree(hPtr));
-}
+}
@@ -27,9 +27,10 @@ This testfile verifies the following scenarios of hipHostRegister API
 2. hipHostRegister and perform hipMemcpy on it.
 */

+#include "hip/hip_runtime_api.h"
 #include <hip_test_common.hh>
 #include <hip_test_helper.hh>
-#include "hip/hip_runtime_api.h"
+#include <utils.hh>

 #define OFFSET 128
 static constexpr auto LEN{1024 * 1024};
@@ -63,9 +64,7 @@ void doMemCopy(size_t numElements, int offset, T* A, T* Bh, T* Bd, bool internal
  HIP_CHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost));

  // Make sure the copy worked
-  for (size_t i = 0; i < numElements; i++) {
-    REQUIRE(Bh[i] == A[i]);
-  }
+  ArrayMismatch(A, Bh, numElements);

  if (internalRegister) {
    HIP_CHECK(hipHostUnregister(A));
@@ -68,6 +68,12 @@ TEST_CASE("Unit_hipHostUnregister_NullPtr") {
  HIP_CHECK_ERROR(hipHostUnregister(nullptr), hipErrorInvalidValue);
 }

+TEST_CASE("Unit_hipHostUnregister_Ptr_Different_Than_Specified_To_Register") {
+  std::vector<int> alloc(2);
+  HIP_CHECK(hipHostRegister(alloc.data(), alloc.size(), 0));
+  HIP_CHECK_ERROR(hipHostUnregister(&alloc.data()[1]), hipErrorHostMemoryNotRegistered);
+}
+
 TEST_CASE("Unit_hipHostUnregister_NotRegisteredPointer") {
  auto x = std::unique_ptr<int>(new int);
  HIP_CHECK_ERROR(hipHostUnregister(x.get()), hipErrorHostMemoryNotRegistered);
@@ -228,6 +228,21 @@ TEST_CASE("Unit_hipMallocPitch_Negative") {
  }
 }

+TEST_CASE("Unit_hipMallocPitch_Zero_Dims") {
+  void* ptr = nullptr;
+  size_t pitch = 0;
+
+  SECTION("width == 0") {
+    HIP_CHECK(hipMallocPitch(&ptr, &pitch, 0, 1));
+    REQUIRE(ptr == nullptr);
+  }
+
+  SECTION("height == 0") {
+    HIP_CHECK(hipMallocPitch(&ptr, &pitch, 1, 0));
+    REQUIRE(ptr == nullptr);
+  }
+}
+
 TEST_CASE("Unit_hipMemAllocPitch_Negative") {
  size_t pitch = 0;
  hipDeviceptr_t ptr{};
@@ -360,42 +375,7 @@ static void MemoryAllocDiffSizes(int gpu) {
 static void threadFunc(int gpu) {
  MemoryAllocDiffSizes<float>(gpu);
 }
-/*
- * This testcase verifies the negative scenarios of hipMallocPitch API
- */
-#if 0 //TODO: Review, fix and re-enable test
-TEST_CASE("Unit_hipMallocPitch_Negative") {
-  float* A_d;
-  size_t pitch_A = 0;
-  size_t width{NUM_W * sizeof(float)};
-#if HT_NVIDIA
-  SECTION("NullPtr to Pitched Ptr") {
-    REQUIRE(hipMallocPitch(nullptr,
-            &pitch_A, width, NUM_H) != hipSuccess);
-  }

-  SECTION("nullptr to pitch") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           nullptr, width, NUM_H) != hipSuccess);
-  }
-#endif
-  SECTION("Width 0 in hipMallocPitch") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           &pitch_A, 0, NUM_H) == hipSuccess);
-  }
-
-  SECTION("Height 0 in hipMallocPitch") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           &pitch_A, width, 0) == hipSuccess);
-  }
-
-  SECTION("Max int values") {
-    REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
-                           &pitch_A, std::numeric_limits<int>::max(),
-                           std::numeric_limits<int>::max()) != hipSuccess);
-  }
-}
-#endif
 /*
 * This testcase verifies the basic scenario of
 * hipMallocPitch API for different datatypes
@@ -408,6 +388,7 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_Basic",
  size_t width{NUM_W * sizeof(TestType)};
  REQUIRE(hipMallocPitch(reinterpret_cast<void**>(&A_d),
          &pitch_A, width, NUM_H) == hipSuccess);
+  REQUIRE(width <= pitch_A);
  HIP_CHECK(hipFree(A_d));
 }

@@ -538,4 +519,3 @@ TEMPLATE_TEST_CASE("Unit_hipMallocPitch_KernelLaunch", ""
  HipTest::freeArrays<TestType>(nullptr, nullptr, nullptr,
                                A_h, B_h, C_h, false);
 }
-
@@ -316,9 +316,8 @@ TEST_CASE("Unit_hipPointerGetAttribute_Negative") {
        == hipErrorInvalidValue);
  }
  SECTION("Pass invalid attribute") {
-    hipPointer_attribute attr{HIP_POINTER_ATTRIBUTE_DEVICE_POINTER};
-    REQUIRE(hipPointerGetAttribute(&data, attr,
-          reinterpret_cast<hipDeviceptr_t>(A_h)) == hipErrorInvalidValue);
+    REQUIRE(hipPointerGetAttribute(&data, static_cast<hipPointer_attribute>(-1),
+                                   reinterpret_cast<hipDeviceptr_t>(A_h)) == hipErrorInvalidValue);
  }
 #if HT_AMD
  SECTION("Pass HIP_POINTER_ATTRIBUTE_IS_GPU_DIRECT_RDMA_CAPABLE"