Merge branch 'amd-develop' into amd-master

Change-Id: I32d41081ac065f2c50531dc2e420802d765665e2 [ROCm/hip commit: 8d40253664]
2016-11-14 06:11:47 +05:30
@@ -1,4 +1,5 @@
 .*
+!.gitignore
 *.o
 *.exe
 *.swp
@@ -1,4 +1 @@
-:set tabstop=4
-:set shiftwidth=4
-:set expandtab
-:set smartindent
+:set makeprg=make\ -C\ build.hcc-LC.db
@@ -142,6 +142,7 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER)
 endif()
 add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER)

+
 #############################
 # Build steps
 #############################
@@ -164,7 +165,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
    set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}")

    # Add remaining flags
-    set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include -stdlib=libc++")
+    set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include -I/usr/local/include/c++/v1 -stdlib=libc++")

    # Set compiler and compiler flags
    set(CMAKE_CXX_COMPILER "${HCC_HOME}/bin/hcc")
@@ -337,7 +338,7 @@ endif()
 add_custom_target(install_for_test COMMAND "${CMAKE_COMMAND}" --build . --target install
    WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
 execute_process(COMMAND getconf _NPROCESSORS_ONLN OUTPUT_VARIABLE DASH_JAY OUTPUT_STRIP_TRAILING_WHITESPACE)
-add_custom_target(test COMMAND ${CMAKE_COMMAND} .
+add_custom_target(test COMMAND ${CMAKE_COMMAND} -DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE} .
    COMMAND make -j ${DASH_JAY}
    COMMAND make test
    WORKING_DIRECTORY ${BUILD_DIR}
@@ -124,6 +124,11 @@ Differences or limitations of HIP APIs as compared to CUDA APIs should be clearl
 - ihipLogStatus should only be called from top-level HIP APIs,and should be called to log and return the error code.  The error code 
  is used by the GetLastError and PeekLastError functions - if a HIP API simply returns, then the error will not be logged correctly.

+- All HIP environment variables should begin with the keyword HIP_
+    Environment variables should be long enough to describe their purpose but short enough so they can be remembered - perhaps 10-20 characters, with 3-4 parts separated by underscores.
+    To see the list of current environment variables, along with their values, set HIP_PRINT_ENV and run any hip applications on ROCM platform .
+    HIPCC or other tools may support additional environment variables which should follow the above convention.  
+


 #### Presubmit Testing:
@@ -13,8 +13,25 @@ Upcoming:

 ## Revision History:

+===================================================================================================
+Release:1.0
+Date: 2016.11.8
+- Initial implementation for FindHIP.cmake
+- HIP library now installs as a static library by default
+- Added support for HIP context and HIP module APIs
+- Major changes to HIP signal & memory management implementation
+- Support for complex data type and math functions
+- clang-hipify is now known as hipify-clang
+- Added several new HIP samples
+- Preliminary support for new APIs: hipMemcpyToSymbol, hipDeviceGetLimit, hipRuntimeGetVersion
+- Added support for async memcpy driver API (for example hipMemcpyHtoDAsync)
+- Support for memory management device functions: malloc, free, memcpy & memset
+- Removed deprecated HIP runtime header locations. Please include "hip/hip_runtime.h" instead of "hip_runtime.h". You can use `find . -type f -exec sed -i 's:#include "hip_runtime.h":#include "hip/hip_runtime.h":g' {} +` to replace all such references
+
+
 ===================================================================================================
 Release:0.92.00
+Date: 2016.8.14
 - hipLaunchKernel supports one-dimensional grid and/or block dims, without explicit cast to dim3 type (actually in 0.90.00)
 - fp16 software support
 - Support for Hawaii dGPUs using environment variable ROCM_TARGET=hawaii
@@ -79,7 +79,7 @@ if ($HIP_PLATFORM eq "hcc") {

    $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";

-    $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'};
+    $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'} // 1;
    $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger";

    $ROCM_TARGET=$ENV{'ROCM_TARGET'} // "fiji";
@@ -116,9 +116,16 @@ if ($HIP_PLATFORM eq "hcc") {
    } else {
        $HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=1 -Wl,--defsym=_binary_kernel_spir_start=1 -Wl,--defsym=_binary_kernel_cl_start=1 -Wl,--defsym=_binary_kernel_cl_end=1";
    }
+    if ($HOST_OSNAME eq "fedora") {
+        $HIPCXXFLAGS .= " -I/usr/local/include/c++/v1";
+    }

    # Satisfy HCC dependencies
-    $HIPLDFLAGS .= " -lc++abi -lsupc++";
+    if ($HOST_OSNAME eq "fedora") {
+        $HIPLDFLAGS .= " -lc++abi";
+    } else {
+        $HIPLDFLAGS .= " -lc++abi -lsupc++";
+    }
    $HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt";

    # Handle ROCm target platform
@@ -273,10 +280,9 @@ foreach $arg (@ARGV)
        # Process HIPCC options here:
        if ($arg =~ m/^--hipcc/) {
            $swallowArg = 1;
-            if ($arg eq "--hipcc_explicit_lib") {
-                # Some environments (ie cmake tests) already link their own hip_hcc.o, so don't add here:
-                $needHipHcc = 0;
-            }
+            #if $arg eq "--hipcc_profile") {  # Example argument here, hipcc
+            #
+            #}
        } else {
            push (@options, $arg);
        }
@@ -314,14 +320,15 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
 }

 if ($needHipHcc) {
-    $HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 0;
+    $HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 1;

+    # TODO - remove the old sea-of-objects solution:
    if ($HIP_LIB_TYPE eq 0) {
-        $HIPLDFLAGS .= " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/hip_ldg.cpp.o $HIP_PATH/lib/hip_fp16.cpp.o $HIP_PATH/lib/hip_context.cpp.o $HIP_PATH/lib/hip_module.cpp.o";
+        substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/hip_ldg.cpp.o $HIP_PATH/lib/hip_fp16.cpp.o $HIP_PATH/lib/hip_context.cpp.o $HIP_PATH/lib/hip_module.cpp.o ";
    } elsif ($HIP_LIB_TYPE eq 1) {
-        $HIPLDFLAGS .= " -L$HIP_PATH/lib -lhip_hcc" ;
+        substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc " ;
    } else {
-        $HIPLDFLAGS .= " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc";
+        substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc ";
    }
 }

@@ -353,7 +360,7 @@ if ($printHipVersion) {
 }
 if ($runCmd) {
    if ($HIP_PLATFORM eq "hcc" and exists($hipConfig{'HCC_VERSION'}) and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) {
-        print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using hcc $HCC_VERSION. Please rebuild HIP including cmake.\n") && die ();
+        print ("HIP ($HIP_PATH) was built using hcc $hipConfig{'HCC_VERSION'}, but you are using $HCC_HOME/hcc with version $HCC_VERSION from hipcc. Please rebuild HIP including cmake or update HCC_HOME variable.\n") && die ();
    }
    system ("$CMD") and die ();
 }
@@ -0,0 +1,19 @@
+#!/bin/bash
+
+# usage: hipdemangleatp.sh ATP_FILE
+
+# HIP kernels
+kernels=$(grep grid_launch_parm $1 | cut -d" " -f1 | sort | uniq)
+for mangled_sym in $kernels; do
+    real_sym=$(c++filt -p $(c++filt _$mangled_sym | cut -d: -f3 | sed 's/_functor//g' | sed 's/ /\\\&nbsp/g'))
+    #echo "$mangled_sym => $real_sym" >> $1.log
+    sed -i "s/$mangled_sym/$real_sym/g" $1
+done
+
+# HC kernels
+kernels=$(grep cxxamp_trampoline $1 | cut -d" " -f1 | sort | uniq)
+for mangled_sym in $kernels; do
+    real_sym=$(echo $mangled_sym |  sed "s/^/_/g; s/_EC_/_$/g" | c++filt -p | cut -d\( -f1 | cut -d" " -f1 --complement | sed 's/ /\\\&nbsp/g')
+    #echo "$mangled_sym => $real_sym" >> $1.log
+    sed -i "s/$mangled_sym/$real_sym/g" $1
+done
@@ -229,43 +229,7 @@ If platform portability is important, use #ifdef __HIP_PLATFORM_HIPCC__ to guard


 ### How do I trace HIP application flow?
-#### Using CodeXL markers for HIP Functions
-HIP can generate markers at function being/end which are displayed on the CodeXL timeline view.
-To do this, you need to install ROCm-Profiler and enable HIP to generate the markers:
-
-1. Install ROCm-Profiler
-Installing HIP from the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) pre-built packages, installs the ROCm-Profiler as well.
-Alternatively, you can build ROCm-Profiler using the instructions [here](https://github.com/RadeonOpenCompute/ROCm-Profiler#building-the-rocm-profiler).
-
-2. Build HIP with ATP markers enabled
-HIP pre-built packages are enabled with ATP marker support by default.
-To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step.
-
-3. Set HIP_ATP_MARKER
-```shell
-export HIP_ATP_MARKER=1
-```
-
-4. Recompile the target application
-
-5. Run with profiler enabled to generate ATP file.
-```shell
-# Use profile to generate timeline view:
-/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>
-
-Or
-/opt/rocm/bin/rocm-profiler -e HIP_ATP_MARKER=1 -o <outputATPFileName> -A <applicationName> <applicationArguments>
-```
-
-#### Using HIP_TRACE_API
-You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided
-by the HIP_DB switch. For example:
-```shell
-# Trace to stderr showing being/end of each function (with arguments) + intermediate debug trace during the execution of each function.
-HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp
-```
-
-Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.
+See the [HIP Profiling Guide](hip_porting_guide.md) for more information.

 ### What if HIP generates error of "symbol multiply defined!" only on AMD machine?
 Unlike CUDA, in HCC, for functions defined in the header files, the keyword of "__forceinline__" does not imply "static".
@@ -233,8 +233,11 @@ typedef struct dim3 {
 ## Memory-Fence Instructions
 HIP supports __threadfence() and  __threadfence_block().

-Applications that use threadfence_system can disable the L1 and L2 caches on the GPU by:
-"export HSA_DISABLE_CACHE=1".  See the hip_porting_guide.md#threadfence_system for more information.
+HIP provides workaround for threadfence_system() under HCC path.
+To enable the workaround, HIP should be built with environment variable HIP_COHERENT_HOST_ALLOC enabled.
+In addition,the kernels that use __threadfence_system() should be modified as follows:
+- The kernel should only operate on finegrained system memory; which should be allocated with hipHostMalloc().
+- Remove all memcpy for those allocated finegrained system memory regions.

 ## Synchronization Functions
 The __syncthreads() built-in function is supported in HIP. The __syncthreads_count(int), __syncthreads_and(int) and __syncthreads_or(int) functions are under development.  
@@ -564,7 +564,7 @@ HIP_LAUNCH_BLOCKING            =  0 : Make HIP APIs 'host-synchronous', so they
 HIP_DB                         =  0 : Print various debug info.  Bitmask, see hip_hcc.cpp for more information.
 HIP_TRACE_API                  =  0 : Trace each HIP API call.  Print function name and return code to stderr as program executes.
 HIP_TRACE_API_COLOR            = green : Color to use for HIP_API.  None/Red/Green/Yellow/Blue/Magenta/Cyan/White
-HIP_ATP_MARKER                 =  0 : Add HIP function begin/end to ATP file generated with CodeXL
+HIP_PROFILE_API                 =  0 : Add HIP function begin/end to ATP file generated with CodeXL
 HIP_VISIBLE_DEVICES            =  0 : Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence
 HIP_NUM_KERNELS_INFLIGHT       = 128 : Number of kernels per stream

@@ -0,0 +1,95 @@
+# Profiling HIP Code
+
+HIP provides several capabilities to support debugging and profiling.  Profiling information can be displayed to stderr or viewed in the CodeXl visualization tool.
+
+### Usign CodeXL to profile a HIP Application
+By defauly, CodeXL can trace all kernel commands, data transfer commands, and HSA Runtime (ROCr) API calls.  
+/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>
+
+### Using CodeXL markers for HIP Functions
+HIP can generate markers at function being/end which are displayed on the CodeXL timeline view.
+HIP 1.0 compiles marker support by default, and you can enable it by setting the HIP_PROFILE_API environment variable and then running the rocm-profiler:
+
+```shell
+
+# Use profile to generate timeline view:
+export HIP_PROFILE_API=1
+/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>
+
+Or
+/opt/rocm/bin/rocm-profiler -e HIP_PROFILE_API=1 -o <outputATPFileName> -A <applicationName> <applicationArguments>
+```
+
+#### Developer Builds
+For developer builds, you must enable marker support manually when compiling HIP.
+
+1. Build HIP with ATP markers enabled
+HIP pre-built packages are enabled with ATP marker support by default.
+To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step.
+
+2. Install ROCm-Profiler
+Installing HIP from the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) pre-built packages, installs the ROCm-Profiler as well.
+Alternatively, you can build ROCm-Profiler using the instructions [here](https://github.com/RadeonOpenCompute/ROCm-Profiler#building-the-rocm-profiler).
+
+3. Recompile the target application
+
+Then follow the steps above to collect a marker-enabled trace.
+
+
+### Using HIP_TRACE_API
+You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided
+by the HIP_DB switch. For example:
+```shell
+# Trace to stderr showing being/end of each function (with arguments) + intermediate debug trace during the execution of each function.
+HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp
+```
+
+#### Color
+Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.  
+You can change the color used for the trace mode with the HIP_TRACE_API_COLOR environment variable.  Possible values are None/Red/Green/Yellow/Blue/Magenta/Cyan/White.
+None will disable use of color control codes and may be useful when saving the trace file or when a pure text trace is desired.
+
+#### 
+
+
+### Using HIP_DB
+
+This flag is primarily targeted to assist HIP development team in the development of the HIP runtime, but in some situations may be useful to HIP application developers as well.
+The HIP debug information is designed to print important information during the execution of a HIP API.  HIP provides
+different color-coded levels of debug informaton:
+  - api  : Print the beginning and end of each HIP API, including the arguments and return codes.
+  - sync : Print multi-thread and other synchronization debug information.
+  - copy : Print which engine is doing the copy, which copy flavor is selected, information on source and destination memory.
+  - mem  : Print information about memory allocation - which pointers are allocated, where they are allocated, peer mappings, and more.
+
+DB_MEM format is flags separated by '+' sign, or a hex code for the bitmask.  Generally the + format is preferred.  
+For example:
+```shell
+HIP_DB=api+copy+mem  my-application
+HIP_DB=0xF  my-application
+```
+HIP_DB=1 same as HIP_TRACE_API=1
+
+
+
+
+Trace provides quick look at API.
+Explain output of 
+Reference the cookbook example.
+Command-line profile.
+/// disable profiling at the start of the application you can start CodeXLGpuProfiler with the --startdisabled flag.
+
+Can use strace interleaved with HSA Debug calls .
+
+HIP_PROFILE_API=1
+HIP_PROFILE_API=2  : Will show the full API in the trace.  This can be useful for lower-level debugging when you want to see all the parameters that are passed to a specific API.
+
+demangle atp
+
+Write how to collect performance counters.
+- include how to compute bandwidth for copy and kernel activity.
+
+- How to disable HSA APIs.
+- Do I need to use profiler with HSA enabled?  Do I need to enable HSA profiling on the command line?
+
+Offline compile, how to visualize.  
@@ -423,7 +423,11 @@ struct cuda2hipMap {
    cuda2hipRename["cuMemHostRegister_v2"]                      = {"hipHostRegister", CONV_MEM, API_DRIVER};
    cuda2hipRename["cuMemHostUnregister"]                       = {"hipHostUnregister", CONV_MEM, API_DRIVER};

-
+    // Profiler
+    // unsupported yet by HIP
+    // cuda2hipRename["cuProfilerInitialize"]                   = {"hipProfilerInitialize", CONV_OTHER, API_DRIVER};
+    cuda2hipRename["cuProfilerStart"]                           = {"hipProfilerStart", CONV_OTHER, API_DRIVER};
+    cuda2hipRename["cuProfilerStop"]                            = {"hipProfilerStop", CONV_OTHER, API_DRIVER};

    /////////////////////////////// CUDA RT API ///////////////////////////////
    // Error API
@@ -1606,7 +1610,11 @@ private:
        }
      }
      XStr.clear();
-      OS << "hipLaunchKernel(HIP_KERNEL_NAME(" << calleeName << "),";
+      if (calleeName.find(',') != StringRef::npos) {
+        SmallString<128> tmpData;
+        calleeName = Twine("HIP_KERNEL_NAME(" + calleeName + ")").toStringRef(tmpData);
+      }
+      OS << "hipLaunchKernel(" << calleeName << ",";
      const CallExpr *config = launchKernel->getConfig();
      DEBUG(dbgs() << "Kernel config arguments:" << "\n");
      SourceManager *SM = Result.SourceManager;
@@ -556,10 +556,19 @@ extern "C" __device__ void __threadfence(void);
 *
 * @param void
 *
- * @warning __threadfence_system is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
+ * @warning __threadfence_system is a stub and map to no-op.
 */
-__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
+__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));

+__device__ unsigned __hip_ds_bpermute(int index, unsigned src);
+__device__ float __hip_ds_bpermutef(int index, float src);
+__device__ unsigned __hip_ds_permute(int index, unsigned src);
+__device__ float __hip_ds_permutef(int index, float src);
+
+__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern);
+__device__ float __hip_ds_swizzlef(float src, int pattern);
+
+__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl);

 // doxygen end Fence Fence
 /**
@@ -1278,7 +1278,11 @@ hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);
 hipError_t  hipDeviceDisablePeerAccess (int peerDeviceId);


-#ifdef PEER_NON_UNIFIED
+#ifndef USE_PEER_NON_UNIFIED
+#define USE_PEER_NON_UNIFIED 1
+#endif
+
+#if USE_PEER_NON_UNIFIED==1
 /**
 * @brief Copies memory from one device to memory on another device.
 *
@@ -1591,6 +1595,16 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device);
 */
 hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device);

+/**
+ * @brief Returns a handle to a compute device.
+ * @param [out] device handle
+ * @param [in] PCI Bus ID
+ *
+ * @returns #hipSuccess, #hipErrorInavlidDevice, #hipErrorInvalidValue
+ */
+hipError_t hipDeviceGetByPCIBusId ( int*  device,const int* pciBusId );
+
+
 /**
 * @brief Returns the total amount of memory on the device.
 * @param [out] bytes
@@ -1737,11 +1751,24 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 *
 *  @warning The cudaProfilerInitialize API format for "configFile" is not supported.
 *
- *  On AMD platforms, hipProfilerStart and hipProfilerStop require installation of AMD's GPU
- *  perf counter API and defining GPU_PERF
 */


+// TODO - expand descriptions:
+/**
+ * @brief Start recording of profiling information
+ * @warning : hipProfilerStart API is under development.
+ */
+hipError_t hipProfilerStart();
+
+
+/**
+ * @brief Stop recording of profiling information.
+ * @warning : hipProfilerStop API is under development.
+ */
+hipError_t hipProfilerStop();
+
+
 /**
 * @}
 */
@@ -0,0 +1,38 @@
+/*
+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.
+*/
+
+#pragma once
+
+#if not defined (ENABLE_HIP_PROFILE)
+#define ENABLE_HIP_PROFILE 1
+#endif
+
+#if defined(__HIP_PLATFORM_HCC__) and (ENABLE_HIP_PROFILE==1)
+#include <CXLActivityLogger.h>
+#define HIP_SCOPED_MARKER(markerName, group) amdtScopedMarker __scopedMarker(markerName, group, nullptr);
+#define HIP_BEGIN_MARKER(markerName, group) amdtBeginMarker(markerName, group, nullptr);
+#define HIP_END_MARKER() amdtEndMarker();
+#else
+#define HIP_SCOPED_MARKER(markerName, group) 
+#define HIP_BEGIN_MARKER(markerName, group)
+#define HIP_END_MARKER() 
+#endif
@@ -24,6 +24,7 @@ THE SOFTWARE.

 #include <cuda_runtime_api.h>
 #include <cuda.h>
+#include <cuda_profiler_api.h>

 #ifdef __cplusplus
 extern "C" {
@@ -83,6 +84,10 @@ typedef CUmodule hipModule_t;
 typedef CUfunction hipFunction_t;
 typedef CUdeviceptr hipDeviceptr_t;

+// Flags that can be used with hipStreamCreateWithFlags
+#define hipStreamDefault            cudaStreamDefault
+#define hipStreamNonBlocking        cudaStreamNonBlocking
+
 //typedef cudaChannelFormatDesc hipChannelFormatDesc;
 #define hipChannelFormatDesc cudaChannelFormatDesc

@@ -106,7 +111,7 @@ switch(cuError) {
    case cudaErrorHostMemoryNotRegistered        : return hipErrorHostMemoryNotRegistered     ;
    case cudaErrorUnsupportedLimit               : return hipErrorUnsupportedLimit            ;
    default                                      : return hipErrorUnknown;  // Note - translated error.
-};
+}
 }

 inline static hipError_t hipCUResultTohipError(CUresult cuError) { //TODO Populate further
@@ -120,7 +125,7 @@ switch(cuError) {
    case CUDA_ERROR_INVALID_CONTEXT              : return hipErrorInvalidContext              ;
    case CUDA_ERROR_NOT_INITIALIZED              : return hipErrorNotInitialized              ;
    default                                      : return hipErrorUnknown;  // Note - translated error.
-};
+}
 }

 // TODO   match the error enum names of hip and cuda
@@ -167,8 +172,8 @@ inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) {
 /**
 * Stream CallBack struct
 */
-typedef void(* hipStreamCallback_t)(hipStream_t stream,  hipError_t status, void* userData);
-
+#define HIPRT_CB CUDART_CB
+typedef void(HIPRT_CB * hipStreamCallback_t)(hipStream_t stream,  hipError_t status, void* userData);
 inline static hipError_t hipInit(unsigned int flags)
 {
    return hipCUResultTohipError(cuInit(flags));
@@ -319,11 +324,11 @@ inline static hipError_t hipDeviceSynchronize() {
 }

 inline static const char* hipGetErrorString(hipError_t error){
-    return cudaGetErrorString( hipErrorToCudaError(error) );
+    return cudaGetErrorString(hipErrorToCudaError(error));
 }

 inline static const char* hipGetErrorName(hipError_t error){
-    return cudaGetErrorName( hipErrorToCudaError(error) );
+    return cudaGetErrorName(hipErrorToCudaError(error));
 }

 inline static hipError_t hipGetDeviceCount(int * count){
@@ -585,8 +590,7 @@ inline static hipError_t hipStreamQuery(hipStream_t stream)

 inline static hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void *userData, unsigned int flags)
 {
-    return hipCUDAErrorTohipError(cudaStreamAddCallback(cudaStream_t stream,
-        cudaStreamCallback_t callback, void *userData, unsigned int flags));
+    return hipCUDAErrorTohipError(cudaStreamAddCallback(stream, (cudaStreamCallback_t)callback, userData, flags));
 }

 inline static hipError_t hipDriverGetVersion(int *driverVersion)
@@ -611,12 +615,12 @@ inline static hipError_t hipDeviceCanAccessPeer ( int* canAccessPeer, int  devic

 inline static hipError_t  hipDeviceDisablePeerAccess ( int  peerDevice )
 {
-    return hipCUDAErrorTohipError(cudaDeviceDisablePeerAccess ( peerDevice ));
-};
+    return hipCUDAErrorTohipError(cudaDeviceDisablePeerAccess(peerDevice));
+}

 inline static hipError_t  hipDeviceEnablePeerAccess ( int  peerDevice, unsigned int  flags )
 {
-    return hipCUDAErrorTohipError(cudaDeviceEnablePeerAccess ( peerDevice, flags ));
+    return hipCUDAErrorTohipError(cudaDeviceEnablePeerAccess(peerDevice, flags));
 }

 inline static hipError_t  hipCtxDisablePeerAccess ( hipCtx_t peerCtx )
@@ -626,22 +630,33 @@ inline static hipError_t  hipCtxDisablePeerAccess ( hipCtx_t peerCtx )

 inline static hipError_t  hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned int  flags )
 {
-    return hipCUResultTohipError(cuCtxEnablePeerAccess ( peerCtx, flags ));
+    return hipCUResultTohipError(cuCtxEnablePeerAccess(peerCtx, flags));
 }

 inline static hipError_t hipMemcpyPeer ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t count )
 {
-    return hipCUDAErrorTohipError(cudaMemcpyPeer ( dst, dstDevice, src, srcDevice, count ));
-};
+    return hipCUDAErrorTohipError(cudaMemcpyPeer(dst, dstDevice, src, srcDevice, count));
+}

 inline static hipError_t hipMemcpyPeerAsync ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t count, hipStream_t stream=0 )
 {
-    return hipCUDAErrorTohipError(cudaMemcpyPeerAsync ( dst, dstDevice, src, srcDevice, count, stream ));
-};
+    return hipCUDAErrorTohipError(cudaMemcpyPeerAsync(dst, dstDevice, src, srcDevice, count, stream));
+}
+
+// Profile APIs:
+inline hipError_t hipProfilerStart()
+{
+    return hipCUDAErrorTohipError(cudaProfilerStart());
+}
+
+inline hipError_t hipProfilerStop()
+{
+    return hipCUDAErrorTohipError(cudaProfilerStop());
+}

 inline static hipError_t hipSetDeviceFlags (unsigned int flags)
 {
-    return hipCUDAErrorTohipError(cudaSetDeviceFlags( flags ));
+    return hipCUDAErrorTohipError(cudaSetDeviceFlags(flags));
 }

 inline static hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned int flags)
@@ -656,62 +671,62 @@ inline static hipError_t hipEventQuery(hipEvent_t event)

 inline static hipError_t  hipCtxCreate(hipCtx_t *ctx, unsigned int flags,  hipDevice_t device)
 {
-    return hipCUResultTohipError(cuCtxCreate ( ctx,flags,device ));
+    return hipCUResultTohipError(cuCtxCreate(ctx,flags,device));
 }

 inline static hipError_t  hipCtxDestroy(hipCtx_t ctx)
 {
-    return hipCUResultTohipError(cuCtxDestroy ( ctx ));
+    return hipCUResultTohipError(cuCtxDestroy(ctx));
 }

 inline static hipError_t  hipCtxPopCurrent(hipCtx_t* ctx)
 {
-    return hipCUResultTohipError(cuCtxPopCurrent ( ctx ));
+    return hipCUResultTohipError(cuCtxPopCurrent(ctx));
 }

 inline static hipError_t  hipCtxPushCurrent(hipCtx_t ctx)
 {
-    return hipCUResultTohipError(cuCtxPushCurrent ( ctx ));
+    return hipCUResultTohipError(cuCtxPushCurrent(ctx));
 }

 inline static hipError_t  hipCtxSetCurrent(hipCtx_t ctx)
 {
-    return hipCUResultTohipError(cuCtxSetCurrent ( ctx ));
+    return hipCUResultTohipError(cuCtxSetCurrent(ctx));
 }

 inline static hipError_t  hipCtxGetCurrent(hipCtx_t* ctx)
 {
-    return hipCUResultTohipError(cuCtxGetCurrent ( ctx ));
+    return hipCUResultTohipError(cuCtxGetCurrent(ctx));
 }

 inline static hipError_t  hipCtxGetDevice(hipDevice_t *device)
 {
-    return hipCUResultTohipError(cuCtxGetDevice ( device ));
+    return hipCUResultTohipError(cuCtxGetDevice(device));
 }

 inline static hipError_t  hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion)
 {
-    return hipCUResultTohipError(cuCtxGetApiVersion ( ctx,(unsigned int*)apiVersion ));
+    return hipCUResultTohipError(cuCtxGetApiVersion(ctx,(unsigned int*)apiVersion));
 }

-inline static hipError_t  hipCtxGetCacheConfig ( hipFuncCache *cacheConfig )
+inline static hipError_t  hipCtxGetCacheConfig (hipFuncCache *cacheConfig)
 {
-    return hipCUResultTohipError(cuCtxGetCacheConfig ( cacheConfig ));
+    return hipCUResultTohipError(cuCtxGetCacheConfig(cacheConfig));
 }

-inline static hipError_t  hipCtxSetCacheConfig ( hipFuncCache cacheConfig )
+inline static hipError_t  hipCtxSetCacheConfig (hipFuncCache cacheConfig)
 {
-    return hipCUResultTohipError(cuCtxSetCacheConfig ( cacheConfig ));
+    return hipCUResultTohipError(cuCtxSetCacheConfig(cacheConfig));
 }

-inline static hipError_t  hipCtxSetSharedMemConfig ( hipSharedMemConfig config )
+inline static hipError_t  hipCtxSetSharedMemConfig (hipSharedMemConfig config)
 {
-    return hipCUResultTohipError(cuCtxSetSharedMemConfig ( config ));
+    return hipCUResultTohipError(cuCtxSetSharedMemConfig(config));
 }

 inline static hipError_t  hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig )
 {
-    return hipCUResultTohipError(cuCtxGetSharedMemConfig ( pConfig ));
+    return hipCUResultTohipError(cuCtxGetSharedMemConfig(pConfig));
 }

 inline static hipError_t  hipCtxSynchronize ( void )
@@ -721,7 +736,7 @@ inline static hipError_t  hipCtxSynchronize ( void )

 inline static hipError_t  hipCtxGetFlags ( unsigned int* flags )
 {
-    return hipCUResultTohipError(cuCtxGetFlags ( flags ));
+    return hipCUResultTohipError(cuCtxGetFlags(flags));
 }

 inline static hipError_t hipCtxDetach(hipCtx_t ctx)
@@ -749,6 +764,11 @@ inline static hipError_t hipDeviceGetPCIBusId(int *pciBusId,int len,hipDevice_t
    return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device));
 }

+inline static hipError_t hipDeviceGetByPCIBusId(int* device, const int *pciBusId)
+{
+    return hipCUDAErrorTohipError(cudaDeviceGetByPCIBusId(device,(char*)pciBusId));
+}
+
 inline static hipError_t hipDeviceGetLimit(size_t *pValue, hipLimit_t limit)
 {
    return hipCUDAErrorTohipError(cudaDeviceGetLimit(pValue, limit));
@@ -798,6 +818,8 @@ inline static hipError_t hipModuleLaunchKernel(hipFunction_t f,
                    sharedMemBytes, stream, kernelParams, extra));
 }

+
+
 #ifdef __cplusplus
 }
 #endif
@@ -9,6 +9,7 @@ else()
    install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
 endif()
 install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
+install(FILES @hip_SOURCE_DIR@/src/hip_ir.ll DESTINATION lib)

 #############################
 # Packaging steps
@@ -0,0 +1,5 @@
+vadd_amp_arrayview
+vadd_hc_am
+vadd_hc_array
+vadd_hc_arrayview
+vadd_hip
@@ -5,8 +5,8 @@ OPT=-O2
 HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` ${OPT}
 HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags` ${OPT}

-CPPAMP_CFLAGS= -std=c++amp -stdlib=libc++ -I$(HCC_HOME)/include
-CPPAMP_LDFLAGS=  -std=c++amp -L$(HCC_HOME)/lib -Wl,--rpath=$(HCC_HOME)/lib -lc++ -lc++abi -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive 
+CPPAMP_CFLAGS= `$(HCC_HOME)/bin/clamp-config --cxxflags`
+CPPAMP_LDFLAGS= `$(HCC_HOME)/bin/clamp-config --ldflags`

 HIP_PATH?= $(wildcard /opt/rocm/hip)
 ifeq (,$(HIP_PATH))
@@ -0,0 +1,5 @@
+runKernel.hip.out
+vcpy_isa.code
+vcpy_isa.hsaco
+vcpy_kernel.co
+vcpy_kernel.code
@@ -1,7 +1,4 @@
 HIP_PATH?= $(wildcard /opt/rocm/hip)
-ifeq (,$(HIP_PATH))
-	HIP_PATH=../../..
-endif
 HIPCC=$(HIP_PATH)/bin/hipcc

 all: square.hip.out
@@ -11,9 +8,10 @@ square.cuda.out : square.cu

 #hipify square.cu > square.cpp
 # Then review & finish port in square.cpp
+#

 square.hip.out: square.hipref.cpp
-	$(HIPCC) square.hipref.cpp -o $@
+	$(HIPCC) $(CXXFLAGS) square.hipref.cpp -o $@



@@ -133,6 +133,15 @@ void printDeviceProp (int deviceId)
        }
    }
    cout << endl;
+    cout << setw(w1) << "non-peers: ";
+    for (int i=0; i<deviceCnt; i++) {
+        int isPeer;
+        hipDeviceCanAccessPeer(&isPeer, i, deviceId);
+        if (!isPeer) {
+            cout << "device#" << i << " ";
+        }
+    }
+    cout << endl;



@@ -1,172 +0,0 @@
-/*
-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<iostream>
-
-// hip header file
-#include "hip/hip_runtime.h"
-
-#define WIDTH     1024
-
-#define NUM       (WIDTH*WIDTH)
-
-#define THREADS_PER_BLOCK_X  4
-#define THREADS_PER_BLOCK_Y  4
-#define THREADS_PER_BLOCK_Z  1
-
-// Device (Kernel) function, it must be void
-// hipLaunchParm provides the execution configuration
-__global__ void matrixTranspose(hipLaunchParm lp,
-                                float *out,
-                                float *in,
-                                const int width)
-{
-    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
-
-    out[y * width + x] = in[x * width + y];
-}
-
-// CPU implementation of matrix transpose
-void matrixTransposeCPUReference(
-    float * output,
-    float * input,
-    const unsigned int width)
-{
-    for(unsigned int j=0; j < width; j++)
-    {
-        for(unsigned int i=0; i < width; i++)
-        {
-            output[i*width + j] = input[j*width + i];
-        }
-    }
-}
-
-int main() {
-
-  float* Matrix;
-  float* TransposeMatrix;
-  float* cpuTransposeMatrix;
-
-  float* gpuMatrix;
-  float* gpuTransposeMatrix;
-
-  hipDeviceProp_t devProp;
-  hipGetDeviceProperties(&devProp, 0);
-
-  std::cout << "Device name " << devProp.name << std::endl;
-
-  hipEvent_t start, stop;
-  hipEventCreate(&start);
-  hipEventCreate(&stop);
-  float eventMs = 1.0f;
-
-  int i;
-  int errors;
-
-  Matrix = (float*)malloc(NUM * sizeof(float));
-  TransposeMatrix = (float*)malloc(NUM * sizeof(float));
-  cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
-
-  // initialize the input data
-  for (i = 0; i < NUM; i++) {
-    Matrix[i] = (float)i*10.0f;
-  }
-
-  // allocate the memory on the device side
-  hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
-  hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
-
-  // Record the start event
-  hipEventRecord(start, NULL);
-
-  // Memory transfer from host to device
-  hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
-
-  // Record the stop event
-  hipEventRecord(stop, NULL);
-  hipEventSynchronize(stop);
-
-  hipEventElapsedTime(&eventMs, start, stop);
-
-  printf ("hipMemcpyHostToDevice time taken  = %6.3fms\n", eventMs);
-
-  // Record the start event
-  hipEventRecord(start, NULL);
-
-  // Lauching kernel from host
-  hipLaunchKernel(matrixTranspose,
-                  dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
-                  dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
-                  0, 0,
-                  gpuTransposeMatrix , gpuMatrix, WIDTH);
-
-  // Record the stop event
-  hipEventRecord(stop, NULL);
-  hipEventSynchronize(stop);
-
-  hipEventElapsedTime(&eventMs, start, stop);
-
-  printf ("kernel Execution time             = %6.3fms\n", eventMs);
-
-  // Record the start event
-  hipEventRecord(start, NULL);
-
-  // Memory transfer from device to host
-  hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
-
-  // Record the stop event
-  hipEventRecord(stop, NULL);
-  hipEventSynchronize(stop);
-
-  hipEventElapsedTime(&eventMs, start, stop);
-
-  printf ("hipMemcpyDeviceToHost time taken  = %6.3fms\n", eventMs);
-
-  // CPU MatrixTranspose computation
-  matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
-
-  // verify the results
-  errors = 0;
-  double eps = 1.0E-6;
-  for (i = 0; i < NUM; i++) {
-    if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
-      errors++;
-    }
-  }
-  if (errors!=0) {
-    printf("FAILED: %d errors\n",errors);
-  } else {
-    printf ("PASSED!\n");
-  }
-
-  //free the resources on device side
-  hipFree(gpuMatrix);
-  hipFree(gpuTransposeMatrix);
-
-  //free the resources on host side
-  free(Matrix);
-  free(TransposeMatrix);
-  free(cpuTransposeMatrix);
-
-  return errors;
-}
@@ -0,0 +1,53 @@
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+
+HIPPROFILER=/opt/rocm/bin/rocm-profiler
+PROFILER_OPT=-A -o MT.atp -e HIP_PROFILE_API=1
+HIPPROFILER_POST_CMD=$(HIP_PATH)/bin/hipdemangleatp MT.atp
+
+TARGET=hcc
+
+SOURCES = MatrixTranspose.cpp
+OBJECTS = $(SOURCES:.cpp=.o)
+
+EXECUTABLE=./MatrixTranspose
+
+.PHONY: test
+
+
+all: $(EXECUTABLE) profile
+
+
+
+OPT =-g 
+CXXFLAGS =$(OPT) 
+CXX=$(HIPCC)
+
+
+$(EXECUTABLE): $(OBJECTS)
+	$(HIPCC) $(OBJECTS) -o $@
+
+
+profile: $(EXECUTABLE)
+	$(HIPPROFILER) $(PROFILER_OPT) $(EXECUTABLE)
+	$(HIPPROFILER_POST_CMD)
+
+
+# Pass option to control start and stop iterations for profiling - see MatrixTranspose.cpp for implementation:
+# Note we start profiler in --startdisabled mode - no timing collected until app enabled it via hipProfilerStart()
+profile_trigger: $(EXECUTABLE)
+	$(HIPPROFILER) $(PROFILER_OPT)  --startdisabled   $(EXECUTABLE) 3 6
+	$(HIPPROFILER_POST_CMD)
+
+
+run: $(EXECUTABLE)
+	$(EXECUTABLE)
+
+
+clean:
+	rm -f $(EXECUTABLE)
+	rm -f $(OBJECTS)
+	rm -f $(HIP_PATH)/src/*.o
+
@@ -0,0 +1,233 @@
+/*
+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<iostream>
+
+// hip header file
+#include "hip/hip_runtime.h"
+#include "hip/hip_profile.h"
+
+#define WIDTH     1024
+
+#define NUM       (WIDTH*WIDTH)
+
+#define THREADS_PER_BLOCK_X  4
+#define THREADS_PER_BLOCK_Y  4
+#define THREADS_PER_BLOCK_Z  1
+
+#define ITERATIONS 10
+
+// Cmdline parms to control start and stop triggers
+int startTriggerIteration=-1;
+int stopTriggerIteration=-1;
+
+// Device (Kernel) function, it must be void
+// hipLaunchParm provides the execution configuration
+__global__ void matrixTranspose(hipLaunchParm lp,
+                                float *out,
+                                float *in,
+                                const int width)
+{
+    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+    out[y * width + x] = in[x * width + y];
+}
+
+// CPU implementation of matrix transpose
+void matrixTransposeCPUReference(
+    float * output,
+    float * input,
+    const unsigned int width)
+{
+    for(unsigned int j=0; j < width; j++)
+    {
+        for(unsigned int i=0; i < width; i++)
+        {
+            output[i*width + j] = input[j*width + i];
+        }
+    }
+}
+
+
+// Use a separate function to demonstrate how to use function name as part of scoped marker:
+void runGPU(float *Matrix, float *TransposeMatrix, 
+            float* gpuMatrix, float* gpuTransposeMatrix)  {
+
+  // __func__ is a standard C++ macro which expands to the name of the function, in this case "runGPU"
+  HIP_SCOPED_MARKER(__func__, "MyGroup");
+
+  for (int i=0; i<ITERATIONS; i++) {
+
+    if (i==startTriggerIteration) {
+      hipProfilerStart();
+    }
+    if (i==stopTriggerIteration) {
+      hipProfilerStop();
+    }
+
+    float eventMs = 0.0f;
+
+    hipEvent_t start, stop;
+    hipEventCreate(&start);
+    hipEventCreate(&stop);
+
+
+    // Record the start event
+    hipEventRecord(start, NULL);
+
+    // Memory transfer from host to device
+    hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
+
+    // Record the stop event
+    hipEventRecord(stop, NULL);
+    hipEventSynchronize(stop);
+
+    hipEventElapsedTime(&eventMs, start, stop);
+
+    printf ("hipMemcpyHostToDevice time taken  = %6.3fms\n", eventMs);
+
+    // Record the start event
+    hipEventRecord(start, NULL);
+
+    // Lauching kernel from host
+    hipLaunchKernel(matrixTranspose,
+                    dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
+                    dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                    0, 0,
+                    gpuTransposeMatrix , gpuMatrix, WIDTH);
+
+    // Record the stop event
+    hipEventRecord(stop, NULL);
+    hipEventSynchronize(stop);
+    hipEventElapsedTime(&eventMs, start, stop);
+
+    printf ("kernel Execution time             = %6.3fms\n", eventMs);
+
+    // Record the start event
+    hipEventRecord(start, NULL);
+
+    // Memory transfer from device to host
+    hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
+
+    // Record the stop event
+    hipEventRecord(stop, NULL);
+    hipEventSynchronize(stop);
+
+    hipEventElapsedTime(&eventMs, start, stop);
+
+    printf ("hipMemcpyDeviceToHost time taken  = %6.3fms\n", eventMs);
+  }
+};
+
+
+int main(int argc, char *argv[]) {
+
+  if (argc >= 2) {
+    startTriggerIteration = atoi(argv[1]);
+    printf ("info : will start tracing at iteration:%d\n", startTriggerIteration);
+  } 
+  if (argc >= 3) {
+    stopTriggerIteration = atoi(argv[2]);
+    printf ("info : will stop tracing at iteration:%d\n", stopTriggerIteration);
+  }
+
+  float* Matrix;
+  float* TransposeMatrix;
+  float* cpuTransposeMatrix;
+
+  float* gpuMatrix;
+  float* gpuTransposeMatrix;
+
+  hipDeviceProp_t devProp;
+  hipGetDeviceProperties(&devProp, 0);
+
+  std::cout << "Device name " << devProp.name << std::endl;
+
+  {
+      // Show example of how to create a "scoped marker".  
+      // The scoped marker records the time spent inside the { scope } of the marker - the begin timestamp is at the
+      // beginning of the code scope, and the end is recorded when the SCOPE exits.  This can be viewed in CodeXL
+      // timeline relative to other GPU and CPU events.
+      // This marker captures the time spent in setup including host allocation, initialization, and device memory allocation.
+      HIP_SCOPED_MARKER("Setup", "MyGroup");
+
+
+
+      Matrix = (float*)malloc(NUM * sizeof(float));
+      TransposeMatrix = (float*)malloc(NUM * sizeof(float));
+      cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
+
+      // initialize the input data
+      for (int i = 0; i < NUM; i++) {
+        Matrix[i] = (float)i*10.0f;
+      }
+
+
+      // allocate the memory on the device side
+      hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
+      hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
+
+      // FYI, the scoped-marker will be destroyed here when the scope exits, and will record its "end" timestamp.
+  }
+
+  runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix);
+
+
+  // show how to use explicit begin/end markers:
+  // We begin the timed region with HIP_BEGIN_MARKER, passing in the markerName and group:
+  // The region will stop when HIP_END_MARKER is called
+  // This is another way to mark begin/end - as an alternative to scoped markers.
+  HIP_BEGIN_MARKER("Check&TearDown", "MyGroup");
+
+  int errors = 0;
+
+  // CPU MatrixTranspose computation
+  matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
+
+  // verify the results
+  double eps = 1.0E-6;
+  for (int i = 0; i < NUM; i++) {
+    if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
+      errors++;
+    }
+  }
+  if (errors!=0) {
+    printf("FAILED: %d errors\n",errors);
+  } else {
+    printf ("PASSED!\n");
+  }
+
+  //free the resources on device side
+  hipFree(gpuMatrix);
+  hipFree(gpuTransposeMatrix);
+
+  //free the resources on host side
+  free(Matrix);
+  free(TransposeMatrix);
+  free(cpuTransposeMatrix);
+
+  // This ends the last marker started in this thread, in this case "Check&TearDown"
+  HIP_END_MARKER();  
+  
+  return errors;
+}
@@ -1,6 +1,6 @@
 ## Using hipEvents to measure performance ###

-This tutorial is follow-up of the previous two tutorial where we learn how to write our first hip program, in which we compute Matrix Transpose and in second one, we added feature to measure time taken for memory transfer and kernel execution. In this tutorial, we won't make amy changes to the source code. We'll explain how to use the  codexl/rocm-profiler for hip timeline tracing.
+This tutorial is follow-up of the previous two tutorial where we learn how to write our first hip program, in which we compute Matrix Transpose and in second one, we added feature to measure time taken for memory transfer and kernel execution. In this tutorial, we'll explain how to use the codexl/rocm-profiler for hip timeline tracing.  Also, we will augment the source code with additional markers so we can see the high-level application flow alongside the information that CodeXL automatically collects.


 ## Introduction:
@@ -24,15 +24,11 @@ HIP can generate markers at function being/end which are displayed on the CodeXL

 1. Install ROCm-Profiler Installing HIP from the rocm pre-built packages, installs the ROCm-Profiler as well. Alternatively, you can build ROCm-Profiler using the instructions given below.

-2. Build HIP with ATP markers enabled HIP pre-built packages are enabled with ATP marker support by default. To enable ATP marker support when building HIP from source, use the option -DCOMPILE_HIP_ATP_MARKER=1 during the cmake configure step.

-3. Set HIP_ATP_MARKER
-`export HIP_ATP_MARKER=1`
-
-4. Recompile the target application
-
-5. Run with profiler enabled to generate ATP file.
-`/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>`
+2. Run with profiler enabled to generate ATP file.
+(These steps are also captured in the Makefile)
+The HIP_PROFILE_API enables display of the HIP APIs on the CodeXL trimeline view.
+`/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> -e HIP_PROFILE_API=1 <applicationArguments>`

 ##Using HIP_TRACE_API

@@ -1,36 +1,36 @@
-HIP_PATH?= $(wildcard /opt/rocm/hip)
-ifeq (,$(HIP_PATH))
-	HIP_PATH=../../..
-endif
-
-HIPCC=$(HIP_PATH)/bin/hipcc
-
-TARGET=hcc
-
-SOURCES = MatrixTranspose.cpp
-OBJECTS = $(SOURCES:.cpp=.o)
-
-EXECUTABLE=./MatrixTranspose
-
-.PHONY: test
-
-
-all: $(EXECUTABLE) test
-
-CXXFLAGS =-g
-CXX=$(HIPCC)
-
-
-$(EXECUTABLE): $(OBJECTS)
-	$(HIPCC) $(OBJECTS) -o $@
-
-
-test: $(EXECUTABLE)
-	$(EXECUTABLE)
-
-
-clean:
-	rm -f $(EXECUTABLE)
-	rm -f $(OBJECTS)
-	rm -f $(HIP_PATH)/src/*.o
-
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+	HIP_PATH=../../..
+endif
+
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+TARGET=hcc
+
+SOURCES = peer2peer.cpp
+OBJECTS = $(SOURCES:.cpp=.o)
+
+EXECUTABLE=./peer2peer
+
+.PHONY: test
+
+
+all: $(EXECUTABLE) test
+
+CXXFLAGS =-g
+CXX=$(HIPCC)
+
+
+$(EXECUTABLE): $(OBJECTS)
+	$(HIPCC) $(OBJECTS) -o $@
+
+
+test: $(EXECUTABLE)
+	$(EXECUTABLE)
+
+
+clean:
+	rm -f $(EXECUTABLE)
+	rm -f $(OBJECTS)
+	rm -f $(HIP_PATH)/src/*.o
+
@@ -0,0 +1,241 @@
+/*
+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 WARRANUMTY OF ANY KIND, EXPRESS OR
+IMPLIED, INUMCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNUMESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANUMY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INUM AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INUM CONUMECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <iostream>
+#include <hip/hip_runtime.h>
+#include <assert.h>
+#define WIDTH     32
+
+#define NUM       (WIDTH*WIDTH)
+
+#define THREADS_PER_BLOCK_X  4
+#define THREADS_PER_BLOCK_Y  4
+#define THREADS_PER_BLOCK_Z  1
+
+using namespace std;
+
+#define KNRM  "\x1B[0m"
+#define KRED  "\x1B[31m"
+
+#define failed(...) \
+    printf ("%serror: ", KRED);\
+    printf (__VA_ARGS__);\
+    printf ("\n");\
+    printf ("error: TEST FAILED\n%s", KNRM );\
+    abort();
+
+#define HIPCHECK(error) \
+{\
+    hipError_t localError = error; \
+    if (localError != hipSuccess) { \
+        printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \
+        KRED, hipGetErrorString(localError), localError,\
+        #error,__FILE__, __LINE__, KNRM); \
+        failed("API returned error code.");\
+    }\
+}
+
+void checkPeer2PeerSupport()
+{
+    int gpuCount;
+    int canAccessPeer;
+    int p2pCapableDeviceCount=0;
+
+    HIPCHECK(hipGetDeviceCount(&gpuCount));
+
+    if (gpuCount < 2)
+        printf("Peer2Peer application requires atleast 2 gpu devices");
+
+    for (int currentGpu=0; currentGpu<gpuCount; currentGpu++)
+    {
+        HIPCHECK(hipSetDevice(currentGpu));
+
+        for (int peerGpu=0; peerGpu<currentGpu; peerGpu++)
+        {
+            if (currentGpu!=peerGpu)
+            {
+                HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu));
+                printf ("currentGpu#%d canAccessPeer: peerGpu#%d=%d\n", currentGpu, peerGpu, canAccessPeer);
+            }
+
+            HIPCHECK(hipSetDevice(peerGpu));
+            HIPCHECK(hipDeviceReset());
+        }
+    HIPCHECK(hipSetDevice(currentGpu));
+    HIPCHECK(hipDeviceReset());
+    }
+}
+
+void enablePeer2Peer(int currentGpu, int peerGpu)
+{
+    int canAccessPeer;
+
+    // Must be on a multi-gpu system:
+    assert (currentGpu != peerGpu);
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
+
+    if(canAccessPeer==1){
+        HIPCHECK(hipDeviceEnablePeerAccess(peerGpu, 0));
+    }
+    else
+    printf("peer2peer transfer not possible between the selected gpu devices");
+}
+
+void disablePeer2Peer(int currentGpu, int peerGpu)
+{
+    int canAccessPeer;
+
+    // Must be on a multi-gpu system:
+    assert (currentGpu != peerGpu);
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
+	
+    if(canAccessPeer==1){
+        HIPCHECK(hipDeviceDisablePeerAccess(peerGpu));
+    }
+    else
+    printf("peer2peer disable not required");
+}
+
+
+__global__ void matrixTranspose_static_shared(hipLaunchParm lp,
+                                float *out,
+                                float *in,
+                                const int width)
+{
+    __shared__ float sharedMem[WIDTH*WIDTH];
+
+    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+    sharedMem[y * width + x] = in[x * width + y];
+
+    __syncthreads();
+
+    out[y * width + x] = sharedMem[y * width + x];
+}
+
+__global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp,
+                                float *out,
+                                float *in,
+                                const int width)
+{
+    // declare dynamic shared memory
+    HIP_DYNAMIC_SHARED(float, sharedMem)
+
+    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+    sharedMem[y * width + x] = in[x * width + y];
+
+    __syncthreads();
+
+    out[y * width + x] = sharedMem[y * width + x];
+}
+
+int main(){
+
+    checkPeer2PeerSupport();
+    
+    int gpuCount;
+    int currentGpu, peerGpu;
+
+    HIPCHECK(hipGetDeviceCount(&gpuCount));
+
+    currentGpu = 0;
+    peerGpu = (currentGpu + 1);
+
+    printf ("currentGpu=%d peerGpu=%d (Total no. of gpu = %d)\n", currentGpu, peerGpu, gpuCount);
+
+    float *data[2], *TransposeMatrix[2], *gpuTransposeMatrix[2], *randArray;
+
+    int width = WIDTH;
+
+    randArray = (float*)malloc(NUM * sizeof(float));
+
+    for(int i = 0; i < NUM; i++)
+    {
+        randArray[i] = (float)i*1.0f;
+    }
+
+    enablePeer2Peer(currentGpu, peerGpu);
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    TransposeMatrix[0] = (float*)malloc(NUM * sizeof(float));
+    hipMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float));
+    hipMalloc((void**)&data[0], NUM * sizeof(float));
+    hipMemcpy(data[0], randArray, NUM * sizeof(float), hipMemcpyHostToDevice);
+
+    hipLaunchKernel(matrixTranspose_static_shared,
+                    dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
+                    dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                    0, 0,
+                    gpuTransposeMatrix[0], data[0], width);
+
+	HIPCHECK(hipSetDevice(peerGpu));
+	TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));
+	hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
+	hipMalloc((void**)&data[1], NUM * sizeof(float));
+    hipMemcpy(data[1], gpuTransposeMatrix[0], NUM * sizeof(float), hipMemcpyDeviceToDevice);
+
+    hipLaunchKernel(matrixTranspose_dynamic_shared,
+                    dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
+                    dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                    sizeof(float)*WIDTH*WIDTH, 0,
+                    gpuTransposeMatrix[1], data[1], width);
+
+    hipMemcpy(TransposeMatrix[1], gpuTransposeMatrix[1], NUM*sizeof(float), hipMemcpyDeviceToHost);
+
+    hipDeviceSynchronize();
+
+    disablePeer2Peer(currentGpu, peerGpu);
+
+    // verify the results
+    int errors = 0;
+    double eps = 1.0E-6;
+    for (int i = 0; i < NUM; i++) {
+        if (std::abs(randArray[i] - TransposeMatrix[1][i]) > eps ) {
+        printf("%d cpu: %f gpu peered data  %f\n",i,randArray[i],TransposeMatrix[1][i]);
+        errors++;
+        }
+    }
+    if (errors!=0) {
+        printf("FAILED: %d errors\n",errors);
+    } else {
+        printf ("Peer2Peer PASSED!\n");
+    }
+
+    free(randArray);
+    for(int i=0;i<2;i++){
+       hipFree(data[i]);
+       hipFree(gpuTransposeMatrix[i]);
+       free(TransposeMatrix[i]);
+    }
+
+    HIPCHECK(hipSetDevice(peerGpu));
+    HIPCHECK(hipDeviceReset());
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    HIPCHECK(hipDeviceReset());
+
+    return 0;
+}
@@ -340,8 +340,28 @@ hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
    return ihipLogStatus(e);
 }

+hipError_t hipDeviceGetByPCIBusId (int*  device, const int* pciBusId )
+{
+    HIP_INIT_API(device,pciBusId);
+    hipDeviceProp_t  tempProp;
+    int deviceCount;
+    hipError_t e = hipErrorInvalidValue;
+    hipGetDeviceCount( &deviceCount );
+    *device = 0;
+    for (int i=0; i< deviceCount; i++) {
+        hipGetDeviceProperties( &tempProp, i );
+        if(tempProp.pciBusID == *pciBusId) {
+            *device =i;
+            e = hipSuccess;
+            break;
+        }
+    }
+    return ihipLogStatus(e);
+}
+
 hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop )
 {
+    HIP_INIT_API(device,prop);
    hipDeviceProp_t  tempProp;
    int deviceCount;
    int inPropCount=0;
@@ -48,6 +48,9 @@ THE SOFTWARE.
 #include "trace_helper.h"


+#ifndef USE_COPY_EXT_V2
+#define USE_COPY_EXT_V2 0
+#endif

 //=================================================================================================
 //Global variables:
@@ -62,12 +65,27 @@ int HIP_LAUNCH_BLOCKING = 0;
 int HIP_PRINT_ENV = 0;
 int HIP_TRACE_API= 0;
 std::string HIP_TRACE_API_COLOR("green");
-int HIP_ATP_MARKER= 0;
+int HIP_PROFILE_API= 0;
+
+// TODO - DB_START/STOP need more testing.
+std::string HIP_DB_START_API;
+std::string HIP_DB_STOP_API;
 int HIP_DB= 0;
 int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */
 int HIP_NUM_KERNELS_INFLIGHT = 128;
 int HIP_WAIT_MODE = 0;

+int HIP_FORCE_P2P_HOST = 0;
+int HIP_DENY_PEER_ACCESS = 0;
+
+// Force async copies to actually use the synchronous copy interface. 
+int HIP_FORCE_SYNC_COPY = 0;
+
+
+
+
+
+
 #define HIP_USE_PRODUCT_NAME 0
 //#define DISABLE_COPY_EXT 1

@@ -84,6 +102,15 @@ std::vector<int> g_hip_visible_devices;
 hsa_agent_t g_cpu_agent;
 unsigned g_numLogicalThreads;

+std::atomic<int> g_lastShortTid(1);
+
+// Indexed by short-tid:
+// 
+std::vector<ProfTrigger> g_dbStartTriggers;
+std::vector<ProfTrigger> g_dbStopTriggers;
+
+
+
 /*
 Implementation of malloc and free device functions.

@@ -163,22 +190,79 @@ __device__ void* __hip_hc_free(void *ptr)
    return nullptr;
 }

+__device__ unsigned __hip_ds_bpermute(int index, unsigned src) {
+    return hc::__amdgcn_ds_bpermute(index, src);
+}

+__device__ float __hip_ds_bpermutef(int index, float src) {
+    return hc::__amdgcn_ds_bpermute(index, src);
+}
+
+__device__ unsigned __hip_ds_permute(int index, unsigned src) {
+    return hc::__amdgcn_ds_permute(index, src);
+}
+
+__device__ float __hip_ds_permutef(int index, float src) {
+    return hc::__amdgcn_ds_permute(index, src);
+}
+
+__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) {
+    return hc::__amdgcn_ds_swizzle(src, pattern);
+}
+
+__device__ float __hip_ds_swizzlef(float src, int pattern) {
+    return hc::__amdgcn_ds_swizzle(src, pattern);
+}
+
+__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) {
+    return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl);
+}
 //=================================================================================================
 // Thread-local storage:
 //=================================================================================================

 // This is the implicit context used by all HIP commands.
 // It can be set by hipSetDevice or by the CTX manipulation commands:
-
 thread_local hipError_t tls_lastHipError = hipSuccess;


+thread_local ShortTid tls_shortTid;
+
+


 //=================================================================================================
 // Top-level "free" functions:
 //=================================================================================================
+void recordApiTrace(std::string *fullStr, const std::string &apiStr)
+{
+    auto apiSeqNum = tls_shortTid.incApiSeqNum();
+    auto tid = tls_shortTid.tid();
+
+    if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) {
+        printf ("info: resume profiling at %lu\n", apiSeqNum);
+        RESUME_PROFILING; 
+        g_dbStartTriggers.pop_back();
+    };
+    if ((tid < g_dbStopTriggers.size()) && (apiSeqNum >= g_dbStopTriggers[tid].nextTrigger())) {
+        printf ("info: stop profiling at %lu\n", apiSeqNum);
+        STOP_PROFILING; 
+        g_dbStopTriggers.pop_back();
+    };
+
+    fullStr->reserve(16 + apiStr.length());
+    *fullStr = std::to_string(tid) + ".";
+    *fullStr += std::to_string(apiSeqNum);
+    *fullStr += " ";
+    *fullStr += apiStr;
+
+
+    if (COMPILE_HIP_DB && HIP_TRACE_API) {
+        fprintf (stderr, "%s<<hip-api tid:%s%s\n" , API_COLOR, fullStr->c_str(), API_COLOR_END);
+    }
+}
+
+
 static inline bool ihipIsValidDevice(unsigned deviceIndex)
 {
    // deviceIndex is unsigned so always > 0
@@ -230,6 +314,23 @@ hipError_t ihipSynchronize(void)
    return (hipSuccess);
 }

+//=================================================================================================
+// ihipStream_t:
+//=================================================================================================
+ShortTid::ShortTid()  :
+    _apiSeqNum(0)
+{ 
+    _shortTid = g_lastShortTid.fetch_add(1); 
+
+    if (HIP_DB & (1<<DB_API)) {
+        std::stringstream tid_ss;
+        std::stringstream tid_ss_num;
+        tid_ss_num << std::this_thread::get_id();
+        tid_ss << std::hex << std::stoull(tid_ss_num.str());
+
+        tprintf(DB_API, "HIP initialized short_tid#%d (maps to full_tid: 0x%s)\n", _shortTid, tid_ss.str().c_str());
+    };
+}

 //=================================================================================================
 // ihipStream_t:
@@ -252,8 +353,6 @@ ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int
    };


-
-
    tprintf(DB_SYNC, " streamCreate: stream=%p\n", this);
 };

@@ -426,7 +525,7 @@ void ihipStream_t::launchModuleKernel(
    (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
    (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);

-    uint16_t setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
+    uint16_t setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
    uint32_t header32 = header | (setup << 16);

    __atomic_store_n((uint32_t*)(dispatch_packet), header32, __ATOMIC_RELEASE);
@@ -451,7 +550,7 @@ void ihipCtxCriticalBase_t<CtxMutex>::recomputePeerAgents()


 template<>
-bool ihipCtxCriticalBase_t<CtxMutex>::isPeer(const ihipCtx_t *peer)
+bool ihipCtxCriticalBase_t<CtxMutex>::isPeerWatcher(const ihipCtx_t *peer)
 {
    auto match = std::find(_peers.begin(), _peers.end(), peer);
    return (match != std::end(_peers));
@@ -459,12 +558,14 @@ bool ihipCtxCriticalBase_t<CtxMutex>::isPeer(const ihipCtx_t *peer)


 template<>
-bool ihipCtxCriticalBase_t<CtxMutex>::addPeer(ihipCtx_t *peer)
+bool ihipCtxCriticalBase_t<CtxMutex>::addPeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peerWatcher)
 {
-    auto match = std::find(_peers.begin(), _peers.end(), peer);
+    auto match = std::find(_peers.begin(), _peers.end(), peerWatcher);
    if (match == std::end(_peers)) {
        // Not already a peer, let's update the list:
-        _peers.push_back(peer);
+        tprintf(DB_COPY, "addPeerWatcher.  Allocations on %s now visible to peerWatcher %s.\n", 
+                          thisCtx->toString().c_str(), peerWatcher->toString().c_str());
+        _peers.push_back(peerWatcher);
        recomputePeerAgents();
        return true;
    }
@@ -475,12 +576,14 @@ bool ihipCtxCriticalBase_t<CtxMutex>::addPeer(ihipCtx_t *peer)


 template<>
-bool ihipCtxCriticalBase_t<CtxMutex>::removePeer(ihipCtx_t *peer)
+bool ihipCtxCriticalBase_t<CtxMutex>::removePeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peerWatcher)
 {
-    auto match = std::find(_peers.begin(), _peers.end(), peer);
+    auto match = std::find(_peers.begin(), _peers.end(), peerWatcher);
    if (match != std::end(_peers)) {
        // Found a valid peer, let's remove it.
-        _peers.remove(peer);
+        tprintf(DB_COPY, "removePeerWatcher.  Allocations on %s no longer visible to former peerWatcher %s.\n", 
+                          thisCtx->toString().c_str(), peerWatcher->toString().c_str());
+        _peers.remove(peerWatcher);
        recomputePeerAgents();
        return true;
    } else {
@@ -490,16 +593,17 @@ bool ihipCtxCriticalBase_t<CtxMutex>::removePeer(ihipCtx_t *peer)


 template<>
-void ihipCtxCriticalBase_t<CtxMutex>::resetPeers(ihipCtx_t *thisDevice)
+void ihipCtxCriticalBase_t<CtxMutex>::resetPeerWatchers(ihipCtx_t *thisCtx)
 {
+    tprintf(DB_COPY, "resetPeerWatchers for context=%s\n", thisCtx->toString().c_str());
    _peers.clear();
    _peerCnt = 0;
-    addPeer(thisDevice); // peer-list always contains self agent.
+    addPeerWatcher(thisCtx, thisCtx); // peer-list always contains self agent.
 }


 template<>
-void ihipCtxCriticalBase_t<CtxMutex>::printPeers(FILE *f) const
+void ihipCtxCriticalBase_t<CtxMutex>::printPeerWatchers(FILE *f) const
 {
    for (auto iter = _peers.begin(); iter!=_peers.end(); iter++) {
        fprintf (f, "%s ", (*iter)->toString().c_str());
@@ -904,7 +1008,7 @@ void ihipCtx_t::locked_reset()


    // Reset peer list to just me:
-    crit->resetPeers(this);
+    crit->resetPeerWatchers(this);

    // Reset and release all memory stored in the tracker:
    // Reset will remove peer mapping so don't need to do this explicitly.
@@ -1006,43 +1110,14 @@ void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, c
        env = getenv(var_name2);
    }

-    // TODO: Refactor this code so it is a separate call rather than being part of ihipReadEnv_I, which should only read integers.
-    // Check if the environment variable is either HIP_VISIBLE_DEVICES or CUDA_LAUNCH_BLOCKING, which
-    // contains a sequence of comma-separated device IDs
-    if (!(strcmp(var_name1,"HIP_VISIBLE_DEVICES") && strcmp(var_name2, "CUDA_VISIBLE_DEVICES")) && env){
-        // Parse the string stream of env and store the device ids to g_hip_visible_devices global variable
-        std::string str = env;
-        std::istringstream ss(str);
-        std::string device_id;
-        // Clean up the defult value
-        g_hip_visible_devices.clear();
-        g_visible_device = true;
-        // Read the visible device numbers
-        while (std::getline(ss, device_id, ',')) {
-            if (atoi(device_id.c_str()) >= 0) {
-                g_hip_visible_devices.push_back(atoi(device_id.c_str()));
-            } else { // Any device number after invalid number will not present
-                break;
-            }
-        }
-        // Print out the number of ids
-        if (HIP_PRINT_ENV) {
-            printf ("%-30s = ", var_name1);
-            for(int i=0;i<g_hip_visible_devices.size();i++)
-                printf ("%2d ", g_hip_visible_devices[i]);
-            printf (": %s\n", description);
-        }
+    // Default is set when variable is initialized (at top of this file), so only override if we find
+    // an environment variable.
+    if (env) {
+        long int v = strtol(env, NULL, 0);
+        *var_ptr = (int) (v);
    }
-    else { // Parse environment variables with sigle value
-        // Default is set when variable is initialized (at top of this file), so only override if we find
-        // an environment variable.
-        if (env) {
-            long int v = strtol(env, NULL, 0);
-            *var_ptr = (int) (v);
-        }
-        if (HIP_PRINT_ENV) {
-            printf ("%-30s = %2d : %s\n", var_name1, *var_ptr, description);
-        }
+    if (HIP_PRINT_ENV) {
+        printf ("%-30s = %2d : %s\n", var_name1, *var_ptr, description);
    }
 }

@@ -1057,7 +1132,7 @@ void ihipReadEnv_S(std::string *var_ptr, const char *var_name1, const char *var_
    }

    if (env) {
-        *var_ptr = env;
+        *static_cast<std::string*>(var_ptr) = env;
    }
    if (HIP_PRINT_ENV) {
        printf ("%-30s = %s : %s\n", var_name1, var_ptr->c_str(), description);
@@ -1065,6 +1140,25 @@ void ihipReadEnv_S(std::string *var_ptr, const char *var_name1, const char *var_
 }


+void ihipReadEnv_Callback(void *var_ptr, const char *var_name1, const char *var_name2, const char *description, std::string (*setterCallback)(void * var_ptr, const char * env))
+{
+    char * env = getenv(var_name1);
+
+    // Check second name if first not defined, used to allow HIP_ or CUDA_ env vars.
+    if ((env == NULL) && strcmp(var_name2, "0")) {
+        env = getenv(var_name2);
+    }
+
+    std::string var_string = "0";
+    if (env) {
+        var_string = setterCallback(var_ptr, env);
+    }
+    if (HIP_PRINT_ENV) {
+        printf ("%-30s = %s : %s\n", var_name1, var_string.c_str(), description);
+    }
+}
+
+
 #if defined (DEBUG)

 #define READ_ENV_I(_build, _ENV_VAR, _ENV_VAR2, _description) \
@@ -1075,6 +1169,10 @@ void ihipReadEnv_S(std::string *var_ptr, const char *var_name1, const char *var_
    if ((_build == release) || (_build == debug) {\
        ihipReadEnv_S(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description);\
    };
+#define READ_ENV_C(_build, _ENV_VAR, _ENV_VAR2, _description, _callback) \
+    if ((_build == release) || (_build == debug) {\
+        ihipReadEnv_Callback(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description, _callback);\
+    };

 #else

@@ -1087,10 +1185,155 @@ void ihipReadEnv_S(std::string *var_ptr, const char *var_name1, const char *var_
    if (_build == release) {\
        ihipReadEnv_S(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description);\
    };
+#define READ_ENV_C(_build, _ENV_VAR, _ENV_VAR2, _description, _callback) \
+    if (_build == release) {\
+        ihipReadEnv_Callback(&_ENV_VAR, #_ENV_VAR, #_ENV_VAR2, _description, _callback);\
+    };

 #endif


+static void tokenize(const std::string &s, char delim, std::vector<std::string> *tokens)
+{
+    std::stringstream ss;
+    ss.str(s);
+    std::string item;
+    while (getline(ss, item, delim)) {
+        item.erase (std::remove (item.begin(), item.end(), ' '), item.end()); // remove whitespace.
+        tokens->push_back(item);
+    }
+}
+
+static void trim(std::string *s)
+{
+    // trim whitespace from beginning and end:
+    const char *t =  "\t\n\r\f\v";
+    s->erase(0, s->find_first_not_of(t));
+    s->erase(s->find_last_not_of(t)+1);
+}
+
+static void ltrim(std::string *s)
+{
+    // trim whitespace from beginning
+    const char *t =  "\t\n\r\f\v";
+    s->erase(0, s->find_first_not_of(t));
+}
+
+
+// TODO - change last arg to pointer.
+void parseTrigger(std::string triggerString, std::vector<ProfTrigger> &profTriggers )
+{
+    std::vector<std::string> tidApiTokens;
+    tokenize(std::string(triggerString), ',', &tidApiTokens);
+    for (auto t=tidApiTokens.begin(); t != tidApiTokens.end(); t++) {
+        std::vector<std::string> oneToken;
+        //std::cout << "token=" << *t << "\n";
+        tokenize(std::string(*t), '.', &oneToken);
+        int tid = 1;
+        uint64_t apiTrigger = 0;
+        if (oneToken.size() == 1) {
+            // the case with just apiNum
+            apiTrigger = std::strtoull(oneToken[0].c_str(), nullptr, 0);
+        } else if (oneToken.size() == 2) {
+            // the case with tid.apiNum
+            tid = std::strtoul(oneToken[0].c_str(), nullptr, 0);
+            apiTrigger = std::strtoull(oneToken[1].c_str(), nullptr, 0);
+        } else {
+            throw ihipException(hipErrorRuntimeOther); // TODO -> bad env var?
+        }
+
+        if (tid > 10000) {
+            throw ihipException(hipErrorRuntimeOther); // TODO -> bad env var?
+        } else {
+            profTriggers.resize(tid+1);
+            //std::cout << "tid:" << tid << " add: " << apiTrigger << "\n";
+            profTriggers[tid].add(apiTrigger);
+        }
+    }
+
+
+    for (int tid=1; tid<profTriggers.size(); tid++) {
+        profTriggers[tid].sort();
+        profTriggers[tid].print(tid);
+    }
+}
+
+std::string HIP_DB_string(unsigned db)
+{
+    std::string dbStr;
+    bool first=true;
+    for (int i=0; i<DB_MAX_FLAG; i++) {
+        if (db & (1<<i)) {
+            if (!first) {
+                dbStr += "+";
+            };
+            dbStr += dbName[i]._color;
+            dbStr += dbName[i]._shortName;
+            dbStr += KNRM;
+            first=false;
+        };
+    }
+
+    return dbStr;
+}
+
+// Callback used to process HIP_DB input, supports either
+// integer or flag names separated by +
+std::string HIP_DB_callback(void *var_ptr, const char *envVarString)
+{
+    int * var_ptr_int = static_cast<int *> (var_ptr);
+
+    std::string e(envVarString);
+    trim(&e);
+    if (!e.empty() && isdigit(e.c_str()[0])) {
+        long int v = strtol(envVarString, NULL, 0);
+        *var_ptr_int = (int) (v);
+    } else {
+        *var_ptr_int = 0;
+        std::vector<std::string> tokens;
+        tokenize(e, '+', &tokens);
+        for (auto t=tokens.begin(); t!= tokens.end(); t++) {
+            for (int i=0; i<DB_MAX_FLAG; i++) {
+                if (!strcmp(t->c_str(), dbName[i]._shortName)) { 
+                    *var_ptr_int |= (1<<i);
+                } // TODO - else throw error?
+            }
+        }
+    }
+
+    return HIP_DB_string(*var_ptr_int);;
+}
+
+
+// Callback used to process list of visible devices.
+std::string HIP_VISIBLE_DEVICES_callback(void *var_ptr, const char *envVarString)
+{
+    // Parse the string stream of env and store the device ids to g_hip_visible_devices global variable
+    std::string str = envVarString;
+    std::istringstream ss(str);
+    std::string device_id;
+    // Clean up the defult value
+    g_hip_visible_devices.clear();
+    g_visible_device = true;
+    // Read the visible device numbers
+    while (std::getline(ss, device_id, ',')) {
+        if (atoi(device_id.c_str()) >= 0) {
+            g_hip_visible_devices.push_back(atoi(device_id.c_str()));
+        } else { // Any device number after invalid number will not present
+            break;
+        }
+    }
+
+    std::string valueString;
+    // Print out the number of ids
+    for(int i=0;i<g_hip_visible_devices.size();i++) {
+        valueString += std::to_string((g_hip_visible_devices[i]));
+        valueString += ' ';
+    }
+
+    return valueString;
+}
+

 //---
 //Function called one-time at initialization time to construct a table of all GPU devices.
@@ -1114,7 +1357,7 @@ void ihipInit()

    // TODO: In HIP/hcc, this variable blocks after both kernel commmands and data transfer.  Maybe should be bit-mask for each command type?
    READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING." );
-    READ_ENV_I(release, HIP_DB, 0,  "Print various debug info.  Bitmask, see hip_hcc.cpp for more information.");
+    READ_ENV_C(release, HIP_DB, 0,  "Print debug info.  Bitmask (HIP_DB=0xff) or flags separated by '+' (HIP_DB=api+sync+mem+copy)", HIP_DB_callback);
    if ((HIP_DB & (1<<DB_API))  && (HIP_TRACE_API == 0)) {
        // Set HIP_TRACE_API default before we read it, so it is printed correctly.
        HIP_TRACE_API = 1;
@@ -1124,41 +1367,34 @@ void ihipInit()

    READ_ENV_I(release, HIP_TRACE_API, 0,  "Trace each HIP API call.  Print function name and return code to stderr as program executes.");
    READ_ENV_S(release, HIP_TRACE_API_COLOR, 0,  "Color to use for HIP_API.  None/Red/Green/Yellow/Blue/Magenta/Cyan/White");
-    READ_ENV_I(release, HIP_ATP_MARKER, 0,  "Add HIP function begin/end to ATP file generated with CodeXL");
-    READ_ENV_I(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence" );
+    READ_ENV_I(release, HIP_PROFILE_API, 0,  "Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, 0x2=full API name including args.");
+    READ_ENV_S(release, HIP_DB_START_API, 0,  "Comma-separted list of tid.api_seq_num for when to start debug and profiling.");
+    READ_ENV_S(release, HIP_DB_STOP_API, 0,  "Comma-separated list of tid.api_seq_num for when to stop debug and profiling.");
+    
+    READ_ENV_C(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the sequence are visible to HIP applications and they are enumerated in the order of sequence.", HIP_VISIBLE_DEVICES_callback );


    READ_ENV_I(release, HIP_WAIT_MODE, 0, "Force synchronization mode. 1= force yield, 2=force spin, 0=defaults specified in application");
-
+    READ_ENV_I(release, HIP_FORCE_P2P_HOST, 0, "Force use of host/staging copy for peer-to-peer copies.1=always use copies, 2=always return false for hipDeviceCanAccessPeer"); 
+    READ_ENV_I(release, HIP_FORCE_SYNC_COPY, 0, "Force all copies (even hipMemcpyAsync) to use sync copies"); 
    READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Max number of inflight kernels per stream before active synchronization is forced.");

    // Some flags have both compile-time and runtime flags - generate a warning if user enables the runtime flag but the compile-time flag is disabled.
    if (HIP_DB && !COMPILE_HIP_DB) {
-        fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0.  (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB);
+        fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0.  (perhaps enable COMPILE_HIP_DB in src code before compiling?)\n", HIP_DB);
    }

    if (HIP_TRACE_API && !COMPILE_HIP_TRACE_API) {
-        fprintf (stderr, "warning: env var HIP_TRACE_API=0x%x but COMPILE_HIP_TRACE_API=0.  (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_DB);
+        fprintf (stderr, "warning: env var HIP_TRACE_API=0x%x but COMPILE_HIP_TRACE_API=0.  (perhaps enable COMPILE_HIP_TRACE_API in src code before compiling?)\n", HIP_DB);
    }

-    if (HIP_ATP_MARKER && !COMPILE_HIP_ATP_MARKER) {
-        fprintf (stderr, "warning: env var HIP_ATP_MARKER=0x%x but COMPILE_HIP_ATP_MARKER=0.  (perhaps enable COMPILE_HIP_DB in src code before compiling?)", HIP_ATP_MARKER);
+    if (HIP_PROFILE_API && !COMPILE_HIP_ATP_MARKER) {
+        fprintf (stderr, "warning: env var HIP_PROFILE_API=0x%x but COMPILE_HIP_ATP_MARKER=0.  (perhaps enable COMPILE_HIP_ATP_MARKER in src code before compiling?)\n", HIP_PROFILE_API);
+        HIP_PROFILE_API = 0;
    }

    if (HIP_DB) {
-        fprintf (stderr, "HIP_DB=0x%x [", HIP_DB);
-        bool first=true;
-        for (int i=0; i<DB_MAX_BITPOS; i++) {
-            if (HIP_DB & (1<<i)) {
-                if (first) {
-                    fprintf (stderr, "%s%s%s", dbName[i]._color, dbName[i]._shortName, KNRM);
-                } else {
-                    fprintf (stderr, "+%s%s%s", dbName[i]._color, dbName[i]._shortName, KNRM);
-                };
-                first=false;
-            };
-        }
-        fprintf (stderr, "]\n");
+        fprintf (stderr, "HIP_DB=0x%x [%s]\n", HIP_DB, HIP_DB_string(HIP_DB).c_str());
    }

    std::transform(HIP_TRACE_API_COLOR.begin(),  HIP_TRACE_API_COLOR.end(), HIP_TRACE_API_COLOR.begin(), ::tolower);
@@ -1184,6 +1420,9 @@ void ihipInit()
        fprintf (stderr, "warning: env var HIP_TRACE_API_COLOR=%s must be None/Red/Green/Yellow/Blue/Magenta/Cyan/White", HIP_TRACE_API_COLOR.c_str());
    };

+    parseTrigger(HIP_DB_START_API, g_dbStartTriggers);
+    parseTrigger(HIP_DB_STOP_API,  g_dbStopTriggers);
+



@@ -1269,19 +1508,28 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)

 void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
 {
-    if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {
+    if (HIP_PROFILE_API || (COMPILE_HIP_DB && HIP_TRACE_API)) {
+        std::stringstream os_pre;
        std::stringstream os;
-        os  << "<<hip-api: hipLaunchKernel '" << kernelName << "'"
+        os_pre  << "<<hip-api tid:";
+        os  << tls_shortTid.tid() << "." << tls_shortTid.incApiSeqNum()
+            << " hipLaunchKernel '" << kernelName << "'"
            << " gridDim:"  << lp->grid_dim
            << " groupDim:" << lp->group_dim
            << " sharedMem:+" << lp->dynamic_group_mem_bytes
            << " " << *stream;

+        if (HIP_PROFILE_API == 0x1) {
+            std::string shortAtpString("hipLaunchKernel:");
+            shortAtpString += kernelName;
+            MARKER_BEGIN(shortAtpString.c_str(), "HIP");
+        } else if (HIP_PROFILE_API == 0x2) {
+            MARKER_BEGIN(os.str().c_str(), "HIP");
+        }

        if (COMPILE_HIP_DB && HIP_TRACE_API) {
            std::cerr << API_COLOR << os.str() << API_COLOR_END << std::endl;
        }
-        SCOPED_MARKER(os.str().c_str(), "HIP", NULL);
    }
 }

@@ -1379,7 +1627,10 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
 //This releases the lock on the stream.
 void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp)
 {
+    tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
+
    stream->lockclose_postKernelCommand(lp.av);
+    MARKER_END();
 }


@@ -1486,29 +1737,44 @@ void ihipSetTs(hipEvent_t e)
 }


-// Returns true if thisCtx can see the memory allocated on dstCtx and srcCtx.
+// Returns true if copyEngineCtx can see the memory allocated on dstCtx and srcCtx.
 // The peer-list for a context controls which contexts have access to the memory allocated on that context.
 // So we check dstCtx's and srcCtx's peerList to see if the both include thisCtx.
-bool ihipStream_t::canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, ihipCtx_t *srcCtx)
+bool ihipStream_t::canSeeMemory(const ihipCtx_t *copyEngineCtx, const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo)
 {
-    tprintf (DB_COPY1, "Checking if direct copy can be used.  thisCtx:%s;  dstCtx:%s ;  srcCtx:%s\n",
-            thisCtx->toString().c_str(), dstCtx->toString().c_str(), srcCtx->toString().c_str());

-    // Use blocks to control scope of critical sections.
-    {
-        LockedAccessor_CtxCrit_t  ctxCrit(dstCtx->criticalData());
-        tprintf(DB_SYNC, "dstCrit lock succeeded\n");
-        if (!ctxCrit->isPeer(thisCtx)) {
-            return false;
-        };
+    // Make sure this is a device-to-device copy with all memory available to the requested copy engine
+    //
+    // TODO - pointer-info stores a deviceID not a context,may have some unusual side-effects here:
+    if (dstPtrInfo->_sizeBytes == 0) {
+        return false;
+    } else {
+        ihipCtx_t *dstCtx = ihipGetPrimaryCtx(dstPtrInfo->_appId);
+        if (copyEngineCtx != dstCtx) {
+            // Only checks peer list if contexts are different
+            LockedAccessor_CtxCrit_t  ctxCrit(dstCtx->criticalData());
+            //tprintf(DB_SYNC, "dstCrit lock succeeded\n");
+            if (!ctxCrit->isPeerWatcher(copyEngineCtx)) {
+                return false;
+            };
+        }
    }

-    {
-        LockedAccessor_CtxCrit_t  ctxCrit(srcCtx->criticalData());
-        tprintf(DB_SYNC, "srcCrit lock succeeded\n");
-        if (!ctxCrit->isPeer(thisCtx)) {
-            return false;
-        };
+
+
+    // TODO - pointer-info stores a deviceID not a context,may have some unusual side-effects here:
+    if (srcPtrInfo->_sizeBytes == 0) {
+        return false;
+    } else {
+        ihipCtx_t *srcCtx = ihipGetPrimaryCtx(srcPtrInfo->_appId);
+        if (copyEngineCtx != srcCtx) {
+            // Only checks peer list if contexts are different
+            LockedAccessor_CtxCrit_t  ctxCrit(srcCtx->criticalData());
+            //tprintf(DB_SYNC, "srcCrit lock succeeded\n");
+            if (!ctxCrit->isPeerWatcher(copyEngineCtx)) {
+                return false;
+            };
+        }
    }

    return true;
@@ -1517,7 +1783,7 @@ bool ihipStream_t::canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx,

 #define CASE_STRING(X)  case X: return #X ;break;

-const char* memcpyStr(unsigned memKind)
+const char* hipMemcpyStr(unsigned memKind)
 {
    switch (memKind) {
        CASE_STRING(hipMemcpyHostToHost);
@@ -1529,38 +1795,81 @@ const char* memcpyStr(unsigned memKind)
    };
 }

+const char* hcMemcpyStr(hc::hcCommandKind memKind)
+{
+    using namespace hc;
+    switch (memKind) {
+        CASE_STRING(hcMemcpyHostToHost);
+        CASE_STRING(hcMemcpyHostToDevice);
+        CASE_STRING(hcMemcpyDeviceToHost);
+        CASE_STRING(hcMemcpyDeviceToDevice);
+        //CASE_STRING(hcMemcpyDefault);
+        default : return ("unknown memcpyKind");
+    };
+}
+


 // Resolve hipMemcpyDefault to a known type.
-// TODO - review why is this so complicated, does this need srcTracked and dstTracked?
-unsigned ihipStream_t::resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem)
+unsigned ihipStream_t::resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem)
 {
    hipMemcpyKind kind = hipMemcpyDefault;
-    if(!srcTracked && !dstTracked)
-    {
-        kind = hipMemcpyHostToHost;
-    }
-    if(!srcTracked && dstTracked)
-    {
-        if(dstInDeviceMem) { kind = hipMemcpyHostToDevice; }
-        else{ kind = hipMemcpyHostToHost; }
-    }
-    if (srcTracked && !dstTracked) {
-        if(srcInDeviceMem) { kind = hipMemcpyDeviceToHost; }
-        else { kind = hipMemcpyHostToHost; }
-    }
-    if (srcTracked && dstTracked) {
-        if(srcInDeviceMem && dstInDeviceMem) { kind = hipMemcpyDeviceToDevice; }
-        if(srcInDeviceMem && !dstInDeviceMem) { kind = hipMemcpyDeviceToHost; }
-        if(!srcInDeviceMem && !dstInDeviceMem) { kind = hipMemcpyHostToHost; }
-        if(!srcInDeviceMem && dstInDeviceMem) { kind = hipMemcpyHostToDevice; }
-    }
+
+    if( srcInDeviceMem  &&  dstInDeviceMem) { kind = hipMemcpyDeviceToDevice; }
+    if( srcInDeviceMem  && !dstInDeviceMem) { kind = hipMemcpyDeviceToHost; }
+    if(!srcInDeviceMem  && !dstInDeviceMem) { kind = hipMemcpyHostToHost; }
+    if(!srcInDeviceMem  &&  dstInDeviceMem) { kind = hipMemcpyHostToDevice; }

    assert (kind != hipMemcpyDefault);
    return kind;
 }


+// hipMemKind must be "resolved" to a specific direction - cannot be default.
+void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind, 
+                                            const hc::AmPointerInfo *dstPtrInfo, 
+                                            const hc::AmPointerInfo *srcPtrInfo, 
+                                            hc::hcCommandKind *hcCopyDir, 
+                                            ihipCtx_t **copyDevice,
+                                            bool *forceUnpinnedCopy)
+{
+    // Ignore what the user tells us and always resolve the direction:
+    // Some apps apparently rely on this.
+    hipMemKind = resolveMemcpyDirection(srcPtrInfo->_isInDeviceMem, dstPtrInfo->_isInDeviceMem);
+
+
+    switch (hipMemKind) {
+        case hipMemcpyHostToHost:     *hcCopyDir = hc::hcMemcpyHostToHost; break;
+        case hipMemcpyHostToDevice:   *hcCopyDir = hc::hcMemcpyHostToDevice; break;
+        case hipMemcpyDeviceToHost:   *hcCopyDir = hc::hcMemcpyDeviceToHost; break;
+        case hipMemcpyDeviceToDevice: *hcCopyDir = hc::hcMemcpyDeviceToDevice; break;
+        default: throw ihipException(hipErrorRuntimeOther);
+    };
+
+    if (srcPtrInfo->_isInDeviceMem) {
+        *copyDevice  = ihipGetPrimaryCtx(srcPtrInfo->_appId);
+    } else if (dstPtrInfo->_isInDeviceMem) {
+        *copyDevice  = ihipGetPrimaryCtx(dstPtrInfo->_appId);
+    } else {
+        *copyDevice = nullptr;
+    }
+
+    *forceUnpinnedCopy = false;
+    if (canSeeMemory(*copyDevice, dstPtrInfo, srcPtrInfo)) {
+
+        if (HIP_FORCE_P2P_HOST & 0x1) {
+            *forceUnpinnedCopy = true;
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", (*copyDevice)->getDeviceNum());
+        } else {
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d) can see src and dst.\n",  (*copyDevice)->getDeviceNum());
+        }
+    } else {
+        *forceUnpinnedCopy = true;
+        tprintf (DB_COPY, "P2P: copy engine(dev:%d) cannot see both host and device pointers - forcing copy with unpinned engine.\n", (*copyDevice)->getDeviceNum());
+    }
+}
+
+
 // TODO - remove kind parm from here or use it below?
 void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes, unsigned kind, bool resolveOn)
 {
@@ -1577,42 +1886,24 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
    bool dstTracked = (hc::am_memtracker_getinfo(&dstPtrInfo, dst) == AM_SUCCESS);
    bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);

-    if (kind == hipMemcpyDefault) {
-        kind = resolveMemcpyDirection(srcTracked, dstTracked, srcPtrInfo._isInDeviceMem, dstPtrInfo._isInDeviceMem);
-    }
+
    hc::hcCommandKind hcCopyDir;
-    switch (kind) {
-        case hipMemcpyHostToHost:     hcCopyDir = hc::hcMemcpyHostToHost; break;
-        case hipMemcpyHostToDevice:   hcCopyDir = hc::hcMemcpyHostToDevice; break;
-        case hipMemcpyDeviceToHost:   hcCopyDir = hc::hcMemcpyDeviceToHost; break;
-        case hipMemcpyDeviceToDevice: hcCopyDir = hc::hcMemcpyDeviceToDevice; break;
-        default: throw ihipException(hipErrorRuntimeOther);
-    };
-
-
-    // If this is P2P access, we need to check to see if the copy agent (specified by the stream where the copy is enqueued)
-    // has peer access enabled to both the source and dest.  If this is true, then the copy agent can see both pointers
-    // and we can perform the access with the copy engine from the current stream.  If not true, then we will copy through the host. (forceHostCopyEngine=true).
-    bool forceHostCopyEngine = false;
-    if (hcCopyDir == hc::hcMemcpyDeviceToDevice) {
-        if (!canSeePeerMemory(ctx, ihipGetPrimaryCtx(dstPtrInfo._appId), ihipGetPrimaryCtx(srcPtrInfo._appId))) {
-            forceHostCopyEngine = true;
-            tprintf (DB_COPY1, "Forcing use of host copy engine.\n");
-        } else {
-            tprintf (DB_COPY1, "Will use SDMA engine on streamDevice=%s.\n", ctx->toString().c_str());
-        }
-    };
-
-    tprintf (DB_COPY1, "locked_copy dir=%s dst=%p src=%p sz=%zu\n", memcpyStr(kind), src, dst, sizeBytes);
+    ihipCtx_t *copyDevice;
+    bool forceUnpinnedCopy;
+    resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, &copyDevice, &forceUnpinnedCopy);

    {
        LockedAccessor_StreamCrit_t crit (_criticalData);
-#if DISABLE_COPY_EXT
-#warning ("Disabled copy_ext path, P2P host staging copies will not work")
-        // Note - peer-to-peer copies which require host staging will not work in this path.
-        crit->_av.copy(src, dst, sizeBytes);
+        tprintf (DB_COPY, "copySync copyDev:%d  dst=%p(home_dev:%d, tracked:%d, isDevMem:%d)  src=%p(home_dev:%d, tracked:%d, isDevMem:%d)   sz=%zu dir=%s forceUnpinnedCopy=%d\n",
+                 copyDevice ? copyDevice->getDeviceNum():-1, 
+                 dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, 
+                 src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem,  
+                 sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
+
+#if USE_COPY_EXT_V2
+        crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? &copyDevice->getDevice()->_acc : nullptr, forceUnpinnedCopy);
 #else
-        crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceHostCopyEngine);
+        crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceUnpinnedCopy);
 #endif
    }
 }
@@ -1624,12 +1915,12 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
    const ihipCtx_t *ctx = this->getCtx();

    if ((ctx == nullptr) || (ctx->getDevice() == nullptr)) {
-        tprintf (DB_COPY1, "locked_copyAsync bad ctx or device\n");
+        tprintf (DB_COPY, "locked_copyAsync bad ctx or device\n");
        throw ihipException(hipErrorInvalidDevice);
    }

    if (kind == hipMemcpyHostToHost) {
-        tprintf (DB_COPY1, "locked_copyAsync: H2H with memcpy");
+        tprintf (DB_COPY, "locked_copyAsync: H2H with memcpy");

        // TODO - consider if we want to perhaps use the GPU SDMA engines anyway, to avoid the host-side sync here and keep everything flowing on the GPU.
        /* As this is a CPU op, we need to wait until all
@@ -1649,24 +1940,39 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
        bool srcTracked = (hc::am_memtracker_getinfo(&srcPtrInfo, src) == AM_SUCCESS);


-        bool copyEngineCanSeeSrcAndDest = true;
-        if ((kind == hipMemcpyDeviceToDevice) || 
-            ((kind == hipMemcpyDefault) && srcTracked && dstTracked)) {
-            copyEngineCanSeeSrcAndDest = canSeePeerMemory(ctx, ihipGetPrimaryCtx(dstPtrInfo._appId), ihipGetPrimaryCtx(srcPtrInfo._appId));
-        }
+        hc::hcCommandKind hcCopyDir;
+        ihipCtx_t *copyDevice;
+        bool forceUnpinnedCopy;
+        resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, &copyDevice, &forceUnpinnedCopy);

-        tprintf (DB_COPY1, "locked_copyAsync:  async memcpy dstTracked=%d srcTracked=%d copyEngineCanSeeSrcAndDest=%d\n",
-                dstTracked, srcTracked, copyEngineCanSeeSrcAndDest);

+        tprintf (DB_COPY, "copyASync copyEngine_dev:%d  dst=%p(home_dev:%d, tracked:%d, isDevMem:%d)  src=%p(home_dev:%d, tracked:%d, isDevMem:%d)   sz=%zu dir=%s.  forceUnpinnedCopy=%d \n",
+                 copyDevice->getDeviceNum(), 
+                 dst, dstPtrInfo._appId, dstTracked, dstPtrInfo._isInDeviceMem, 
+                 src, srcPtrInfo._appId, srcTracked, srcPtrInfo._isInDeviceMem,  
+                 sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);

        // "tracked" really indicates if the pointer's virtual address is available in the GPU address space.
        // If both pointers are not tracked, we need to fall back to a sync copy.
-        if (dstTracked && srcTracked && copyEngineCanSeeSrcAndDest) {
+        if (dstTracked && srcTracked && !forceUnpinnedCopy && copyDevice/*code below assumes this is !nullptr*/) {
            LockedAccessor_StreamCrit_t crit(_criticalData);

-            // Perform asynchronous copy:
+            // Perform fast asynchronous copy - we know copyDevice != NULL based on check above
            try {
-                crit->_av.copy_async(src, dst, sizeBytes);
+                if (HIP_FORCE_SYNC_COPY) {
+#if USE_COPY_EXT_V2
+                    crit->_av.copy_ext      (src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, &copyDevice->getDevice()->_acc, forceUnpinnedCopy);
+#else
+                    crit->_av.copy_ext      (src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceUnpinnedCopy);
+#endif
+
+                } else {
+#if USE_COPY_EXT_V2
+                    crit->_av.copy_async_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, &copyDevice->getDevice()->_acc);
+#else
+                    crit->_av.copy_async(src, dst, sizeBytes);
+#endif
+                }
            } catch (Kalmar::runtime_exception) {
                throw ihipException(hipErrorRuntimeOther);
            };
@@ -1678,12 +1984,42 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
            }

        } else {
-            // TODO - call copy_ext directly here?
-            locked_copySync(dst, src, sizeBytes, kind);
+            LockedAccessor_StreamCrit_t crit(_criticalData);
+#if USE_COPY_EXT_V2
+            crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, copyDevice ? &copyDevice->getDevice()->_acc : nullptr, forceUnpinnedCopy);
+#else
+            crit->_av.copy_ext(src, dst, sizeBytes, hcCopyDir, srcPtrInfo, dstPtrInfo, forceUnpinnedCopy);
+#endif
        }
    }
 }

+
+//-------------------------------------------------------------------------------------------------
+//-------------------------------------------------------------------------------------------------
+//Profiler, really these should live elsewhere:
+hipError_t hipProfilerStart()
+{
+    HIP_INIT_API();
+#if COMPILE_HIP_ATP_MARKER
+    amdtResumeProfiling(AMDT_ALL_PROFILING);
+#endif
+
+    return ihipLogStatus(hipSuccess);
+};
+
+
+hipError_t hipProfilerStop()
+{
+    HIP_INIT_API();
+#if COMPILE_HIP_ATP_MARKER
+    amdtStopProfiling(AMDT_ALL_PROFILING);
+#endif
+
+    return ihipLogStatus(hipSuccess);
+};
+
+
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
 // HCC-specific accessor functions:
@@ -43,26 +43,70 @@ THE SOFTWARE.
 //static const int debug = 0;
 extern const int release;

+// TODO - this blocks both kernels and memory ops.  Perhaps should have separate env var for kernels?
 extern int HIP_LAUNCH_BLOCKING;

 extern int HIP_PRINT_ENV;
-extern int HIP_ATP_MARKER;
+extern int HIP_PROFILE_API;
 //extern int HIP_TRACE_API;
 extern int HIP_ATP;
 extern int HIP_DB;
 extern int HIP_STAGING_SIZE;   /* size of staging buffers, in KB */
 extern int HIP_STREAM_SIGNALS;  /* number of signals to allocate at stream creation */
 extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU identifiers */
+extern int HIP_FORCE_P2P_HOST;


 //---
 // Chicken bits for disabling functionality to work around potential issues:
 extern int HIP_DISABLE_HW_KERNEL_DEP;

+
+// Class to assign a short TID to each new thread, for HIP debugging purposes.
+class ShortTid {
+public:
+
+    ShortTid() ;
+
+    int      tid()       const { return _shortTid; };
+    uint64_t incApiSeqNum() { return ++_apiSeqNum; };
+    uint64_t apiSeqNum() const { return _apiSeqNum; };
+
+private:
+    int      _shortTid;
+
+    // monotonically increasing API sequence number for this threa.
+    uint64_t _apiSeqNum;
+};
+
+struct ProfTrigger {
+
+    static const uint64_t MAX_TRIGGER = std::numeric_limits<uint64_t>::max();
+
+    void print (int tid) {
+        std::cout << "Enabling tracing for ";
+        for (auto iter=_profTrigger.begin(); iter != _profTrigger.end(); iter++) {
+            std::cout << "tid:" << tid << "." << *iter << ",";
+        }
+        std::cout << "\n";
+    };
+
+    uint64_t nextTrigger() { return _profTrigger.empty() ? MAX_TRIGGER : _profTrigger.back(); };
+    void add(uint64_t trigger) { _profTrigger.push_back(trigger); };
+    void sort() { std::sort (_profTrigger.begin(), _profTrigger.end(), std::greater<int>()); };
+private:
+    std::vector<uint64_t> _profTrigger;
+};
+
+
+
 //---
 //Extern tls
 extern thread_local hipError_t tls_lastHipError;
+extern thread_local ShortTid tls_shortTid;

+extern std::vector<ProfTrigger> g_dbStartTriggers;
+extern std::vector<ProfTrigger> g_dbStopTriggers;

 //---
 //Forward defs:
@@ -112,41 +156,42 @@ extern const char *API_COLOR_END;
 #endif


-#define DB_SHOW_TID 0
-
-#if DB_SHOW_TID
-#define COMPUTE_TID_STR \
-    std::stringstream tid_ss;\
-    std::stringstream tid_ss_num;\
-    tid_ss_num << std::this_thread::get_id();\
-    tid_ss << " tid:" << std::hex << std::stoull(tid_ss_num.str());
-#else
-#define COMPUTE_TID_STR std::stringstream tid_ss;
+// Compile code that force hipHostMalloc only allocates finegrained system memory.
+#ifndef HIP_COHERENT_HOST_ALLOC
+#define HIP_COHERENT_HOST_ALLOC 0
 #endif


+
 // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary.
 // TODO - currently we print the trace message at the beginning. if we waited, we could also include return codes, and any values returned
 // through ptr-to-args (ie the pointers allocated by hipMalloc).
 #if COMPILE_HIP_ATP_MARKER
 #include "CXLActivityLogger.h"
-#define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString)
+#define MARKER_BEGIN(markerName,group) amdtBeginMarker(markerName, group, nullptr);
+#define MARKER_END() amdtEndMarker();
+#define RESUME_PROFILING amdtResumeProfiling(AMDT_ALL_PROFILING);
+#define STOP_PROFILING   amdtStopProfiling(AMDT_ALL_PROFILING);
 #else
 // Swallow scoped markers:
-#define SCOPED_MARKER(markerName,group,userString)
+#define MARKER_BEGIN(markerName,group)
+#define MARKER_END()
+#define RESUME_PROFILING
+#define STOP_PROFILING
 #endif


+extern void recordApiTrace(std::string *fullStr, const std::string &apiStr);
+
 #if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1)
 #define API_TRACE(...)\
 {\
-    if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {\
-        std::string s = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\
-        if (COMPILE_HIP_DB && HIP_TRACE_API) {\
-            COMPUTE_TID_STR\
-            fprintf (stderr, "%s<<hip-api:%s %s\n%s" , API_COLOR, tid_ss.str().c_str(), s.c_str(), API_COLOR_END);\
-        }\
-        SCOPED_MARKER(s.c_str(), "HIP", NULL);\
+    if (HIP_PROFILE_API || (COMPILE_HIP_DB && HIP_TRACE_API)) {\
+        std::string apiStr = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')';\
+        std::string fullStr;\
+        recordApiTrace(&fullStr, apiStr);\
+        if      (HIP_PROFILE_API == 0x1) {MARKER_BEGIN(__func__, "HIP") }\
+        else if (HIP_PROFILE_API == 0x2) {MARKER_BEGIN(fullStr.c_str(), "HIP"); }\
    }\
 }
 #else
@@ -174,8 +219,9 @@ extern const char *API_COLOR_END;
        tls_lastHipError = localHipStatus;\
        \
        if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API) {\
-            fprintf(stderr, "  %ship-api: %-30s ret=%2d (%s)>>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\
+            fprintf(stderr, "  %ship-api tid:%d.%lu %-30s ret=%2d (%s)>>%s\n", (localHipStatus == 0) ? API_COLOR:KRED, tls_shortTid.tid(),tls_shortTid.apiSeqNum(),  __func__, localHipStatus, ihipErrorString(localHipStatus), API_COLOR_END);\
        }\
+        if (HIP_PROFILE_API) { MARKER_END(); }\
        localHipStatus;\
    })

@@ -187,10 +233,9 @@ extern const char *API_COLOR_END;
 #define DB_API    0 /* 0x01 - shortcut to enable HIP_TRACE_API on single switch */
 #define DB_SYNC   1 /* 0x02 - trace synchronization pieces */
 #define DB_MEM    2 /* 0x04 - trace memory allocation / deallocation */
-#define DB_COPY1  3 /* 0x08 - trace memory copy commands. . */
+#define DB_COPY   3 /* 0x08 - trace memory copy and peer commands. . */
 #define DB_SIGNAL 4 /* 0x10 - trace signal pool commands */
-#define DB_COPY2  5 /* 0x20 - trace memory copy commands. Detailed. */
-#define DB_MAX_BITPOS 5
+#define DB_MAX_FLAG 5
 // When adding a new debug flag, also add to the char name table below.
 //

@@ -204,20 +249,18 @@ static const DbName dbName [] =
    {KGRN, "api"}, // not used,
    {KYEL, "sync"},
    {KCYN, "mem"},
-    {KMAG, "copy1"},
+    {KMAG, "copy"},
    {KRED, "signal"},
-    {KNRM, "copy2"},
 };

- 
+

 #if COMPILE_HIP_DB
 #define tprintf(trace_level, ...) {\
    if (HIP_DB & (1<<(trace_level))) {\
        char msgStr[1000];\
        snprintf(msgStr, 2000, __VA_ARGS__);\
-        COMPUTE_TID_STR\
-        fprintf (stderr, "  %ship-%s%s:%s%s", dbName[trace_level]._color, dbName[trace_level]._shortName, tid_ss.str().c_str(), msgStr, KNRM); \
+        fprintf (stderr, "  %ship-%s tid:%d:%s%s", dbName[trace_level]._color, dbName[trace_level]._shortName, tls_shortTid.tid(), msgStr, KNRM); \
    }\
 }
 #else
@@ -431,7 +474,7 @@ public:
    void launchModuleKernel(hc::accelerator_view av, hsa_signal_t signal,
                            uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
                            uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
-														uint32_t groupSegmentSize, uint32_t sharedMemBytes, 
+														uint32_t groupSegmentSize, uint32_t sharedMemBytes,
 														void *kernarg, size_t kernSize, uint64_t kernel);


@@ -453,9 +496,14 @@ private:


    // The unsigned return is hipMemcpyKind
-    unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem);
+    unsigned resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem);
+    void resolveHcMemcpyDirection(unsigned hipMemKind, 
+                                  const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo, 
+                                  hc::hcCommandKind *hcCopyDir, 
+                                  ihipCtx_t **copyDevice,
+                                  bool *forceUnpinnedCopy);

-    bool canSeePeerMemory(const ihipCtx_t *thisCtx, ihipCtx_t *dstCtx, ihipCtx_t *srcCtx);
+    bool canSeeMemory(const ihipCtx_t *thisCtx, const hc::AmPointerInfo *dstInfo, const hc::AmPointerInfo *srcInfo);


 private: // Data
@@ -559,16 +607,18 @@ public:


    // Peer Accessor classes:
-    bool isPeer(const ihipCtx_t *peer); // returns True if peer has access to memory physically located on this device.
-    bool addPeer(ihipCtx_t *peer);
-    bool removePeer(ihipCtx_t *peer);
-    void resetPeers(ihipCtx_t *thisDevice);
-    void printPeers(FILE *f) const;
+    bool isPeerWatcher(const ihipCtx_t *peer); // returns True if peer has access to memory physically located on this device.
+    bool addPeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peer);
+    bool removePeerWatcher(const ihipCtx_t *thisCtx, ihipCtx_t *peer);
+    void resetPeerWatchers(ihipCtx_t *thisDevice);
+    void printPeerWatchers(FILE *f) const;

    uint32_t peerCnt() const { return _peerCnt; };
    hsa_agent_t *peerAgents() const { return _peerAgents; };


+    // TODO - move private
+    std::list<ihipCtx_t*>     _peers;     // list of enabled peer devices.

    friend class LockedAccessor<ihipCtxCriticalBase_t>;
 private:
@@ -580,7 +630,6 @@ private:
    // These reflect the currently Enabled set of peers for this GPU:
    // Enabled peers have permissions to access the memory physically allocated on this device.
    // Note the peers always contain the self agent for easy interfacing with HSA APIs.
-    std::list<ihipCtx_t*>     _peers;     // list of enabled peer devices.
    uint32_t                  _peerCnt;     // number of enabled peers
    hsa_agent_t              *_peerAgents;  // efficient packed array of enabled agents (to use for allocations.)
 private:
@@ -619,11 +668,12 @@ public: // Functions:
    ihipCtxCritical_t  &criticalData() { return _criticalData; }; // TODO, move private.  Fix P2P.

    const ihipDevice_t *getDevice() const { return _device; };
+    int                 getDeviceNum() const { return _device->_deviceId; };

    // TODO - review uses of getWriteableDevice(), can these be converted to getDevice()
    ihipDevice_t *getWriteableDevice() const { return _device; };

-    std::string toString() const; 
+    std::string toString() const;

 public:  // Data
    // The NULL stream is used if no other stream is specified.
@@ -712,10 +762,16 @@ inline std::ostream& operator<<(std::ostream& os, const hipEvent_t& e)

 inline std::ostream& operator<<(std::ostream& os, const ihipCtx_t* c)
 {
-    os << "ctx:" << static_cast<const void*> (c) 
-       << " dev:" << c->getDevice()->_deviceId;
+    os << "ctx:" << static_cast<const void*> (c)
+       << ".dev:" << c->getDevice()->_deviceId;
    return os;
 }


+// Helper functions that are used across src files:
+namespace hip_internal {
+    hipError_t memcpyAsync (void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream);
+};
+
+
 #endif
@@ -119,14 +119,31 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
        const unsigned am_flags = 0;
        *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags);

+
        if (sizeBytes && (*ptr == NULL)) {
            hip_status = hipErrorMemoryAllocation;
        } else {
            hc::am_memtracker_update(*ptr, device->_deviceId, 0);
+            int peerCnt=0;
            {
                LockedAccessor_CtxCrit_t crit(ctx->criticalData());
-                if (crit->peerCnt()) {
-                    hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+                // the peerCnt always stores self so make sure the trace actually
+                peerCnt = crit->peerCnt();
+                tprintf(DB_MEM, " allocated device_mem ptr:%p size:%zu on dev:%d and allowed %d other peer(s) access\n",
+                        *ptr, sizeBytes, device->_deviceId, peerCnt-1);
+                if (peerCnt > 1) {
+
+                    //printf ("peer self access\n");
+
+                    // TODOD - remove me:
+                    for (auto iter = crit->_peers.begin(); iter!=crit->_peers.end(); iter++) {
+                        tprintf (DB_MEM, "   allow access to peer: %s%s\n", (*iter)->toString().c_str(), (iter == crit->_peers.begin()) ? " (self)":"");
+                    };
+
+                    hsa_status_t e = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+                    if (e != HSA_STATUS_SUCCESS) {
+                        hip_status = hipErrorMemoryAllocation;
+                    }
                }
            }
        }
@@ -134,11 +151,25 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
        hip_status = hipErrorMemoryAllocation;
    }

-    //printf ("  hipMalloc allocated %p\n", *ptr);

    return ihipLogStatus(hip_status);
 }

+void ihipReadSingleEnv(int *var_ptr, const char *var_name1, const char *description)
+{
+    char * env = getenv(var_name1);
+
+    // Default is set when variable is initialized (at top of this file), so only override if we find
+    // an environment variable.
+    if (env) {
+        long int v = strtol(env, NULL, 0);
+        *var_ptr = (int) (v);
+    }
+    if (HIP_PRINT_ENV) {
+        printf ("%-30s = %2d : %s\n", var_name1, *var_ptr, description);
+    }
+}
+
 hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 {
    HIP_INIT_API(ptr, sizeBytes, flags);
@@ -147,31 +178,58 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)

    auto ctx = ihipGetTlsDefaultCtx();

-    if(ctx){
-        // am_alloc requires writeable __acc, perhaps could be refactored?
-        auto device = ctx->getWriteableDevice();
-        if((flags == hipHostMallocDefault)|| (flags == hipHostMallocPortable)){
-            *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
-            if(sizeBytes < 1 && (*ptr == NULL)){
-                hip_status = hipErrorMemoryAllocation;
-            } else {
-                hc::am_memtracker_update(*ptr, device->_deviceId, amHostPinned);
-            }
-            tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
-        } else if(flags & hipHostMallocMapped){
-            *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
-            if(sizeBytes && (*ptr == NULL)){
-                hip_status = hipErrorMemoryAllocation;
-            }else{
-                hc::am_memtracker_update(*ptr, device->_deviceId, flags);
-                {
-                    LockedAccessor_CtxCrit_t crit(ctx->criticalData());
-                    if (crit->peerCnt()) {
-                        hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+    if (sizeBytes == 0) {
+        hip_status = hipSuccess;
+        // TODO - should size of 0 return err or be siliently ignored?
+    } else if ((ctx==nullptr) || (ptr == nullptr)) {
+        hip_status = hipErrorInvalidValue;
+    } else {
+        unsigned trueFlags = flags;
+        if (flags == hipHostMallocDefault) {
+            trueFlags = hipHostMallocMapped | hipHostMallocWriteCombined;
+        }
+
+        const unsigned supportedFlags = hipHostMallocPortable | hipHostMallocMapped | hipHostMallocWriteCombined;
+
+        // Read from environment variable of HIP_COHERENT_HOST_ALLOC
+        int coherent_alloc=0;
+        ihipReadSingleEnv(&coherent_alloc, "HIP_COHERENT_HOST_ALLOC", "Flag to force allocate finegrained system memory");
+
+        if (flags & ~supportedFlags) {
+            hip_status = hipErrorInvalidValue;
+        }
+        else {
+            auto device = ctx->getWriteableDevice();
+            if(coherent_alloc){
+                // Force to allocate finedgrained system memory
+                *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
+                if(sizeBytes < 1 && (*ptr == NULL)){
+                    hip_status = hipErrorMemoryAllocation;
+                } else {
+                    hc::am_memtracker_update(*ptr, device->_deviceId, amHostCoherent);
+                }
+                tprintf(DB_MEM, " %s: finegrained system memory ptr=%p\n", __func__, *ptr);
+             }
+            else{
+                // TODO - am_alloc requires writeable __acc, perhaps could be refactored?
+                // TODO - hipHostMallocMapped is be ignored on ROCM - all memory is mapped to host address space as WC.
+                *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
+                if (*ptr == NULL) {
+                    hip_status = hipErrorMemoryAllocation;
+                } else {
+                    hc::am_memtracker_update(*ptr, device->_deviceId, flags);
+                    // TODO-hipHostMallocPortable should map the host memory into all contexts, regardless of peer status.
+                    int peerCnt=0;
+                    {
+                        LockedAccessor_CtxCrit_t crit(ctx->criticalData());
+                        peerCnt = crit->peerCnt();
+                        if (peerCnt > 1) {
+                            hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
+                        }
                    }
+                    tprintf(DB_MEM, "allocated pinned_host ptr:%p size:%zu on dev:%d and allow access to %d other peer(s)\n", *ptr, sizeBytes, device->_deviceId, peerCnt-1);
                }
            }
-            tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
        }
    }
    return ihipLogStatus(hip_status);
@@ -354,6 +412,8 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
                    vecAcc.push_back(ihipGetDevice(i)->_acc);
                }
                am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size());
+
+                tprintf(DB_MEM, " %s registered ptr=%p\n", __func__, hostPtr);
                if(am_status == AM_SUCCESS){
                    hip_status = hipSuccess;
                } else {
@@ -377,6 +437,7 @@ hipError_t hipHostUnregister(void *hostPtr)
    }else{
        auto device = ctx->getWriteableDevice();
        am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr);
+        tprintf(DB_MEM, " %s unregistered ptr=%p\n", __func__, hostPtr);
        if(am_status != AM_SUCCESS){
            hip_status = hipErrorHostMemoryNotRegistered;
        }
@@ -398,6 +459,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou
    hc::accelerator acc = ctx->getDevice()->_acc;

    void *ptr = acc.get_symbol_address(symbolName);
+    tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr);

    if(ptr == nullptr)
    {
@@ -427,6 +489,7 @@ hipError_t hipMemcpyToSymbolAsync(const char* symbolName, const void *src, size_
    hc::accelerator acc = ctx->getDevice()->_acc;

    void *ptr = acc.get_symbol_address(symbolName);
+    tprintf(DB_MEM, " symbol '%s' resolved to address:%p\n", symbolName, ptr);

    if(ptr == nullptr)
    {
@@ -555,10 +618,13 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes)
    return ihipLogStatus(e);
 }

-hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream)
-{
-    HIP_INIT_API(dst, src, sizeBytes, kind, stream);

+
+// Internal copy sync:
+namespace hip_internal {
+
+hipError_t memcpyAsync (void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream)
+{
    hipError_t e = hipSuccess;

    stream = ihipSyncAndResolveStream(stream);
@@ -577,86 +643,39 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp
        e = hipErrorInvalidValue;
    }

-    return ihipLogStatus(e);
+    return e;
 }
+} // end namespace hip_internal
+
+
+hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream)
+{
+    HIP_INIT_API(dst, src, sizeBytes, kind, stream);
+
+    return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream));
+
+}
+

 hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream)
 {
    HIP_INIT_API(dst, src, sizeBytes, stream);

-    hipError_t e = hipSuccess;
-
-    stream = ihipSyncAndResolveStream(stream);
-
-    hipMemcpyKind kind = hipMemcpyHostToDevice;
-
-    if ((dst == NULL) || (src == NULL)) {
-        e= hipErrorInvalidValue;
-    } else if (stream) {
-        try {
-            stream->locked_copyAsync((void*)dst, src, sizeBytes, kind);
-        }
-        catch (ihipException ex) {
-            e = ex._code;
-        }
-    } else {
-        e = hipErrorInvalidValue;
-    }
-
-    return ihipLogStatus(e);
+    return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyHostToDevice, stream));
 }

 hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
 {
    HIP_INIT_API(dst, src, sizeBytes, stream);

-    hipError_t e = hipSuccess;
-
-    hipMemcpyKind kind = hipMemcpyDeviceToDevice;
-
-    stream = ihipSyncAndResolveStream(stream);
-
-
-    if ((dst == NULL) || (src == NULL)) {
-        e= hipErrorInvalidValue;
-    } else if (stream) {
-        try {
-            stream->locked_copyAsync((void*)dst, (void*)src, sizeBytes, kind);
-        }
-        catch (ihipException ex) {
-            e = ex._code;
-        }
-    } else {
-        e = hipErrorInvalidValue;
-    }
-
-    return ihipLogStatus(e);
+    return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream));
 }

 hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream)
 {
    HIP_INIT_API(dst, src, sizeBytes, stream);

-    hipError_t e = hipSuccess;
-
-    stream = ihipSyncAndResolveStream(stream);
-
-    hipMemcpyKind kind = hipMemcpyDeviceToHost;
-
-    if ((dst == NULL) || (src == NULL)) {
-        e= hipErrorInvalidValue;
-    } else if (stream) {
-        try {
-            stream->locked_copyAsync(dst, (void*)src, sizeBytes, kind);
-        }
-        catch (ihipException ex) {
-            e = ex._code;
-        }
-    } else {
-        e = hipErrorInvalidValue;
-    }
-
-    return ihipLogStatus(e);
+    return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream));
 }

 // TODO - review and optimize
@@ -35,22 +35,27 @@ THE SOFTWARE.
 // public APIs are thin wrappers which call into this internal implementations.
 // TODO - actually not yet - currently the integer deviceId flavors just call the context APIs.  need to fix.

-/**
- * HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
- */
-//---
-hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx)
-{
-    HIP_INIT_API(canAccessPeer, thisCtx, peerCtx);

+
+hipError_t ihipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx)
+{
    hipError_t err = hipSuccess;


    if ((thisCtx != NULL) && (peerCtx != NULL)) {
+
        if (thisCtx == peerCtx) {
            *canAccessPeer = 0;
+            tprintf(DB_MEM, "Can't be peer to self. (this=%s, peer=%s)\n", 
+                    thisCtx->toString().c_str(), peerCtx->toString().c_str()); 
+        } else  if (HIP_FORCE_P2P_HOST & 0x2) {
+            *canAccessPeer = false;
+            tprintf(DB_MEM, "HIP_FORCE_P2P_HOST denies peer access this=%s peer=%s  canAccessPeer=%d\n", 
+                    thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer); 
        } else {
-             *canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc);
+            *canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc);
+            tprintf(DB_MEM, "deviceCanAccessPeer this=%s peer=%s  canAccessPeer=%d\n", 
+                    thisCtx->toString().c_str(), peerCtx->toString().c_str(), *canAccessPeer); 
        }

    } else {
@@ -59,7 +64,19 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_
    }


-    return ihipLogStatus(err);
+    return err;
+}
+
+
+/**
+ * HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
+ */
+//---
+hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx)
+{
+    HIP_INIT_API(canAccessPeer, thisCtx, peerCtx);
+
+    return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, thisCtx, peerCtx));
 }


@@ -80,8 +97,10 @@ hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx)
            err = hipErrorInvalidDevice;  // Can't disable peer access to self.
        } else {
            LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
-            bool changed = peerCrit->removePeer(thisCtx);
+            bool changed = peerCrit->removePeerWatcher(peerCtx, thisCtx);
            if (changed) {
+                tprintf(DB_MEM, "device %s disable access to memory allocated on peer:%s\n", 
+                                  thisCtx->toString().c_str(), peerCtx->toString().c_str()); 
                // Update the peers for all memory already saved in the tracker:
                am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
            } else {
@@ -112,8 +131,10 @@ hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
        } else if ((thisCtx != NULL) && (peerCtx != NULL)) {
            LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData());
            // Add thisCtx to peerCtx's access list so that new allocations on peer will be made visible to this device:
-            bool isNewPeer = peerCrit->addPeer(thisCtx);
+            bool isNewPeer = peerCrit->addPeerWatcher(peerCtx, thisCtx);
            if (isNewPeer) {
+                tprintf(DB_MEM, "device=%s can now see all memory allocated on peer=%s\n", 
+                                  thisCtx->toString().c_str(), peerCtx->toString().c_str()); 
                am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents());
            } else {
                err = hipErrorPeerAccessAlreadyEnabled;
@@ -134,7 +155,7 @@ hipError_t hipMemcpyPeer (void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t

    // TODO - move to ihip memory copy implementaion.
    // HCC has a unified memory architecture so device specifiers are not required.
-    return hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault);
+    return ihipLogStatus(hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault));
 };


@@ -145,7 +166,7 @@ hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t dstDevice, const void* src, h

    // TODO - move to ihip memory copy implementaion.
    // HCC has a unified memory architecture so device specifiers are not required.
-    return hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream);
+    return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream));
 };


@@ -158,7 +179,7 @@ hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t dstDevice, const void* src, h
 hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId)
 {
    HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId);
-    return hipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId), ihipGetPrimaryCtx(peerDeviceId));
+    return ihipLogStatus(ihipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId), ihipGetPrimaryCtx(peerDeviceId)));
 }


@@ -181,14 +202,14 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
 hipError_t hipMemcpyPeer (void* dst, int  dstDevice, const void* src, int  srcDevice, size_t sizeBytes)
 {
    HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes);
-    return hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes);
+    return ihipLogStatus(hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes));
 }


 hipError_t hipMemcpyPeerAsync (void* dst, int  dstDevice, const void* src, int  srcDevice, size_t sizeBytes, hipStream_t stream)
 {
    HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
-    return hipMemcpyPeerAsync(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes, stream);
+    return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDefault, stream));
 }

 hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
@@ -53,11 +53,11 @@ ctest -R Memcpy

 ### If a test fails - how to debug a test

-Extract the commandline from the testing log:
+Find the test and commandline that fail:

 (From the test build directory, perhaps hip/tests/build)
-$ grep -A3 -m2  hipMemcpy-size  Testing/Temporary/LastTest.log
-36/47 Testing: hipMemcpy-size
-36/47 Test: hipMemcpy-size
-Command: "/home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory/hipMemcpy" "--tests" "0x6"
-Directory: /home/bensander/git/compute/external/hip/hip/tests/b6.hcc-LC.debug/runtimeApi/memory
+grep -IR hipMemcpy-modes -IR ../tests/
+../tests/src/runtimeApi/memory/hipMemcpy.cpp: * RUN_NAMED: %t hipMemcpy-modes --tests 0x1
+
+
+
@@ -45,11 +45,15 @@ int main(int argc, char *argv[])
 { int warpSize, pshift;
  hipDeviceProp_t devProp;
  hipGetDeviceProperties(&devProp, 0);
-  if(strncmp(devProp.name,"Fiji",1)==0)
-{ warpSize =64;
-  pshift =6;
-}
-  else {warpSize =32; pshift=5;}
+  warpSize = devProp.warpSize;
+
+  int w = warpSize;
+  pshift = 0; 
+  while (w >>= 1) ++pshift;
+
+  printf ("warpSize=%d pshift=%d\n", warpSize, pshift);
+
+
  int anycount =0;
  int allcount =0;
  int Num_Threads_per_Block      = 1024;
@@ -48,9 +48,11 @@ int main(int argc, char *argv[])
  hipDeviceProp_t devProp;
  hipGetDeviceProperties(&devProp, 0);
  
-  if(strncmp(devProp.name,"Fiji",1)==0)  
-  {warpSize = 64; pshift =6;}
-  else {warpSize =32; pshift =5;}
+  warpSize = devProp.warpSize;
+
+  int w = warpSize;
+  pshift = 0; 
+  while (w >>= 1) ++pshift;
  
  unsigned int Num_Threads_per_Block      = 512;
  unsigned int Num_Blocks_per_Grid        = 1;
@@ -35,6 +35,7 @@ THE SOFTWARE.

 __global__ void Iter(hipLaunchParm lp, int *Ad, int num){
    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    // Kernel loop designed to execute very slowly... ... ...   so we can test timing-related behavior below
    if(tx == 0){
        for(int i = 0; i<num;i++){
            Ad[tx] += 1;
@@ -61,7 +62,11 @@ int main(){
    for(int i=0;i<NUM_STREAMS;i++){
        HIPCHECK(hipMemcpyAsync(A[i], Ad[i], _SIZE, hipMemcpyDeviceToHost, stream[i]));
    }
-    
+
+   
+   // This first check but relies on the kernel running for so long that the D2H async memcopy has not started yet.   
+   // This will be true in an optimal asynchronous implementation.  
+   // Conservative implementations which synchronize the hipMemcpyAsync will fail, ie if HIP_LAUNCH_BLOCKING=true
    HIPASSERT(1<<30 != A[NUM_STREAMS-1][0]-1);
    HIPCHECK(hipDeviceSynchronize());
    HIPASSERT(1<<30 == A[NUM_STREAMS-1][0]-1);
@@ -76,7 +76,7 @@ dim3 dimBlock(512,1,1);

 hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);

-HIPCHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
+HIPCHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));  // Note this really HostToHost not DeviceToHost, since memory is mapped...
 HIPASSERT(C[10] == 3.0f);
 HIPASSERT(flagA == FlagA);
 HIPASSERT(flagB == FlagB);
@@ -45,11 +45,23 @@ void help(char *argv[])
 {
    printf ("usage: %s [OPTIONS]\n", argv[0]);
    printf (" --memcpyWithPeer : Perform memcpy with peer.\n");
-    printf (" --mirrorPeersi   : Mirror memory onto both default device and peerdevice.  If 0, memory is mapped only on the default device.\n");
+    printf (" --mirrorPeers    : Mirror memory onto both default device and peerdevice.  If 0, memory is mapped only on the default device.\n");
    printf (" --peerDevice N   : Set peer device.\n");
 };


+static hipError_t myHipMemcpy(void *dest, const void *src, size_t sizeBytes, hipMemcpyKind kind,  hipStream_t stream, bool async) 
+{
+    if (async) {
+		hipError_t e = hipMemcpyAsync(dest, src, sizeBytes, kind, stream);
+        //HIPCHECK(hipStreamSynchronize(stream));
+        return (e);
+    } else {
+		return hipMemcpy(dest, src, sizeBytes, kind);
+    };
+}
+
+
 void parseMyArguments(int argc, char *argv[])
 {
    int more_argc = HipTest::parseStandardArguments(argc, argv, false);
@@ -74,6 +86,19 @@ void parseMyArguments(int argc, char *argv[])
    };
 };

+void syncBothDevices()
+{
+    int saveDevice;
+    HIPCHECK(hipGetDevice(&saveDevice));
+    HIPCHECK(hipSetDevice(g_currentDevice));
+    HIPCHECK(hipDeviceSynchronize());
+
+    HIPCHECK(hipSetDevice(g_peerDevice));
+    HIPCHECK(hipDeviceSynchronize());
+
+    HIPCHECK(hipSetDevice(saveDevice));
+}
+

 // Sets globals g_currentDevice, g_peerDevice
 void setupPeerTests() 
@@ -104,9 +129,9 @@ void setupPeerTests()

 //---
 // Test which enables peer2peer first, then allocates the memory.
-void enablePeerFirst()
+void enablePeerFirst(bool useAsyncCopy)
 {
-    printf ("\n==testing: %s\n", __func__);
+    printf ("\n==testing: %s useAsyncCopy=%d\n", __func__, useAsyncCopy);

    setupPeerTests();

@@ -147,11 +172,17 @@ void enablePeerFirst()
    // NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a 
    //        a host staging copy for the P2P access.
    HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
-    HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // This is P2P copy.
+    HIPCHECK (myHipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault, 0/*stream*/, useAsyncCopy)); // This is P2P copy.

    // Copy data back to host:
+    // Have to wait for previous operation to finish, since we are switching to another one:
+    HIPCHECK(hipDeviceSynchronize());
+    
    HIPCHECK (hipSetDevice(g_peerDevice));
-    HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));
+    HIPCHECK (myHipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost, 0/*stream*/, useAsyncCopy));
+    HIPCHECK(hipDeviceSynchronize());
+
+    HIPCHECK (hipSetDevice(g_currentDevice));

    // Check host data:
    for (int i=0; i<N; i++) {
@@ -160,16 +191,16 @@ void enablePeerFirst()
        }
    }

-    printf ("==done: %s\n\n", __func__);
+    printf ("==done: %s useAsyncCopy:%d\n\n", __func__, useAsyncCopy);
 }


 //---
 // Test which allocated memory first, then enables peer2peer.
 // Enabling peer needs to scan all allocated memory and enable peer access.
-void allocMemoryFirst()
+void allocMemoryFirst(bool useAsyncCopy)
 {
-    printf ("\n==testing: %s\n", __func__);
+    printf ("\n==testing: %s useAsyncCopy=%d\n", __func__, useAsyncCopy);

    setupPeerTests();

@@ -211,11 +242,15 @@ void allocMemoryFirst()
    // Copies to test functionality:
    // Device0 push to device1, using P2P:
    HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
-    HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault));
+    HIPCHECK (myHipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault, 0/*stream*/, useAsyncCopy));
+
+    syncBothDevices(); // TODO - remove me, should handle this in implementation.

    // Copy data back to host:
    HIPCHECK (hipSetDevice(g_peerDevice));
-    HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));
+    HIPCHECK (myHipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost, 0/*stream*/, useAsyncCopy));
+
+    syncBothDevices(); // TODO - remove me, should handle this in implementation.


    //---
@@ -225,9 +260,97 @@ void allocMemoryFirst()
            failed("mismatch at index:%d computed:0x%02x, golden memsetval:0x%02x\n", i, (int)A_h[i], (int)memsetval);
        }
    }
-    printf ("==done: %s\n\n", __func__);
+    printf ("==done: %s useAsyncCopy=%d\n\n", __func__, useAsyncCopy);
 }

+
+//---
+// Test which tests peer H2D copy - ie: copy-engine=1, dst=1, src=0 (Host)
+// A_d0 is pinned host on dev0 (this)
+// A_d1 is device memory on dev1 (peer)
+// 
+void testPeerHostToDevice(bool useAsyncCopy)
+{
+    printf ("\n==testing: %s useAsyncCopy=%d\n", __func__, useAsyncCopy);
+
+    setupPeerTests();
+
+    // Always enable g_currentDevice to see the allocations on peerDevice.
+    HIPCHECK(hipSetDevice(g_currentDevice));
+    HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0));
+
+    if (p_mirrorPeers) {
+        // Mirror peers allows the peer device to see the allocations on currentDevice.
+        int canAccessPeer;
+        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_peerDevice, g_currentDevice));
+        assert(canAccessPeer);
+
+        HIPCHECK(hipSetDevice(g_peerDevice));
+        HIPCHECK(hipDeviceEnablePeerAccess(g_currentDevice, 0));
+    }
+
+    size_t Nbytes = N*sizeof(char);
+
+    char *A_host_d0, *A_d1;
+    char *A_h;
+
+    A_h = (char*)malloc(Nbytes);
+
+    // allocate and initialize memory on device0
+    HIPCHECK (hipSetDevice(g_currentDevice));
+    HIPCHECK (hipHostMalloc(&A_host_d0, Nbytes) );
+    HIPCHECK (hipMemset(A_host_d0, memsetval, Nbytes) ); 
+
+    // allocate and initialize memory on peer device
+    HIPCHECK (hipSetDevice(g_peerDevice));
+    HIPCHECK (hipMalloc(&A_d1, Nbytes) );
+    HIPCHECK (hipMemset(A_d1, 0x13, Nbytes) ); 
+
+    bool firstAsyncCopy = useAsyncCopy; /*TODO - should be useAsyncCopy*/
+
+    syncBothDevices();
+
+
+
+    // Device0 push to device1, using P2P:
+    // NOTE : if p_mirrorPeers=0 and p_memcpyWithPeer=1, then peer device does not have mapping for A_d1 and we need to use a 
+    //        a host staging copy for the P2P access.
+    if (p_memcpyWithPeer) {
+        // p_memcpyWithPeer=1 case is HostToDevice.
+        // if p_mirrorPeers = 1, this is accelerated copy over PCIe.  
+        // if p_mirrorPeers = 0, this should fall back to host (because peer can't see A_host_d0)
+        HIPCHECK (hipSetDevice(g_peerDevice));
+        HIPCHECK (myHipMemcpy(A_d1, A_host_d0, Nbytes, hipMemcpyHostToDevice, 0/*stream*/, firstAsyncCopy)); // This is P2P copy.
+    } else {
+        // p_memcpyWithPeer=0 case is HostToDevice.
+        // if p_mirrorPeers = 1, this is accelerated copy over PCIe.  
+        // if p_mirrorPeers = 0, this should fall back to host (because device0 can't see A_d1)
+        HIPCHECK (hipSetDevice(g_currentDevice));
+        HIPCHECK (myHipMemcpy(A_d1, A_host_d0, Nbytes, hipMemcpyHostToDevice, 0/*stream*/, firstAsyncCopy)); // This is P2P copy.
+    }
+
+    syncBothDevices();
+
+    // Copy data back to host:
+    HIPCHECK (hipSetDevice(g_peerDevice));
+    HIPCHECK (myHipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost, 0/*stream*/, useAsyncCopy));
+    HIPCHECK(hipDeviceSynchronize());
+
+    HIPCHECK (hipSetDevice(g_currentDevice));
+    HIPCHECK(hipDeviceSynchronize());
+
+    // Check host data:
+    for (int i=0; i<N; i++) {
+        if (A_h[i] != memsetval) {
+            failed("mismatch at index:%d computed:0x%02x, golden memsetval:0x%02x\n", i, (int)A_h[i], (int)memsetval);
+        }
+    }
+
+    printf ("==done: %s useAsyncCopy:%d\n\n", __func__, useAsyncCopy);
+}
+
+
+
 void simpleNegative()
 {
    printf ("\n==testing: %s\n", __func__);
@@ -269,17 +392,30 @@ int main(int argc, char *argv[])
 {
    parseMyArguments(argc, argv);

+
+    if (p_tests & 0x100) {
+        testPeerHostToDevice(false/*useAsyncCopy*/);
+    }
+    testPeerHostToDevice(true/*useAsyncCopy*/);
+
    if (p_tests & 0x1) {
-        enablePeerFirst();
+        enablePeerFirst(false/*useAsyncCopy*/);
    }

    if (p_tests & 0x2) {
-        allocMemoryFirst();
+        allocMemoryFirst(false/*useAsyncCopy*/);
    }

    if (p_tests & 0x4) {
        simpleNegative();
    }

+    if (p_tests & 0x8) {
+        enablePeerFirst(true/*useAsyncCopy*/);
+    }
+    if (p_tests & 0x10) {
+        allocMemoryFirst(true/*useAsyncCopy*/);
+    }
+
    passed();
 }
@@ -0,0 +1,61 @@
+/*
+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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * HIT_END
+ */
+
+
+#include <cstdio>
+#include "hip/hip_runtime.h"
+
+__global__ void Kernel(hipLaunchParm lp,volatile float* hostRes)
+{
+    int tid = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    hostRes[tid] = tid + 1;
+    __threadfence_system();
+    // expecting that the data is getting flushed to host here!
+    // time waster for-loop (sleep)
+    for (int timeWater = 0; timeWater  < 100000000; timeWater++);
+}
+
+int main()
+{
+    size_t blocks = 2;
+    volatile float* hostRes;
+    hipHostMalloc((void**)&hostRes,blocks*sizeof(float),hipHostMallocMapped);
+    hostRes[0]=0;
+    hostRes[1]=0;
+    hipLaunchKernel(HIP_KERNEL_NAME(Kernel), dim3(1), dim3(blocks), 0, 0, hostRes);
+    int eleCounter = 0;
+    while (eleCounter < blocks)
+    {
+        // blocks until the value changes
+        while(hostRes[eleCounter] == 0);
+        printf("%f\n", hostRes[eleCounter]);;
+        eleCounter++;
+    }
+    hipHostFree((void *)hostRes);
+    return 0;
+}
+
@@ -0,0 +1,60 @@
+/* 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. */
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
+ * RUN: %t
+ * HIT_END
+ */
+
+#include <iostream>
+#include <sstream>
+#include <vector>
+#include <stdio.h>
+#include <stdlib.h>
+#include <assert.h>
+#include <string>
+#include "hip/hip_runtime.h"
+using namespace std;
+
+string getRes(){
+    FILE *in;
+    char buff[512], buff_2[512];
+    string str = "./hipMemoryAllocateCoherent";
+    if(!(in = popen(str.c_str(), "r"))){
+        exit(1);
+    }
+    fgets(buff, sizeof(buff), in);
+    fgets(buff_2, sizeof(buff_2), in);
+    string str_buff = buff;
+    str_buff += buff_2;
+    pclose(in);
+    return str_buff;
+}
+
+int main() {
+    setenv("HIP_COHERENT_HOST_ALLOC","1000,0,1",1);
+    string output = getRes();
+    istringstream buffer(output);
+    double res1, res2;
+    buffer >> res1;
+    buffer >> res2;
+    if((res2-res1*2)>0.000001)
+        exit(1);
+    std::cout << "PASSED" << std::endl;
+    return 0;
+}