diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 4d75a1d7a1..b78eb7af6e 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -238,8 +238,9 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
 	}
 	if(device){
 	if(flags == hipHostRegisterDefault){
-		hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr);
-		if(hsa_status == HSA_STATUS_SUCCESS){
+                am_status_t am_status = hc::am_memtracker_host_memory_lock(device->_acc, hostPtr, sizeBytes);
+//		hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr);
+		if(am_status == AM_SUCCESS){
 			hip_status = hipSuccess;	
 		}else{
 			hip_status = hipErrorMemoryAllocation;
diff --git a/tests/src/hipHostRegister.cpp b/tests/src/hipHostRegister.cpp
index b9e4632369..7505ac661c 100644
--- a/tests/src/hipHostRegister.cpp
+++ b/tests/src/hipHostRegister.cpp
@@ -28,11 +28,8 @@ Ad[tx] = Ad[tx] + float(1);
 int main(){
 	float *A, *Ad;
 	const size_t size = N * sizeof(float);
-#ifdef __HIP_PLATFORM_NVCC__
 	A = (float*)malloc(size*2);
-#else
-	A = (float*)memalign(64, size);
-#endif
+
 	HIPCHECK(hipHostRegister(A, size, 0));
 
 	for(int i=0;i<N;i++){
@@ -40,6 +37,17 @@ int main(){
 	}
 
 	HIPCHECK(hipMalloc(&Ad, size));
+        hipStream_t stream;
+        HIPCHECK(hipStreamCreate(&stream));
+
+	HIPCHECK(hipMemcpyAsync(Ad, A, size, hipMemcpyHostToDevice, stream));
+
+	hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, stream, Ad);
+	HIPCHECK(hipDeviceSynchronize());
+
+	HIPCHECK(hipMemcpyAsync(A, Ad, size, hipMemcpyDeviceToHost, stream));
+
+	HIPASSERT(A[10] == 2.0f);
 
 	HIPCHECK(hipMemcpy(Ad, A, size, hipMemcpyHostToDevice));
 
@@ -48,7 +56,10 @@ int main(){
 
 	HIPCHECK(hipMemcpy(A, Ad, size, hipMemcpyDeviceToHost));
 
-	HIPASSERT(A[10] == 2.0f);
+	HIPASSERT(A[10] == 3.0f);
+
+
 	HIPCHECK(hipHostUnregister(A));
+        HIPCHECK(hipStreamDestroy(stream));
 	passed();
 }