touchup

[ROCm/clr commit: be0d3510f5]

projects/clr/hipamd/bin/hipconfig

@@ -67,7 +67,14 @@ if ($p_cpp_config) {
     $printed = 1;
 }
 
-if ($p_full) {
+
+if ($p_compiler or $p_platform) {
+    print $HIP_PLATFORM;
+    $printed = 1;
+}
+
+
+if (!$printed or $p_full) {
     print "== hipconfig\n";
     print "HIP_PATH     : ", $HIP_PATH, "\n";
     print "HIP_PLATFORM : ", $HIP_PLATFORM, "\n";
@@ -76,6 +83,7 @@ if ($p_full) {
     {
         print "\n" ;
         print "== hcc\n";
+        print ("HSA_HOME     : $HCC_HOME\n");
         print ("HCC_HOME     : $HCC_HOME\n");
         system("$HCC_HOME/bin/hcc --version");
         print ("HCC-cxxflags: "); 
@@ -107,10 +115,6 @@ if ($p_full) {
 }
 
 
-if (!$printed or $p_compiler or $p_platform) {
-    print $HIP_PLATFORM;
-}
-
 if ($p_newline) {
     print "\n";
 }

projects/clr/hipamd/src/hip_hcc.cpp

@@ -46,7 +46,7 @@ THE SOFTWARE.
 
 
 #define USE_AM_TRACKER 1  /* >0 = use new AM memory tracker features. */
-#define USE_ROCR_V2    0  /* use the ROCR v2 async copy API with dst and src agents */
+#define USE_ROCR_V2    1  /* use the ROCR v2 async copy API with dst and src agents */
 
 #if (USE_AM_TRACKER) and (__hcc_workweek__ < 16074)
 #error (USE_AM_TRACKER requries HCC version of 16074 or newer)
@@ -2370,6 +2370,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
 /**
  * @result #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidMemcpyDirection, #hipErrorInvalidValue
  * @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies.
+ * @warning on HCC hipMemcpyAsync requires that any host pointers are pinned (ie via the hipMallocHost call).
  */
 //---
 hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream)