From 20315571bd49dfdfb454a5d89ad2d9a9eb0c643f Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Mon, 24 Apr 2017 21:05:29 -0500
Subject: [PATCH] Refactor hipHostRegister test.

- Add more testing for offsets.
- Parse cmdline options and use --tests.


[ROCm/clr commit: e7af4ef641709ad8c91d4c9ebd701a4bb38d571e]
---
 .../src/runtimeApi/memory/hipHostRegister.cpp | 32 ++++++++++++++-----
 1 file changed, 24 insertions(+), 8 deletions(-)

diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
index efa23b4068..8cf0979261 100644
--- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
+++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
@@ -21,10 +21,11 @@ THE SOFTWARE.
  * BUILD: %t %s ../../test_common.cpp
  * RUN: %t --tests 0x1
  * RUN: %t --tests 0x2
+ * RUN: %t --tests 0x4
  * HIT_END
  */
 
-// TODO - bug if run both back-to-back
+// TODO - bug if run both back-to-back, once fixed should just need one command line
 
 #include"test_common.h"
 #include<malloc.h>
@@ -36,14 +37,16 @@ __global__ void Inc(hipLaunchParm lp, float *Ad){
 
 
 template<typename T>
-void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd) 
+void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd, bool internalRegister) 
 {
     A = A + offset;
     numElements -= offset;
 
     size_t sizeBytes = numElements * sizeof(T);
 
-	HIPCHECK(hipHostRegister(A, sizeBytes, 0));
+    if (internalRegister) {
+        HIPCHECK(hipHostRegister(A, sizeBytes, 0));
+    }
 
 
     // Reset
@@ -67,7 +70,9 @@ void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd)
         };
     }
 
-    HIPCHECK(hipHostUnregister(A));
+    if (internalRegister) {
+        HIPCHECK(hipHostUnregister(A));
+    }
 
 }
 
@@ -112,7 +117,7 @@ int main(int argc, char *argv[])
     }
 
 
-    if (p_tests & 0x2) { 
+    if (p_tests & 0x6) { 
         // Sensitize HIP bug if device does not match where the memory was registered.
         HIPCHECK(hipSetDevice(0));
 
@@ -125,11 +130,22 @@ int main(int argc, char *argv[])
         Bh = (float*)malloc(size);
         HIPCHECK(hipMalloc(&Bd, size));
 
-        // TODO - change to 256:
+        // TODO - set to 128
 #define OFFSETS_TO_TRY 1 
         assert (N>OFFSETS_TO_TRY);
-        for (size_t i=0; i<OFFSETS_TO_TRY; i++) {
-            doMemCopy(N, i, A, Bh, Bd);
+
+        if (p_tests & 0x2) {
+            for (size_t i=0; i<OFFSETS_TO_TRY; i++) {
+                doMemCopy(N, i, A, Bh, Bd, true/*internalRegister*/);
+            }
+        }
+
+        if (p_tests & 0x4) {
+            HIPCHECK(hipHostRegister(A, size, 0));
+            for (size_t i=0; i<OFFSETS_TO_TRY; i++) {
+                doMemCopy(N, i, A, Bh, Bd, false/*internalRegister*/);
+            }
+            HIPCHECK(hipHostUnregister(A));
         }