From d8ea9afa666e9d3b6252a9d18f5e9a91b30ce174 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 6 Oct 2021 18:45:15 +0530
Subject: [PATCH] Update hipMallocManaged.cc

* Fix integer and size_t comparisions.
* Split basic scenario into two - all devices and HMM enabled devices only.
---
 catch/unit/memory/hipMallocManaged.cc | 25 ++++++++++++++++++++++---
 1 file changed, 22 insertions(+), 3 deletions(-)

diff --git a/catch/unit/memory/hipMallocManaged.cc b/catch/unit/memory/hipMallocManaged.cc
index 0b61d0d922..e8bc69997c 100644
--- a/catch/unit/memory/hipMallocManaged.cc
+++ b/catch/unit/memory/hipMallocManaged.cc
@@ -155,13 +155,28 @@ static unsigned blocksPerCU{6};
 static unsigned threadsPerBlock{256};
 
 /*
-   This testcase verifies the hipMallocManaged basic scenario
+   This testcase verifies the hipMallocManaged basic scenario - supported on all devices
  */
 
 TEST_CASE("Unit_hipMallocManaged_Basic") {
     int numElements = (N < (64 * 1024 * 1024)) ? 64 * 1024 * 1024 : N;
     float *A, *B, *C;
 
+    HIP_CHECK(hipMallocManaged(&A, numElements*sizeof(float)));
+    HIP_CHECK(hipMallocManaged(&B, numElements*sizeof(float)));
+    HIP_CHECK(hipMallocManaged(&C, numElements*sizeof(float)));
+}
+
+/*
+   This testcase verifies the hipMallocManaged basic scenario - supported only on HMM enabled devices
+ */
+
+TEST_CASE("Unit_hipMallocManaged_Advanced") {
+  int managed =  HmmAttrPrint();
+  if (managed == 1) {
+    int numElements = (N < (64 * 1024 * 1024)) ? 64 * 1024 * 1024 : N;
+    float *A, *B, *C;
+
     HIP_CHECK(hipMallocManaged(&A, numElements*sizeof(float)));
     HIP_CHECK(hipMallocManaged(&B, numElements*sizeof(float)));
     HIP_CHECK(hipMallocManaged(&C, numElements*sizeof(float)));
@@ -220,6 +235,10 @@ TEST_CASE("Unit_hipMallocManaged_Basic") {
     HIP_CHECK(hipFree(A));
     HIP_CHECK(hipFree(B));
     REQUIRE(maxError != 0.0f);
+  } else {
+    SUCCEED("GPU 0 doesn't support hipDeviceAttributeManagedMemory "
+           "attribute. Hence skipping the testing with Pass result.\n");
+  }
 }
 
 
@@ -237,7 +256,7 @@ TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
     HIP_CHECK(hipStreamCreate(&strm));
     HIP_CHECK(hipMallocManaged(&Hmm, sizeof(int) * NumElms));
     HIP_CHECK(hipMalloc(&Dptr, sizeof(int) * NumElms));
-    for (int i = 0; i < NumElms; ++i) {
+    for (size_t i = 0; i < NumElms; ++i) {
       Hmm[i] = 0;
       Hptr[i] = InitVal;
     }
@@ -248,7 +267,7 @@ TEST_CASE("Unit_hipMallocManaged_KrnlWth2MemTypes") {
     KrnlWth2MemTypes<<<dimGrid, dimBlock, 0, strm>>>(Hmm, Dptr, NumElms);
     HIP_CHECK(hipStreamSynchronize(strm));
     // Verifying the results
-    for (int k = 0; k < NumElms; ++k) {
+    for (size_t k = 0; k < NumElms; ++k) {
       if (Hmm[k] != (InitVal + 10)) {
         DataMismatch++;
       }