added support pinned dma memcpy between host and device

hipamd/tests/src/CMakeLists.txt

@@ -145,7 +145,7 @@ make_hip_executable (hip_popc hip_popc.cpp)
 make_hip_executable (hip_clz hip_clz.cpp) 
 make_hip_executable (hip_brev hip_brev.cpp) 
 make_hip_executable (hip_ffs hip_ffs.cpp) 
-make_hip_executable (hip_ldg hip_ldg.cpp)
+make_hip_executable (hip_test_ldg hip_test_ldg.cpp)
 make_hip_executable (hipGetDeviceAttribute hipGetDeviceAttribute.cpp) 
 make_hip_executable (hipEnvVar hipEnvVar.cpp) 
 make_hip_executable (hipEnvVarDriver hipEnvVarDriver.cpp) 
@@ -178,6 +178,7 @@ make_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp)
 make_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) 
 make_hip_executable (hipMemcpyAll hipMemcpyAll.cpp)
 make_hip_executable (hipMultiThreadDevice hipMultiThreadDevice.cpp) 
+make_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp)
 
 make_test(hip_ballot " " )
 make_test(hip_anyall " " )
@@ -185,7 +186,7 @@ make_test(hip_popc " " )
 make_test(hip_brev " " )
 make_test(hip_clz " " )
 make_test(hip_ffs " " )
-make_test(hip_ldg " " )
+make_test(hip_test_ldg " " )
 make_test(hipEventRecord --iterations 10)
 make_test(hipMemset " " )
 make_test(hipMemset --N 10    --memsetval 0x42 )  # small copy, just 10 bytes.
@@ -222,6 +223,7 @@ make_test(hipFuncSetDeviceFlags " ")
 make_test(hipFuncGetDevice " ")
 make_test(hipFuncSetDevice " ")
 make_test(hipFuncDeviceSynchronize " ")
+make_test(hipTestMemcpyPin " ")
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-serial" --tests 0x1) 
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4) 
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) 

hipamd/tests/src/hipTestMemcpyPin.cpp

@@ -0,0 +1,33 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include<hip_runtime.h>
+#include"test_common.h"
+
+#define len 1024*1024
+#define size len * sizeof(float)
+
+int main(){
+    float *Ad, *A;
+    hipHostMalloc((void**)&A, size);
+    hipMalloc((void**)&Ad, size);
+    assert(hipSuccess == hipMemcpy(Ad, A, size, hipMemcpyHostToDevice));
+    assert(hipSuccess == hipMemcpy(A, Ad, size, hipMemcpyDeviceToHost));
+    passed();
+}

hipamd/tests/src/hip_test_ldg.cpp

@@ -25,7 +25,7 @@ THE SOFTWARE.
 #include <stdlib.h>
 #include<iostream>
 #include "hip_runtime.h"
-
+#include "test_common.h"
 
 #define HIP_ASSERT(x) (assert((x)==hipSuccess))
 
@@ -125,7 +125,6 @@ bool dataTypesRun(){
     printf("FAILED: %d errors\n",errors);
     ret = false;
   } else {
-      printf ("PASSED!\n");
       ret = true;
   }
 
@@ -150,13 +149,14 @@ int main() {
   cout << " System major " << devProp.major << endl;
   cout << " agent prop name " << devProp.name << endl;
 
-  int errors;
-    errors = dataTypesRun<char>();
-    errors = dataTypesRun<signed char>();
-    errors = dataTypesRun<short>();
-    errors = dataTypesRun<int>();
-  cout << "__ldg " << endl ;
+   int errors = dataTypesRun<char>() &
+           dataTypesRun<signed char>() &
+           dataTypesRun<short>() &
+           dataTypesRun<int>();
   //hipResetDefaultAccelerator();
+    if(errors == 1){
+        passed();
+        return 0;
+    }
 
-  return errors;
 }