From 8ba993e36c60275ef23324513f5cb09b95103c81 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Mon, 24 Apr 2017 20:38:37 -0500
Subject: [PATCH] Refactor hipHostRegister to cover misaligned cases.

---
 .../src/runtimeApi/memory/hipHostRegister.cpp | 129 ++++++++++++------
 1 file changed, 84 insertions(+), 45 deletions(-)

diff --git a/tests/src/runtimeApi/memory/hipHostRegister.cpp b/tests/src/runtimeApi/memory/hipHostRegister.cpp
index 1a1319c500..efa23b4068 100644
--- a/tests/src/runtimeApi/memory/hipHostRegister.cpp
+++ b/tests/src/runtimeApi/memory/hipHostRegister.cpp
@@ -19,87 +19,126 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * RUN: %t
+ * RUN: %t --tests 0x1
+ * RUN: %t --tests 0x2
  * HIT_END
  */
 
+// TODO - bug if run both back-to-back
+
 #include"test_common.h"
 #include<malloc.h>
 
 __global__ void Inc(hipLaunchParm lp, float *Ad){
-int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
-Ad[tx] = Ad[tx] + float(1);
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    Ad[tx] = Ad[tx] + float(1);
 }
 
-int main(){
-	float *A, **Ad;
-	int num_devices;
-	HIPCHECK(hipGetDeviceCount(&num_devices));
-	Ad = new float*[num_devices];
-	const size_t size = N * sizeof(float);
-	A = (float*)malloc(size);
-	HIPCHECK(hipHostRegister(A, size, 0));
+
+template<typename T>
+void doMemCopy(size_t numElements, int offset, T *A, T *Bh, T *Bd) 
+{
+    A = A + offset;
+    numElements -= offset;
+
+    size_t sizeBytes = numElements * sizeof(T);
+
+	HIPCHECK(hipHostRegister(A, sizeBytes, 0));
 
 
-	for(int i=0;i<N;i++){
-		A[i] = float(1);
-	}
+    // Reset
+    for(size_t i=0;i<numElements;i++){
+        A[i] = float(i);
+        Bh[i] = 0.0f;
+    }
+
+    HIPCHECK(hipMemset(Bd, 13.0f, sizeBytes));
 
 
-	for(int i=0;i<num_devices;i++){
-        HIPCHECK(hipSetDevice(i));
-        HIPCHECK(hipHostGetDevicePointer((void**)&Ad[i], A, 0));
-	}
+    // 
+    HIPCHECK(hipMemcpy(Bd, A, sizeBytes,  hipMemcpyHostToDevice));
+    HIPCHECK(hipMemcpy(Bh, Bd, sizeBytes, hipMemcpyDeviceToHost));
 
-    // Reference the registered device pointer Ad from inside the kernel:
-	for(int i=0;i<num_devices;i++){
-        HIPCHECK(hipSetDevice(i));
-        hipLaunchKernel(Inc, dim3(N/512), dim3(512), 0, 0, Ad[i]);
+    // Make sure the copy worked
+    for(size_t i=0;i<numElements;i++){
+        if (Bh[i] != A[i]) {
+            printf ("mismatch at Bh[%zu]=%f, A[%zu]=%f\n", i, Bh[i], i, A[i]);
+            failed("mismatch");
+        };
+    }
 
-        HIPCHECK(hipDeviceSynchronize());
-	}
-	HIPASSERT(A[10] == 1.0f + float(num_devices));
+    HIPCHECK(hipHostUnregister(A));
+
+}
+
+int main(int argc, char *argv[])
+{
+
+    HipTest::parseStandardArguments(argc, argv, true);
+
+    const size_t size = N * sizeof(float);
+
+    if (p_tests & 0x1) { 
+        float *A, **Ad;
+        int num_devices;
+        HIPCHECK(hipGetDeviceCount(&num_devices));
+        Ad = new float*[num_devices];
+        A = (float*)malloc(size);
+        HIPCHECK(hipHostRegister(A, size, 0));
 
 
-    { 
+        for(int i=0;i<N;i++){
+            A[i] = float(1);
+        }
+
+
+        for(int i=0;i<num_devices;i++){
+            HIPCHECK(hipSetDevice(i));
+            HIPCHECK(hipHostGetDevicePointer((void**)&Ad[i], A, 0));
+        }
+
+        // Reference the registered device pointer Ad from inside the kernel:
+        for(int i=0;i<num_devices;i++){
+            HIPCHECK(hipSetDevice(i));
+            hipLaunchKernel(Inc, dim3(N/512), dim3(512), 0, 0, Ad[i]);
+
+            HIPCHECK(hipDeviceSynchronize());
+        }
+        HIPASSERT(A[10] == 1.0f + float(num_devices));
+
+        HIPCHECK(hipHostUnregister(A));
+
+        free (A);
+    }
+
+
+    if (p_tests & 0x2) { 
         // Sensitize HIP bug if device does not match where the memory was registered.
         HIPCHECK(hipSetDevice(0));
 
 
+        float * A = (float*)malloc(size);
 
         // Copy to B, this should be optimal pinned malloc copy:
         // Note we are using the host pointer here:
         float *Bh, *Bd;
         Bh = (float*)malloc(size);
         HIPCHECK(hipMalloc(&Bd, size));
-        HIPCHECK(hipMemset(Bd, 13.0f, size));
 
-        for(int i=0;i<N;i++){
-            A[i] = float(i);
-            Bh[i] = 0.0f;
+        // TODO - change to 256:
+#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);
         }
 
-        HIPCHECK(hipMemcpy(Bd, A, size,  hipMemcpyHostToDevice));
-
-        HIPCHECK(hipMemcpy(Bh, Bd, size, hipMemcpyDeviceToHost));
-
-#if 0
-        //TODO - disable check HCC patch for registered/locked memory usin device pointers is merged.
-        for(int i=0;i<N;i++){
-            if (Bh[i] != A[i]) {
-                printf ("mismatch at Bh[%d]=%f, A[%d]=%f\n", i, Bh[i], i, A[i]);
-                failed("mismatch");
-            };
-        }
-#endif
 
 
+        free (A);
 
-        // Make sure the copy worked
     }
 
 
 
-	HIPCHECK(hipHostUnregister(A));
 	passed();
 }