From 72b420bab473001bb28c6aa0d4ebfd1ae7660b20 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sun, 12 Mar 2017 09:51:33 -0500
Subject: [PATCH] Update hiphostregister test.

Move check to correct place.


[ROCm/clr commit: 9adbbd2980b326a9e55a4021e3252e3ba3fe29d5]
---
 .../tests/src/runtimeApi/memory/hipHostRegister.cpp | 13 ++++++++-----
 1 file changed, 8 insertions(+), 5 deletions(-)

diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
index eae73e1a65..1a1319c500 100644
--- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
+++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipHostRegister.cpp
@@ -51,24 +51,28 @@ int main(){
         HIPCHECK(hipHostGetDevicePointer((void**)&Ad[i], A, 0));
 	}
 
-    // Use device pointer inside a kernel:
+    // Reference the registered device pointer Ad from inside the kernel:
 	for(int i=0;i<num_devices;i++){
         HIPCHECK(hipSetDevice(i));
-        hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, 0, Ad[i]);
+        hipLaunchKernel(Inc, dim3(N/512), dim3(512), 0, 0, Ad[i]);
 
         HIPCHECK(hipDeviceSynchronize());
 	}
+	HIPASSERT(A[10] == 1.0f + float(num_devices));
 
 
     { 
-        // Senstizes HIP bug if device does not match where the memory was registered.
+        // Sensitize HIP bug if device does not match where the memory was registered.
         HIPCHECK(hipSetDevice(0));
 
+
+
         // 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);
@@ -80,7 +84,7 @@ int main(){
         HIPCHECK(hipMemcpy(Bh, Bd, size, hipMemcpyDeviceToHost));
 
 #if 0
-        //TODO - disable check until we update HCC to deal with registered memory pointers.
+        //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]);
@@ -96,7 +100,6 @@ int main(){
 
 
 
-	HIPASSERT(A[10] == 1.0f + float(num_devices));
 	HIPCHECK(hipHostUnregister(A));
 	passed();
 }