From 2fbbf2b136b3672b814974de30f34eb6e17ae325 Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Mon, 28 Nov 2016 10:47:18 -0600
Subject: [PATCH 01/28] Change the parameter type of hipDeviceGetPCIBusID to
 char*

Change-Id: Ia72f403126e95f65da53208fc246f45d1417381f
---
 include/hip/hcc_detail/hip_runtime_api.h  | 2 +-
 include/hip/nvcc_detail/hip_runtime_api.h | 2 +-
 src/hip_device.cpp                        | 8 ++++----
 3 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index ee703c4eec..688cfe9437 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -1597,7 +1597,7 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device);
  *
  * @returns #hipSuccess, #hipErrorInavlidDevice
  */
-hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device);
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
 
 /**
  * @brief Returns a handle to a compute device.
diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h
index 625448094b..b304218883 100644
--- a/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/include/hip/nvcc_detail/hip_runtime_api.h
@@ -765,7 +765,7 @@ inline static hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return hipCUResultTohipError(cuDeviceGetName(name,len,device));
 }
 
-inline static hipError_t hipDeviceGetPCIBusId(int *pciBusId,int len,hipDevice_t device)
+inline static hipError_t hipDeviceGetPCIBusId(char* pciBusId,int len,hipDevice_t device)
 {
     return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device));
 }
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 4f5729791e..1275c8f19b 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -306,7 +306,7 @@ hipError_t hipSetDeviceFlags( unsigned int flags)
            }
        }
 
-       unsigned supportedFlags = hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax; 
+       unsigned supportedFlags = hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax;
 
        if (flags & (~supportedFlags)) {
           e = hipErrorInvalidValue;
@@ -338,12 +338,12 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return ihipLogStatus(e);
 }
 
-hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device)
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
 {
-    HIP_INIT_API(pciBusId,len, device);
+    HIP_INIT_API(pciBusId, len, device);
     hipError_t e = hipSuccess;
     int deviceId= device->_deviceId;
-    e = ihipDeviceGetAttribute(pciBusId, hipDeviceAttributePciBusId, deviceId);
+    e = ihipDeviceGetAttribute((int*)pciBusId, hipDeviceAttributePciBusId, deviceId);
     return ihipLogStatus(e);
 }
 

From fe6ba656c9e1c99e056d94da0a19f61abfda1ab2 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Tue, 29 Nov 2016 22:04:09 +0530
Subject: [PATCH 02/28] Added support for hipMemGetAddressRange

Change-Id: I99a796a4eb765152cf15a12d6a86b58684d34f50
---
 include/hip/hcc_detail/hip_runtime_api.h  | 12 ++++++++++++
 include/hip/nvcc_detail/hip_runtime_api.h |  5 +++++
 src/hip_memory.cpp                        | 15 +++++++++++++++
 3 files changed, 32 insertions(+)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 688cfe9437..83c60f3e35 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -1281,6 +1281,18 @@ hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);
  */
 hipError_t  hipDeviceDisablePeerAccess (int peerDeviceId);
 
+/**
+ * @brief Get information on memory allocations.
+ *
+ * @param [out] pbase - BAse pointer address
+ * @param [out] psize - Size of allocation
+ * @param [in]  dptr- Device Pointer
+ *
+ * @returns #hipSuccess, #hipErrorInvalidDevicePointer
+ *
+ * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
+ */
+hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr );
 
 #ifndef USE_PEER_NON_UNIFIED
 #define USE_PEER_NON_UNIFIED 1
diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h
index b304218883..1a24615475 100644
--- a/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/include/hip/nvcc_detail/hip_runtime_api.h
@@ -639,6 +639,11 @@ inline static hipError_t  hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned in
     return hipCUResultTohipError(cuCtxEnablePeerAccess(peerCtx, flags));
 }
 
+inline static hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr )
+{
+    return hipCUResultTohipError(cuMemGetAddressRange( pbase , psize , dptr));
+}
+
 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));
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 089d1a626e..314890d167 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -28,6 +28,7 @@ THE SOFTWARE.
 #include "hip_hcc.h"
 #include "trace_helper.h"
 #include "hip/hcc_detail/hip_texture.h"
+#include <hc_am.hpp>
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
@@ -1037,3 +1038,17 @@ hipError_t hipFreeArray(hipArray* array)
     return ihipLogStatus(hipStatus);
 }
 
+hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr )
+{
+    HIP_INIT_API ( pbase , psize , dptr );
+    hipError_t hipStatus = hipSuccess;
+    hc::accelerator acc;
+    hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 );
+    am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , dptr );
+    if (status == AM_SUCCESS) {
+        *pbase = amPointerInfo._devicePointer;
+        *psize = amPointerInfo._sizeBytes;
+    }
+    hipStatus = hipErrorInvalidDevicePointer;
+    return hipStatus;
+}

From 0dfcd3e66474ea31dedc5ea67047e482e7d402a3 Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Tue, 29 Nov 2016 11:33:51 -0600
Subject: [PATCH 03/28] Change to use produce device name by default

Change-Id: Ie2cee2a2e94a08b5874a2f5abee5d1ab6c9fdf47
---
 src/hip_hcc.cpp | 46 +++++++++++++++++++++++-----------------------
 1 file changed, 23 insertions(+), 23 deletions(-)

diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 36790271de..06402f9a67 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -78,7 +78,7 @@ 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. 
+// Force async copies to actually use the synchronous copy interface.
 int HIP_FORCE_SYNC_COPY = 0;
 
 
@@ -86,7 +86,7 @@ int HIP_FORCE_SYNC_COPY = 0;
 
 
 
-#define HIP_USE_PRODUCT_NAME 0
+#define HIP_USE_PRODUCT_NAME 1
 //#define DISABLE_COPY_EXT 1
 
 
@@ -105,7 +105,7 @@ unsigned g_numLogicalThreads;
 std::atomic<int> g_lastShortTid(1);
 
 // Indexed by short-tid:
-// 
+//
 std::vector<ProfTrigger> g_dbStartTriggers;
 std::vector<ProfTrigger> g_dbStopTriggers;
 
@@ -133,12 +133,12 @@ void recordApiTrace(std::string *fullStr, const std::string &apiStr)
 
     if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) {
         printf ("info: resume profiling at %lu\n", apiSeqNum);
-        RESUME_PROFILING; 
+        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; 
+        STOP_PROFILING;
         g_dbStopTriggers.pop_back();
     };
 
@@ -211,8 +211,8 @@ hipError_t ihipSynchronize(void)
 //=================================================================================================
 ShortTid::ShortTid()  :
     _apiSeqNum(0)
-{ 
-    _shortTid = g_lastShortTid.fetch_add(1); 
+{
+    _shortTid = g_lastShortTid.fetch_add(1);
 
     if (COMPILE_HIP_DB && HIP_TRACE_API) {
         std::stringstream tid_ss;
@@ -282,7 +282,7 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
         } else if (HIP_WAIT_MODE == 2) {
             waitMode = hc::hcWaitModeActive;
         }
-            
+
         crit->_av.wait(waitMode);
     }
 
@@ -455,7 +455,7 @@ bool ihipCtxCriticalBase_t<CtxMutex>::addPeerWatcher(const ihipCtx_t *thisCtx, i
     auto match = std::find(_peers.begin(), _peers.end(), peerWatcher);
     if (match == std::end(_peers)) {
         // Not already a peer, let's update the list:
-        tprintf(DB_COPY, "addPeerWatcher.  Allocations on %s now visible to peerWatcher %s.\n", 
+        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();
@@ -473,7 +473,7 @@ bool ihipCtxCriticalBase_t<CtxMutex>::removePeerWatcher(const ihipCtx_t *thisCtx
     auto match = std::find(_peers.begin(), _peers.end(), peerWatcher);
     if (match != std::end(_peers)) {
         // Found a valid peer, let's remove it.
-        tprintf(DB_COPY, "removePeerWatcher.  Allocations on %s no longer visible to former peerWatcher %s.\n", 
+        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();
@@ -813,7 +813,7 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop)
     prop->arch.hasFloatAtomicAdd           = 0;
     prop->arch.hasGlobalInt64Atomics       = 1;
     prop->arch.hasSharedInt64Atomics       = 1;
-    prop->arch.hasDoubles                  = 1; 
+    prop->arch.hasDoubles                  = 1;
     prop->arch.hasWarpVote                 = 1;
     prop->arch.hasWarpBallot               = 1;
     prop->arch.hasWarpShuffle              = 1;
@@ -1181,7 +1181,7 @@ std::string HIP_DB_callback(void *var_ptr, const char *envVarString)
         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)) { 
+                if (!strcmp(t->c_str(), dbName[i]._shortName)) {
                     *var_ptr_int |= (1<<i);
                 } // TODO - else throw error?
             }
@@ -1257,13 +1257,13 @@ void ihipInit()
     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-separated 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_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.
@@ -1716,10 +1716,10 @@ unsigned ihipStream_t::resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDev
 
 
 // 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, 
+void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,
+                                            const hc::AmPointerInfo *dstPtrInfo,
+                                            const hc::AmPointerInfo *srcPtrInfo,
+                                            hc::hcCommandKind *hcCopyDir,
                                             ihipCtx_t **copyDevice,
                                             bool *forceUnpinnedCopy)
 {
@@ -1749,11 +1749,11 @@ void ihipStream_t::resolveHcMemcpyDirection(unsigned hipMemKind,
 
         if (HIP_FORCE_P2P_HOST & 0x1) {
             *forceUnpinnedCopy = true;
-            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n", 
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst but HIP_FORCE_P2P_HOST=0, forcing copy through staging buffers.\n",
                     (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
 
         } else {
-            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst.\n",  
+            tprintf (DB_COPY, "P2P.  Copy engine (dev:%d agent=0x%lx) can see src and dst.\n",
                     (*copyDevice)->getDeviceNum(), (*copyDevice)->getDevice()->_hsaAgent.handle);
         }
     } else {
@@ -1789,7 +1789,7 @@ void ihipStream_t::locked_copySync(void* dst, const void* src, size_t sizeBytes,
     {
         LockedAccessor_StreamCrit_t crit (_criticalData);
         tprintf (DB_COPY, "copySync copyDev:%d  dst=%p (phys_dev:%d, isDevMem:%d)  src=%p(phys_dev:%d, isDevMem:%d)   sz=%zu dir=%s forceUnpinnedCopy=%d\n",
-                 copyDevice ? copyDevice->getDeviceNum():-1, 
+                 copyDevice ? copyDevice->getDeviceNum():-1,
                  dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
                  src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
                  sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);
@@ -1846,7 +1846,7 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
         bool forceUnpinnedCopy;
         resolveHcMemcpyDirection(kind, &dstPtrInfo, &srcPtrInfo, &hcCopyDir, &copyDevice, &forceUnpinnedCopy);
         tprintf (DB_COPY, "copyASync copyDev:%d  dst=%p (phys_dev:%d, isDevMem:%d)  src=%p(phys_dev:%d, isDevMem:%d)   sz=%zu dir=%s forceUnpinnedCopy=%d\n",
-                 copyDevice ? copyDevice->getDeviceNum():-1, 
+                 copyDevice ? copyDevice->getDeviceNum():-1,
                  dst, dstPtrInfo._appId, dstPtrInfo._isInDeviceMem,
                  src, srcPtrInfo._appId, srcPtrInfo._isInDeviceMem,
                  sizeBytes, hcMemcpyStr(hcCopyDir), forceUnpinnedCopy);

From 7b0650773c5db409d8f6ade0e35da9b58e08c6e3 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 29 Nov 2016 19:46:01 -0600
Subject: [PATCH 04/28] added half add and fma intrinsic

Change-Id: Ifa60c1a7065f524f069291bb00d987b11c836cc4
---
 src/hip_ir.ll | 22 ++++++++++++++++++++++
 1 file changed, 22 insertions(+)

diff --git a/src/hip_ir.ll b/src/hip_ir.ll
index 6850293778..d0e2a879a3 100644
--- a/src/hip_ir.ll
+++ b/src/hip_ir.ll
@@ -12,4 +12,26 @@ define void @__threadfence_block()  #1 {
     ret void
 }
 
+define linkonce_odr spir_func i32 @__rocm_dp4a(i32 %in1, i32 %in2, i32 %in3) {
+  %val1 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0","=v,v,v"(i32 %in1, i32 %in2)
+  %ret1 = add i32 %val1, %in3
+  %val2 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:BYTE_1","=v,v,v"(i32 %in1, i32 %in2)
+  %ret2 = add i32 %ret1, %val2
+  %val3 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_2 src1_sel:BYTE_2","=v,v,v"(i32 %in1, i32 %in2)
+  %ret3 = add i32 %val3, %ret2
+  %val4 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:BYTE_3","=v,v,v"(i32 %in1, i32 %in2)
+	%ret4 = add i32 %val4, %ret3
+  ret i32 %ret4
+}
+
+define linkonce_odr spir_func i32 @__rocm_hfma(i32 %in1, i32 %in2, i32 %in3) {
+  tail call void asm "v_mac_f16 $0, $1, $2","v,v,v"(i32 %in1, i32 %in2, i32 %in3)
+  ret i32 %in3
+}
+
+define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) {
+	%val = tail call i32 asm "v_add_f16 $0, $1, $2","=v,v,v"(i32 %in1, i32 %in2)
+  ret i32 %val
+}
+
 attributes #1 = { alwaysinline nounwind }

From 1e9dc5b521dd2c22408e1b64f4d1203386d05f00 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 1 Dec 2016 12:47:37 +0530
Subject: [PATCH 05/28] Fix some broken directed tests

Change-Id: I40f2661a74951f8d91824df8fd9ef0cc2312d183
---
 tests/src/context/hipCtx_simple.cpp        | 2 +-
 tests/src/hipHcc.cpp                       | 2 +-
 tests/src/kernel/hipLanguageExtensions.cpp | 2 +-
 3 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/tests/src/context/hipCtx_simple.cpp b/tests/src/context/hipCtx_simple.cpp
index 086cc77549..bb2b32a56b 100644
--- a/tests/src/context/hipCtx_simple.cpp
+++ b/tests/src/context/hipCtx_simple.cpp
@@ -21,7 +21,7 @@ THE SOFTWARE.
 */
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp HCC_OPTIONS --stdlib=libc++
+ * BUILD: %t %s ../test_common.cpp HCC_OPTIONS -stdlib=libc++
  * RUN: %t
  * HIT_END
  */
diff --git a/tests/src/hipHcc.cpp b/tests/src/hipHcc.cpp
index 1c54103b08..b09898309e 100644
--- a/tests/src/hipHcc.cpp
+++ b/tests/src/hipHcc.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 // Test the HCC-specific API extensions for HIP:
 
 /* HIT_START
- * BUILD: %t %s HCC_OPTIONS --stdlib=libc++
+ * BUILD: %t %s HCC_OPTIONS -stdlib=libc++
  * RUN: %t
  * HIT_END
  */
diff --git a/tests/src/kernel/hipLanguageExtensions.cpp b/tests/src/kernel/hipLanguageExtensions.cpp
index 6ad0c382a0..cc170f9c8a 100644
--- a/tests/src/kernel/hipLanguageExtensions.cpp
+++ b/tests/src/kernel/hipLanguageExtensions.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 // Collection of code to make sure that various features in the hip kernel language compile.
 
 /* HIT_START
- * BUILD: %t %s ../test_common.cpp HCC_OPTIONS --stdlib=libc++
+ * BUILD: %t %s ../test_common.cpp HCC_OPTIONS -stdlib=libc++
  * RUN: %t
  * HIT_END
  */

From df9faffe936b8bc296a55f94d40eae25b8a979c6 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 1 Dec 2016 12:51:58 +0530
Subject: [PATCH 06/28] hipcc: HCC workaround no longer needs env var

Change-Id: I81f4eafddbda5e9e2f1082932dd502ab451cfc24
---
 bin/hipcc | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/bin/hipcc b/bin/hipcc
index 09c4d813d0..a384a86fbf 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -76,6 +76,7 @@ if ($HIP_PLATFORM eq "hcc") {
     $HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc";
 
     $HCC_VERSION=`${HCC_HOME}/bin/hcc --version | cut -d" " -f9 | tr -d "\n"`;
+    $HCC_VERSION_MAJOR=`${HCC_HOME}/bin/hcc --version | cut -d" " -f9 | cut -d"." -f1 | tr -d "\n"`;
 
     $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
 
@@ -92,14 +93,14 @@ if ($HIP_PLATFORM eq "hcc") {
     $HIPCXXFLAGS = $HCCFLAGS;
 
     #### GCC system includes workaround ####
-    $WA = $ENV{'HIP_HCC_SYS_INCLUDES_WA'} // 0;
-    if (${WA}) {
+    $HCC_WA_FLAGS = " ";
+    if ($HCC_VERSION_MAJOR eq 1) {
         my $GCC_CUR_VER = `gcc -dumpversion`;
         my $GPP_CUR_VER = `g++ -dumpversion`;
         $GCC_CUR_VER =~ s/\R//g;
         $GPP_CUR_VER =~ s/\R//g;
         if (${GCC_CUR_VER} eq ${GPP_CUR_VER}) {
-            $HIPCXXFLAGS .= "-I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER}";
+            $HCC_WA_FLAGS .= " -I/usr/include/x86_64-linux-gnu -I/usr/include/x86_64-linux-gnu/c++/${GCC_CUR_VER} -I/usr/include/c++/${GCC_CUR_VER} ";
         }
     }
 
@@ -260,6 +261,7 @@ foreach $arg (@ARGV)
     if(($arg eq '-stdlib=libstdc++') and ($setStdLib eq 0))
     {
         $HIPCXXFLAGS .= " -stdlib=libstdc++";
+        $HIPCXXFLAGS .= $HCC_WA_FLAGS;
         $setStdLib = 1;
     }
     if($arg eq '--version') {
@@ -322,6 +324,7 @@ if ($buildDeps and $HIP_PLATFORM eq 'nvcc') {
 if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
 {
     $HIPCXXFLAGS .= " -stdlib=libstdc++";
+    $HIPCXXFLAGS .= $HCC_WA_FLAGS;
 }
 
 if ($needHipHcc) {

From ef046c7098b1a445ed362e112edbee8f8845bc52 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 1 Dec 2016 15:33:12 +0530
Subject: [PATCH 07/28] Improve shared and static library support

hipcc accepts new parameter -use-staticlib and -use-sharedlib to
control linking behavior. Default is still static library.

Change-Id: I28fb9a939f8177c75abefd8b77d8118a6666d1f4
---
 CMakeLists.txt        | 40 ++++++++++------------------------------
 bin/hipcc             | 23 ++++++++++++++---------
 packaging/hip_hcc.txt | 11 +++--------
 3 files changed, 27 insertions(+), 47 deletions(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index 1b95145b96..6f9a819c4d 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -115,16 +115,6 @@ else()
     message(FATAL_ERROR "Don't know where to install HIP. Please specify absolute path using -DCMAKE_INSTALL_PREFIX")
 endif()
 
-# Set if we need to build shared or static library
-if(NOT DEFINED HIP_LIB_TYPE)
-    if(NOT DEFINED ENV{HIP_LIB_TYPE})
-        set(HIP_LIB_TYPE 1)
-    else()
-        set(HIP_LIB_TYPE $ENV{HIP_LIB_TYPE})
-    endif()
-endif()
-add_to_config(_buildInfo HIP_LIB_TYPE)
-
 # Check if we need to build hipify-clang
 if(NOT DEFINED HIPIFY_CLANG_LLVM_DIR)
     if(DEFINED ENV{HIPIFY_CLANG_LLVM_DIR})
@@ -173,7 +163,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
     set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${HIP_HCC_BUILD_FLAGS}")
 
-    set(SOURCE_FILES_SHARED
+    set(SOURCE_FILES_RUNTIME
         src/hip_hcc.cpp
         src/hip_context.cpp
         src/hip_device.cpp
@@ -184,27 +174,23 @@ if(HIP_PLATFORM STREQUAL "hcc")
         src/hip_stream.cpp
         src/hip_module.cpp)
 
-    set(SOURCE_FILES_STATIC
+    set(SOURCE_FILES_DEVICE
         src/device_util.cpp
         src/hip_ldg.cpp
         src/hip_fp16.cpp)
 
-    if(${HIP_LIB_TYPE} EQUAL 0)
-        add_library(hip_hcc OBJECT ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
-    elseif(${HIP_LIB_TYPE} EQUAL 1)
-        add_library(hip_hcc STATIC ${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC})
-    else()
-        set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
-        add_library(hip_hcc SHARED ${SOURCE_FILES_SHARED})
-        add_library(hip_device STATIC ${SOURCE_FILES_STATIC})
-        add_dependencies(hip_device hip_hcc)
-    endif()
+    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L${HCC_HOME}/lib -lmcwamp -Wl,-Bsymbolic")
+    add_library(hip_hcc SHARED ${SOURCE_FILES_RUNTIME})
+    add_library(hip_hcc_static STATIC ${SOURCE_FILES_RUNTIME})
+    add_dependencies(hip_hcc_static hip_hcc)
+    add_library(hip_device STATIC ${SOURCE_FILES_DEVICE})
+    add_dependencies(hip_device hip_hcc)
 
     # Generate hcc_version.txt
     add_custom_target(query_hcc_version COMMAND ${HCC_HOME}/bin/hcc --version > ${PROJECT_BINARY_DIR}/hcc_version.tmp)
     add_custom_target(check_hcc_version COMMAND ${CMAKE_COMMAND} -E copy_if_different ${PROJECT_BINARY_DIR}/hcc_version.tmp ${PROJECT_BINARY_DIR}/hcc_version.txt DEPENDS query_hcc_version)
     set_source_files_properties(${PROJECT_BINARY_DIR}/hcc_version.txt PROPERTIES GENERATED TRUE)
-    set_source_files_properties(${SOURCE_FILES_SHARED} ${SOURCE_FILES_STATIC} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
+    set_source_files_properties(${SOURCE_FILES_RUNTIME} ${SOURCE_FILES_DEVICE} PROPERTIES OBJECT_DEPENDS ${PROJECT_BINARY_DIR}/hcc_version.txt)
     add_dependencies(hip_hcc check_hcc_version update_build_and_version_info)
 
     # Generate .hipInfo
@@ -223,13 +209,7 @@ add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMA
 #############################
 # Install hip_hcc if platform is hcc
 if(HIP_PLATFORM STREQUAL "hcc")
-    if(${HIP_LIB_TYPE} EQUAL 0)
-        install(DIRECTORY ${PROJECT_BINARY_DIR}/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
-    elseif(${HIP_LIB_TYPE} EQUAL 1)
-        install(TARGETS hip_hcc DESTINATION lib)
-    else()
-        install(TARGETS hip_hcc hip_device DESTINATION lib)
-    endif()
+    install(TARGETS hip_hcc_static hip_hcc hip_device DESTINATION lib)
     install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_ir.ll DESTINATION lib)
 
     # Install .hipInfo
diff --git a/bin/hipcc b/bin/hipcc
index a384a86fbf..7d2675d0ba 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -193,7 +193,8 @@ my $needHipHcc = ($HIP_PLATFORM eq 'hcc');      # set if we need to link hip_hcc
 my $printHipVersion = 0;    # print HIP version
 my $runCmd = 1;
 my $buildDeps = 0;
-
+my $linkType = 0;
+my $setLinkType = 0;
 
 my @options = ();
 my @inputs  = ();
@@ -275,11 +276,20 @@ foreach $arg (@ARGV)
         $compileOnly = 1;
         $buildDeps = 1;
     }
-
     if($arg eq '-use_fast_math') {
         print "In fast Math";
         $HIPCXXFLAGS .= " -DHIP_FAST_MATH ";
     }
+    if(($arg eq '-use-staticlib') and ($setLinkType eq 0))
+    {
+        $linkType = 0;
+        $setLinkType = 1;
+    }
+    if(($arg eq '-use-sharedlib') and ($setLinkType eq 0))
+    {
+        $linkType = 1;
+        $setLinkType = 1;
+    }
 
     if ($arg =~ m/^-/) {
         # options start with -
@@ -328,13 +338,8 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
 }
 
 if ($needHipHcc) {
-    $HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 1;
-
-    # TODO - remove the old sea-of-objects solution:
-    if ($HIP_LIB_TYPE eq 0) {
-        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) {
-        substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc " ;
+    if ($linkType eq 0) {
+        substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -lhip_hcc_static -lhip_device " ;
     } else {
         substr($HIPLDFLAGS,0,0) = " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc -lhip_device ";
     }
diff --git a/packaging/hip_hcc.txt b/packaging/hip_hcc.txt
index 78717247f0..00a62ab14c 100644
--- a/packaging/hip_hcc.txt
+++ b/packaging/hip_hcc.txt
@@ -1,14 +1,9 @@
 cmake_minimum_required(VERSION 2.8.3)
 project(hip_hcc)
 
-if(@HIP_LIB_TYPE@ EQUAL 0)
-    install(DIRECTORY @PROJECT_BINARY_DIR@/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
-elseif(@HIP_LIB_TYPE@ EQUAL 1)
-    install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.a DESTINATION lib)
-else()
-    install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
-    install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
-endif()
+install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
+install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib)
+install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
 install(FILES @hip_SOURCE_DIR@/src/hip_ir.ll DESTINATION lib)
 

From ff2f54c1bf50ffdd1dfb121ef93681fe8f517063 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Fri, 2 Dec 2016 23:57:22 +0000
Subject: [PATCH 08/28] Add additional controls for forcing serialization and
 blocking.

Move HIP_COHERENT_HOST_ALLOC so it is read once at init time.
Add HIP_LAUNCH_BLOCKING_KERNELS, HIP_API_BLOCKING.
Update docs on debug and chicken bits.

Conflicts:
	src/hip_hcc.cpp
---
 docs/markdown/hip_profiling.md       | 13 +++++--
 include/hip/hcc_detail/hip_runtime.h |  4 +--
 src/hip_hcc.cpp                      | 51 ++++++++++++++++++++++------
 src/hip_hcc.h                        | 10 +++---
 src/hip_memory.cpp                   | 30 ++++------------
 src/hip_module.cpp                   |  2 +-
 6 files changed, 65 insertions(+), 45 deletions(-)

diff --git a/docs/markdown/hip_profiling.md b/docs/markdown/hip_profiling.md
index db5d0fc425..21133100ec 100644
--- a/docs/markdown/hip_profiling.md
+++ b/docs/markdown/hip_profiling.md
@@ -325,9 +325,18 @@ Some key information from the trace above.
 Chicken bits are environment variables which cause the HIP, HCC, or HSA driver to disable some feature or optimization.
 These are not intended for production but can be useful diagnose synchronization problems in the application (or driver).
 
-Some of the most useful chicken bits are described here:
+Some of the most useful chicken bits are described here. These bits are supported on the ROCm path:
+
+HIP provides 3 environment variables in the HIP_*_BLOCKING family.  These introduce additional synchronization and can be useful to isolate synchronization problems. Specifically, if the code works with this flag set, then it indicates the kernels are executing correctly, and any failures likely are causes by improper or missing synchronization.  These flags will have performance impact and are not intended for production use.
+
+- HIP_LAUNCH_BLOCKING=1 : Waits on the host after each kernel launch.  Equivalent to setting CUDA_LAUNCH_BLOCKING.
+- HIP_LAUNCH_BLOCKING_KERNELS: A comma-separated list of kernel names.  The HIP runtime will wait on the host after one of the named kernels executes.  This provides a more targeted version of HIP_LAUNCH_BLOCKING and may be useful to isolate exactly which kernel needs further analysis if HIP_LAUNCH_BLOCKING=1 improves functionality.  There is no indication if kernel names are spelled incorrectly.  One mechanism to verify that the blocking is working is to run with HIP_DB=api+sync and search for debug messages with "LAUNCH_BLOCKING".
+- HIP_API_BLOCKING : Forces hipMemcpyAsync and hipMemsetAsync to be host-synchronous, meaning they will wait for the requested operation to complete before returning to the caller.
+
+These options cause HCC to serialize.  Useful if you have libraries or code which is calling HCC kernels directly rather than using HIP.  
+- HCC_SERIALZIE_KERNELS : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize.
+- HCC_SERIALIZE_COPY    : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize.0
 
-- HIP_LAUNCH_BLOCKING=1 : On ROCm, this flag waits on the host after each kernel launches and after each memory copy command.  On CUDA, the waits are only enforced after each kernel launch.  This is useful to isolate synchronization problems.  Specifically, if the code works with this flag set, then it indicates the kernels and memory management code are correct, and any failures likely are causes by improper or missing synchronization.
 - HSA_ENABLE_SDMA=0     : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines.  Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine.  This flag is useful to isolate issues with the hardware copy engines.
 - HSA_ENABLE_INTERRUPT=0 : Causes completion signals to be detected with memory-based polling rather than interrupts.  Can be useful to diagnose interrupt storm issues in the driver.
 - HSA_DISABLE_CACHE=1  : Disables the GPU L2 data cache.
diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index b1877ed0b3..55a0485365 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -828,7 +828,7 @@ extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block
 extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
 extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm *lp, const char *kernelNameStr);
 extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm *lp, const char *kernelNameStr);
-extern void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
+extern void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp);
 
 
 // Due to multiple overloaded versions of ihipPreLaunchKernel, the numBlocks3D and blockDim3D can be either size_t or dim3 types
@@ -838,7 +838,7 @@ do {\
   lp.dynamic_group_mem_bytes = _groupMemBytes; \
   hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp, #_kernelName)); \
   _kernelName (lp, ##__VA_ARGS__);\
-  ihipPostLaunchKernel(trueStream, lp);\
+  ihipPostLaunchKernel(#_kernelName, trueStream, lp);\
 } while(0)
 
 
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 06402f9a67..abd260762f 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -61,6 +61,9 @@ const char *API_COLOR = KGRN;
 const char *API_COLOR_END = KNRM;
 
 int HIP_LAUNCH_BLOCKING = 0;
+std::string HIP_LAUNCH_BLOCKING_KERNELS;
+std::vector<std::string> g_hipLaunchBlockingKernels;
+int HIP_API_BLOCKING = 0;
 
 int HIP_PRINT_ENV = 0;
 int HIP_TRACE_API= 0;
@@ -81,6 +84,8 @@ int HIP_DENY_PEER_ACCESS = 0;
 // Force async copies to actually use the synchronous copy interface.
 int HIP_FORCE_SYNC_COPY = 0;
 
+int HIP_COHERENT_HOST_ALLOC = 0;
+
 
 
 
@@ -358,13 +363,24 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
 
 //---
 // Must be called after kernel finishes, this releases the lock on the stream so other commands can submit.
-void ihipStream_t::lockclose_postKernelCommand(hc::accelerator_view *av)
+void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::accelerator_view *av)
 {
+    bool blockThisKernel = false;
 
-    if (HIP_LAUNCH_BLOCKING) {
+    if (!g_hipLaunchBlockingKernels.empty()) {
+        std::string kernelNameString(kernelName);
+        for (auto o=g_hipLaunchBlockingKernels.begin(); o!=g_hipLaunchBlockingKernels.end(); o++) {
+            if ((*o == kernelNameString)) {
+                //printf ("force blocking for kernel %s\n", o->c_str());
+                blockThisKernel = true;
+            }
+        }
+    }
+
+    if (HIP_LAUNCH_BLOCKING || blockThisKernel) {
         // TODO - fix this so it goes through proper stream::wait() call.// direct wait OK since we know the stream is locked.
         av->wait(hc::hcWaitModeActive);
-        tprintf(DB_SYNC, " %s LAUNCH_BLOCKING for kernel completion\n", ToString(this).c_str());
+        tprintf(DB_SYNC, "%s LAUNCH_BLOCKING for kernel '%s' completion\n", ToString(this).c_str(), kernelName);
     }
 
     _criticalData.unlock(); // paired with lock from lockopen_preKernelCommand.
@@ -1243,7 +1259,15 @@ void ihipInit()
     //-- READ HIP_PRINT_ENV env first, since it has impact on later env var reading
 
     // 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_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP kernel launches 'host-synchronous', so they block until any kernel launches. Alias: CUDA_LAUNCH_BLOCKING." );
+    READ_ENV_S(release, HIP_LAUNCH_BLOCKING_KERNELS, 0, "Comma-separated list of kernel names to make host-synchronous, so they block until completed.");
+    if (!HIP_LAUNCH_BLOCKING_KERNELS.empty()) {
+        tokenize(HIP_LAUNCH_BLOCKING_KERNELS, ',', &g_hipLaunchBlockingKernels);
+    }
+    READ_ENV_I(release, HIP_API_BLOCKING, 0, "Make HIP APIs 'host-synchronous', so they block until completed.  Impacts hipMemcpyAsync, hipMemsetAsync." );
+    
+
+
     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.
@@ -1252,6 +1276,8 @@ 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_PROFILE_API, 0,  "Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, 0x2=full API name including args.");
@@ -1262,10 +1288,14 @@ void ihipInit()
 
 
     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_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"); 
+
+    // TODO - review, can we remove this?
     READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Max number of inflight kernels per stream before active synchronization is forced.");
 
+    READ_ENV_I(release, HIP_COHERENT_HOST_ALLOC, 0, "If set, all host memory will be allocated as fine-grained system memory.  This allows threadfence_system to work but prevents host memory from being cached on GPU which may have performance impact.");
+
     // 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?)\n", HIP_DB);
@@ -1398,6 +1428,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
 
 void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
 {
+
     if (HIP_PROFILE_API || (COMPILE_HIP_DB && HIP_TRACE_API)) {
         std::stringstream os_pre;
         std::stringstream os;
@@ -1515,11 +1546,11 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
 //---
 //Called after kernel finishes execution.
 //This releases the lock on the stream.
-void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp)
+void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp)
 {
     tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
 
-    stream->lockclose_postKernelCommand(lp.av);
+    stream->lockclose_postKernelCommand(kernelName, lp.av);
     MARKER_END();
 }
 
@@ -1883,8 +1914,8 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
             };
 
 
-            if (HIP_LAUNCH_BLOCKING) {
-                tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
+            if (HIP_API_BLOCKING) {
+                tprintf(DB_SYNC, "%s LAUNCH_BLOCKING for completion of hipMemcpyAsync(sz=%zu)\n", ToString(this).c_str(), sizeBytes);
                 this->wait(crit);
             }
 
diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index 0040263194..b01d41be14 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -45,6 +45,7 @@ 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_API_BLOCKING;
 
 extern int HIP_PRINT_ENV;
 extern int HIP_PROFILE_API;
@@ -56,6 +57,8 @@ extern int HIP_STREAM_SIGNALS;  /* number of signals to allocate at stream creat
 extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU identifiers */
 extern int HIP_FORCE_P2P_HOST;
 
+extern int HIP_COHERENT_HOST_ALLOC;
+
 
 //---
 // Chicken bits for disabling functionality to work around potential issues:
@@ -156,11 +159,6 @@ extern const char *API_COLOR_END;
 #endif
 
 
-// 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.
@@ -455,7 +453,7 @@ public:
     //---
     // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.
     LockedAccessor_StreamCrit_t  lockopen_preKernelCommand();
-    void                 lockclose_postKernelCommand(hc::accelerator_view *av);
+    void                 lockclose_postKernelCommand(const char *kernelName, hc::accelerator_view *av);
 
 
     void                 locked_wait(bool assertQueueEmpty=false);
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 314890d167..4b91228032 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -157,20 +157,6 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
     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)
 {
@@ -193,16 +179,12 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 
         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){
+            if(HIP_COHERENT_HOST_ALLOC){
                 // Force to allocate finedgrained system memory
                 *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
                 if(sizeBytes < 1 && (*ptr == NULL)){
@@ -853,13 +835,13 @@ hipError_t hipMemsetAsync(void* dst, int  value, size_t sizeBytes, hipStream_t s
             }
         }
 
-        stream->lockclose_postKernelCommand(&crit->_av);
+        stream->lockclose_postKernelCommand("hipMemsetAsync", &crit->_av);
 
 
-        if (HIP_LAUNCH_BLOCKING) {
-            tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset [stream:%p].\n", __func__, (void*)stream);
+        if (HIP_API_BLOCKING) {
+            tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str());
             cf.wait();
-            tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
+            //tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
         }
     } else {
         e = hipErrorInvalidValue;
@@ -906,7 +888,7 @@ hipError_t hipMemset(void* dst, int  value, size_t sizeBytes )
         // TODO - is hipMemset supposed to be async?
         cf.wait();
 
-        stream->lockclose_postKernelCommand(&crit->_av);
+        stream->lockclose_postKernelCommand("hipMemset", &crit->_av);
 
 
         if (HIP_LAUNCH_BLOCKING) {
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index f7ac35c77b..606d99f2fd 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -333,7 +333,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 #endif // USE_DISPATCH_HSA_KERNEL
 
 
-        ihipPostLaunchKernel(hStream, lp);
+        ihipPostLaunchKernel(f->_kernelName, hStream, lp);
 
     }
 

From 097e4eb9d8a003f4dc00814b4c437cf9896df40e Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sun, 4 Dec 2016 00:13:19 -0600
Subject: [PATCH 09/28] Enable USE_DISPATCH_HSA_KERNEL.

Optimize hipLaunchModule dispatch latency.
---
 src/hip_hcc.cpp    | 55 ----------------------------------------------
 src/hip_hcc.h      |  2 --
 src/hip_module.cpp | 27 -----------------------
 3 files changed, 84 deletions(-)

diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index abd260762f..223e7a9243 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -389,61 +389,6 @@ void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::acce
 
 
 
-#if USE_DISPATCH_HSA_KERNEL==0
-// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av.
-void ihipStream_t::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 privateSegmentSize,
-                        void *kernarg,
-                        size_t kernSize,
-                        uint64_t kernel){
-    hsa_status_t status;
-    void *kern;
-
-    hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(av.get_hsa_kernarg_region());
-    status = hsa_amd_memory_pool_allocate(*pool, kernSize, 0, &kern);
-    status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)av.get_hsa_agent(), 0, kern);
-    memcpy(kern, kernarg, kernSize);
-    hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue();
-    const uint32_t queue_mask = Queue->size-1;
-    uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue);
-    hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]);
-
-    dispatch_packet->completion_signal = signal;
-    dispatch_packet->workgroup_size_x = blockDimX;
-    dispatch_packet->workgroup_size_y = blockDimY;
-    dispatch_packet->workgroup_size_z = blockDimZ;
-    dispatch_packet->grid_size_x = blockDimX * gridDimX;
-    dispatch_packet->grid_size_y = blockDimY * gridDimY;
-    dispatch_packet->grid_size_z = blockDimZ * gridDimZ;
-    dispatch_packet->group_segment_size = groupSegmentSize;
-    dispatch_packet->private_segment_size = privateSegmentSize;
-    dispatch_packet->kernarg_address = kern;
-    dispatch_packet->kernel_object = kernel;
-    uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
-    (1 << HSA_PACKET_HEADER_BARRIER) |
-    (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
-    (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
-
-    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);
-
-    hsa_queue_store_write_index_relaxed(Queue, packet_index + 1);
-    hsa_signal_store_relaxed(Queue->doorbell_signal, packet_index);
-}
-#endif
-
-
 //=============================================================================
 // Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted.
 // The packed _peerAgents can efficiently be used on each memory allocation.
diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index b01d41be14..5ab51b5ea3 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -32,8 +32,6 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif
 
-#define USE_DISPATCH_HSA_KERNEL 0
-//
 
 
 //---
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 606d99f2fd..2daa251004 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -282,8 +282,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         grid_launch_parm lp;
         hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f->_kernelName);
 
-#if USE_DISPATCH_HSA_KERNEL
-
         hsa_kernel_dispatch_packet_t aql;
 
         memset(&aql, 0, sizeof(aql));
@@ -307,31 +305,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
                                     (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
 
         lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize);
-#else
-
-        /*
-          Create signal
-        */
-
-        hsa_signal_t signal;
-        status = hsa_signal_create(1, 0, NULL, &signal);
-
-
-        /*
-          Launch AQL packet
-        */
-        hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
-                  gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, config[1], kernArgSize, f->_kernel);
-
-
-        /*
-          Wait for signal
-        */
-
-        hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
-
-#endif // USE_DISPATCH_HSA_KERNEL
-
 
         ihipPostLaunchKernel(f->_kernelName, hStream, lp);
 

From 1cf9332c3f5a6a6c19301f0cdd9418c107be64cc Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Mon, 5 Dec 2016 11:03:45 +0530
Subject: [PATCH 10/28] Don't share g_malloc_heap_size between libraries

Change-Id: Ic70bf83d4f865bc5c453941fdbc1814c77f0ad9d
---
 src/device_util.cpp | 10 +---------
 src/device_util.h   | 38 ++++++++++++++++++++++++++++++++++++++
 src/hip_device.cpp  |  5 ++---
 3 files changed, 41 insertions(+), 12 deletions(-)
 create mode 100644 src/device_util.h

diff --git a/src/device_util.cpp b/src/device_util.cpp
index 203e7a7826..19a23e419b 100644
--- a/src/device_util.cpp
+++ b/src/device_util.cpp
@@ -23,6 +23,7 @@ THE SOFTWARE.
 #include <hc.hpp>
 #include <grid_launch.h>
 #include <hc_math.hpp>
+#include "device_util.h"
 
 #include "hip/hip_runtime.h"
 
@@ -33,15 +34,6 @@ THE SOFTWARE.
  This is the best place to put them because the device
  global variables need to be initialized at the start.
 */
-
-#define NUM_PAGES_PER_THREAD  16
-#define SIZE_OF_PAGE          64
-#define NUM_THREADS_PER_CU    64
-#define NUM_CUS_PER_GPU       64
-#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
-#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
-#define SIZE_OF_HEAP SIZE_MALLOC
-
 size_t g_malloc_heap_size = SIZE_OF_HEAP;
 
 __attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
diff --git a/src/device_util.h b/src/device_util.h
new file mode 100644
index 0000000000..8ccf5a540e
--- /dev/null
+++ b/src/device_util.h
@@ -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.
+*/
+
+#ifndef DEVICE_UTIL_H
+#define DEVICE_UTIL_H
+
+/*
+ Heap size computation for malloc and free device functions.
+*/
+
+#define NUM_PAGES_PER_THREAD  16
+#define SIZE_OF_PAGE          64
+#define NUM_THREADS_PER_CU    64
+#define NUM_CUS_PER_GPU       64
+#define NUM_PAGES NUM_PAGES_PER_THREAD * NUM_THREADS_PER_CU * NUM_CUS_PER_GPU
+#define SIZE_MALLOC NUM_PAGES * SIZE_OF_PAGE
+#define SIZE_OF_HEAP SIZE_MALLOC
+
+#endif
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 1275c8f19b..bc8a2c297f 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -23,6 +23,7 @@ THE SOFTWARE.
 #include "hip/hip_runtime.h"
 #include "hip_hcc.h"
 #include "trace_helper.h"
+#include "device_util.h"
 
 //-------------------------------------------------------------------------------------------------
 //Devices
@@ -97,8 +98,6 @@ hipError_t hipDeviceGetCacheConfig(hipFuncCache_t *cacheConfig)
     return ihipLogStatus(hipSuccess);
 }
 
-extern "C" size_t g_malloc_heap_size;
-
 hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit_t limit)
 {
     HIP_INIT_API(pValue, limit);
@@ -106,7 +105,7 @@ hipError_t hipDeviceGetLimit (size_t *pValue, hipLimit_t limit)
         return ihipLogStatus(hipErrorInvalidValue);
     }
     if(limit == hipLimitMallocHeapSize) {
-        *pValue = g_malloc_heap_size;
+        *pValue = (size_t)SIZE_OF_HEAP;
         return ihipLogStatus(hipSuccess);
     }else{
         return ihipLogStatus(hipErrorUnsupportedLimit);

From 46ffc69557f2d5357fba2d8aea9a384742f9c6d7 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Mon, 5 Dec 2016 16:55:26 +0530
Subject: [PATCH 11/28] Revert "Enable USE_DISPATCH_HSA_KERNEL."

This reverts commit 097e4eb9d8a003f4dc00814b4c437cf9896df40e.
---
 src/hip_hcc.cpp    | 55 ++++++++++++++++++++++++++++++++++++++++++++++
 src/hip_hcc.h      |  2 ++
 src/hip_module.cpp | 27 +++++++++++++++++++++++
 3 files changed, 84 insertions(+)

diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 223e7a9243..abd260762f 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -389,6 +389,61 @@ void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::acce
 
 
 
+#if USE_DISPATCH_HSA_KERNEL==0
+// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av.
+void ihipStream_t::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 privateSegmentSize,
+                        void *kernarg,
+                        size_t kernSize,
+                        uint64_t kernel){
+    hsa_status_t status;
+    void *kern;
+
+    hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(av.get_hsa_kernarg_region());
+    status = hsa_amd_memory_pool_allocate(*pool, kernSize, 0, &kern);
+    status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)av.get_hsa_agent(), 0, kern);
+    memcpy(kern, kernarg, kernSize);
+    hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue();
+    const uint32_t queue_mask = Queue->size-1;
+    uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue);
+    hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]);
+
+    dispatch_packet->completion_signal = signal;
+    dispatch_packet->workgroup_size_x = blockDimX;
+    dispatch_packet->workgroup_size_y = blockDimY;
+    dispatch_packet->workgroup_size_z = blockDimZ;
+    dispatch_packet->grid_size_x = blockDimX * gridDimX;
+    dispatch_packet->grid_size_y = blockDimY * gridDimY;
+    dispatch_packet->grid_size_z = blockDimZ * gridDimZ;
+    dispatch_packet->group_segment_size = groupSegmentSize;
+    dispatch_packet->private_segment_size = privateSegmentSize;
+    dispatch_packet->kernarg_address = kern;
+    dispatch_packet->kernel_object = kernel;
+    uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+    (1 << HSA_PACKET_HEADER_BARRIER) |
+    (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+    (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
+
+    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);
+
+    hsa_queue_store_write_index_relaxed(Queue, packet_index + 1);
+    hsa_signal_store_relaxed(Queue->doorbell_signal, packet_index);
+}
+#endif
+
+
 //=============================================================================
 // Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted.
 // The packed _peerAgents can efficiently be used on each memory allocation.
diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index 5ab51b5ea3..b01d41be14 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -32,6 +32,8 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif
 
+#define USE_DISPATCH_HSA_KERNEL 0
+//
 
 
 //---
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 2daa251004..606d99f2fd 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -282,6 +282,8 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         grid_launch_parm lp;
         hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f->_kernelName);
 
+#if USE_DISPATCH_HSA_KERNEL
+
         hsa_kernel_dispatch_packet_t aql;
 
         memset(&aql, 0, sizeof(aql));
@@ -305,6 +307,31 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
                                     (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
 
         lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize);
+#else
+
+        /*
+          Create signal
+        */
+
+        hsa_signal_t signal;
+        status = hsa_signal_create(1, 0, NULL, &signal);
+
+
+        /*
+          Launch AQL packet
+        */
+        hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
+                  gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, config[1], kernArgSize, f->_kernel);
+
+
+        /*
+          Wait for signal
+        */
+
+        hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
+
+#endif // USE_DISPATCH_HSA_KERNEL
+
 
         ihipPostLaunchKernel(f->_kernelName, hStream, lp);
 

From 778c6626fdf5791b2442fe47079584075d0837fa Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Mon, 5 Dec 2016 20:21:33 -0600
Subject: [PATCH 12/28] HIP resource leaks fix from Jack

Change-Id: I93f3ad7cb94ff1cba1577bd8acc90e826693d12e
---
 src/hip_hcc.cpp    | 10 ++-------
 src/hip_hcc.h      | 34 ++++++++++++++++++++----------
 src/hip_module.cpp | 52 +++++++++++++++++++++++++++++++++++++++++++---
 3 files changed, 74 insertions(+), 22 deletions(-)

diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index abd260762f..84794b02fb 100644
--- a/src/hip_hcc.cpp
+++ b/src/hip_hcc.cpp
@@ -401,17 +401,11 @@ void ihipStream_t::launchModuleKernel(
                         uint32_t gridDimY,
                         uint32_t gridDimZ,
                         uint32_t groupSegmentSize,
-						uint32_t privateSegmentSize,
+			uint32_t privateSegmentSize,
                         void *kernarg,
                         size_t kernSize,
                         uint64_t kernel){
     hsa_status_t status;
-    void *kern;
-
-    hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(av.get_hsa_kernarg_region());
-    status = hsa_amd_memory_pool_allocate(*pool, kernSize, 0, &kern);
-    status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)av.get_hsa_agent(), 0, kern);
-    memcpy(kern, kernarg, kernSize);
     hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue();
     const uint32_t queue_mask = Queue->size-1;
     uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue);
@@ -426,7 +420,7 @@ void ihipStream_t::launchModuleKernel(
     dispatch_packet->grid_size_z = blockDimZ * gridDimZ;
     dispatch_packet->group_segment_size = groupSegmentSize;
     dispatch_packet->private_segment_size = privateSegmentSize;
-    dispatch_packet->kernarg_address = kern;
+    dispatch_packet->kernarg_address = kernarg;
     dispatch_packet->kernel_object = kernel;
     uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
     (1 << HSA_PACKET_HEADER_BARRIER) |
diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index b01d41be14..3074a4f121 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -367,17 +367,6 @@ struct LockedBase {
     MUTEX_TYPE  _mutex;
 };
 
-
-class ihipModule_t{
-public:
-  hsa_executable_t executable;
-  hsa_code_object_t object;
-  std::string fileName;
-  void *ptr;
-  size_t size;
-};
-
-
 class ihipFunction_t{
 public:
     ihipFunction_t(const char *name) {
@@ -399,6 +388,29 @@ public:
     uint64_t _kernel;
 };
 
+class ihipModule_t {
+public:
+  hsa_executable_t executable;
+  hsa_code_object_t object;
+  std::string fileName;
+  void *ptr;
+  size_t size;
+
+  ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0), hipFunctionTable() {}
+  ~ihipModule_t() {
+    for (int i = 0; i < hipFunctionTable.size(); ++i) {
+      ihipFunction_t *func = hipFunctionTable[i];
+      delete func;
+    }
+    hipFunctionTable.clear();
+  }
+
+  void registerFunction(ihipFunction_t* func) {
+    hipFunctionTable.push_back(func);
+  }
+private:
+  std::vector<ihipFunction_t*> hipFunctionTable;
+};
 
 template <typename MUTEX_TYPE>
 class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index 606d99f2fd..d58a476f77 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -126,7 +126,6 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
 
         } else {
 
-            *module = new ihipModule_t;
             size_t size = std::string::size_type(in.tellg());
             void *p = NULL;
             hsa_agent_t agent = currentDevice->_hsaAgent;
@@ -172,6 +171,11 @@ hipError_t hipModuleUnload(hipModule_t hmod){
 				ret = hipErrorInvalidValue;
 		}
     status = hsa_code_object_destroy(hmod->object);
+    if(status != HSA_STATUS_SUCCESS)
+		{
+				ret = hipErrorInvalidValue;
+		}
+    status = hsa_memory_free(hmod->ptr);
     if(status != HSA_STATUS_SUCCESS)
 		{
 				ret = hipErrorInvalidValue;
@@ -193,6 +197,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
 
     }else{
         *func = new ihipFunction_t(name);
+        hmod->registerFunction(*func);
         int deviceId = ctx->getDevice()->_deviceId;
         ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
         hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
@@ -268,11 +273,17 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol,
                                                              HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
                                                              &groupSegmentSize);
+        if(status != HSA_STATUS_SUCCESS){
+            return ihipLogStatus(hipErrorNotFound);
+        }
 
         uint32_t privateSegmentSize;
         status = hsa_executable_symbol_get_info(f->_kernelSymbol,
                                                 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
                                                 &privateSegmentSize);
+        if(status != HSA_STATUS_SUCCESS){
+            return ihipLogStatus(hipErrorNotFound);
+        }
         privateSegmentSize += sharedMemBytes;
 
 
@@ -300,7 +311,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         aql.group_segment_size = groupSegmentSize;
         aql.private_segment_size = privateSegmentSize;
         aql.kernel_object = f->_kernel;
-        aql.setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
+        aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
         aql.header =   (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
                                     (1 << HSA_PACKET_HEADER_BARRIER) |
                                     (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
@@ -315,13 +326,32 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 
         hsa_signal_t signal;
         status = hsa_signal_create(1, 0, NULL, &signal);
+        if(status != HSA_STATUS_SUCCESS){
+            return ihipLogStatus(hipErrorNotFound);
+        }
+
+        /*
+           Allocate kernarg
+        */
+        void *kern = nullptr;
+
+        hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(lp.av->get_hsa_kernarg_region());
+        status = hsa_amd_memory_pool_allocate(*pool, kernArgSize, 0, &kern);
+        if(status != HSA_STATUS_SUCCESS){
+            return ihipLogStatus(hipErrorNotFound);
+        }
+        status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)lp.av->get_hsa_agent(), 0, kern);
+        if(status != HSA_STATUS_SUCCESS){
+            return ihipLogStatus(hipErrorNotFound);
+        }
+        memcpy(kern, config[1], kernArgSize);
 
 
         /*
           Launch AQL packet
         */
         hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
-                  gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, config[1], kernArgSize, f->_kernel);
+                  gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_kernel);
 
 
         /*
@@ -330,6 +360,22 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 
         hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
 
+        /*
+           Destroy kernarg
+        */
+        status = hsa_amd_memory_pool_free(kern);
+        if(status != HSA_STATUS_SUCCESS){
+            return ihipLogStatus(hipErrorNotFound);
+        }
+
+        /*
+          Destroy the signal
+        */
+        status = hsa_signal_destroy(signal);
+        if(status != HSA_STATUS_SUCCESS){
+            return ihipLogStatus(hipErrorNotFound);
+        }
+
 #endif // USE_DISPATCH_HSA_KERNEL
 
 

From 6d5145eba47189282bbb74109e27a642b06070c1 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Tue, 6 Dec 2016 10:19:03 +0530
Subject: [PATCH 13/28] Remove redundant variable g_malloc_heap_size

Change-Id: Idaf47be70488f0deb3eab05a86d9c5a413d3fff7
---
 src/device_util.cpp | 2 --
 1 file changed, 2 deletions(-)

diff --git a/src/device_util.cpp b/src/device_util.cpp
index 19a23e419b..7b881751ae 100644
--- a/src/device_util.cpp
+++ b/src/device_util.cpp
@@ -34,8 +34,6 @@ THE SOFTWARE.
  This is the best place to put them because the device
  global variables need to be initialized at the start.
 */
-size_t g_malloc_heap_size = SIZE_OF_HEAP;
-
 __attribute__((address_space(1))) char gpuHeap[SIZE_OF_HEAP];
 __attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
 

From 27072b897210810d01534f06e8c5d9b0e15d099e Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Tue, 6 Dec 2016 16:55:17 +0530
Subject: [PATCH 14/28] Changed hipDeviceGetPCIBusId to return Bus ID as string

Change-Id: I6d5aa7362084109d34bc015d948f8723b2a38ee9
---
 src/hip_device.cpp | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index bc8a2c297f..085314d461 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -342,7 +342,12 @@ hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
     HIP_INIT_API(pciBusId, len, device);
     hipError_t e = hipSuccess;
     int deviceId= device->_deviceId;
-    e = ihipDeviceGetAttribute((int*)pciBusId, hipDeviceAttributePciBusId, deviceId);
+    int tempPciBusId = 0;
+    e = ihipDeviceGetAttribute( tempPciBusId, hipDeviceAttributePciBusId, deviceId);
+    if( e == hipSuccess) {
+        std::string tempPciStr = std::to_string(tempPciBusId);
+        memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );
+    }
     return ihipLogStatus(e);
 }
 

From ca06747e1faa77818c09a984d902cbf888c41aa3 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Tue, 6 Dec 2016 17:09:21 +0530
Subject: [PATCH 15/28] Build Error correction in hipDeviceGetPCIBusId

Change-Id: I50ff4d95b7a732924c7a991cba60400b1c93c0de
---
 src/hip_device.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 085314d461..868dd0b05d 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -343,7 +343,7 @@ hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
     hipError_t e = hipSuccess;
     int deviceId= device->_deviceId;
     int tempPciBusId = 0;
-    e = ihipDeviceGetAttribute( tempPciBusId, hipDeviceAttributePciBusId, deviceId);
+    e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, deviceId);
     if( e == hipSuccess) {
         std::string tempPciStr = std::to_string(tempPciBusId);
         memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );

From 9aebbe00def4ada9fab497c79604c048fe7cd7e1 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Tue, 6 Dec 2016 17:31:54 +0530
Subject: [PATCH 16/28] Simple test case for hipDeviceGetPCIBusId

Change-Id: I0fb6d1bef4739a5b6280928f7b349d95c1656431
---
 tests/src/hipDrvGetPCIBusId.cpp | 34 +++++++++++++++++++++++++++++++++
 1 file changed, 34 insertions(+)
 create mode 100644 tests/src/hipDrvGetPCIBusId.cpp

diff --git a/tests/src/hipDrvGetPCIBusId.cpp b/tests/src/hipDrvGetPCIBusId.cpp
new file mode 100644
index 0000000000..21c49c194c
--- /dev/null
+++ b/tests/src/hipDrvGetPCIBusId.cpp
@@ -0,0 +1,34 @@
+/*
+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 WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+#include "hip/hip_runtime_api.h"
+#include<iostream>
+#include<fstream>
+#include<vector>
+
+int main(){
+    hipInit(0);
+	hipDevice_t device;
+	hipDeviceGet(&device, 0);
+    char pciBusId[100];
+    hipDeviceGetPCIBusId(pciBusId,100,device);
+    printf("PCI Bus ID= %s\n",pciBusId);
+    return 0;
+}

From 6209565ec05363205f2dd07ff86de31df36bdf28 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Tue, 6 Dec 2016 16:05:46 +0000
Subject: [PATCH 17/28] Enabled USE_DISPATCH_HSA_KERNEL, with serialization in
 hipModuleUnload.

---
 src/hip_hcc.h      | 2 +-
 src/hip_module.cpp | 7 ++++++-
 2 files changed, 7 insertions(+), 2 deletions(-)

diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index 3074a4f121..c4b092aa5f 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -32,7 +32,7 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif
 
-#define USE_DISPATCH_HSA_KERNEL 0
+#define USE_DISPATCH_HSA_KERNEL 1
 //
 
 
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index d58a476f77..badeac2dcd 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -163,7 +163,12 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
     return ihipLogStatus(ret);
 }
 
-hipError_t hipModuleUnload(hipModule_t hmod){
+hipError_t hipModuleUnload(hipModule_t hmod)
+{
+    // TODO - improve this synchronization so it is thread-safe.
+    // Currently we want for all inflight activity to complete, but don't prevent another
+    // thread from launching new kernels before we finish this operation.
+    ihipSynchronize();
     hipError_t ret = hipSuccess;
     hsa_status_t status = hsa_executable_destroy(hmod->executable);
     if(status != HSA_STATUS_SUCCESS)

From eba2595611467b336d36cde42f194c4af59d49ef Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Thu, 1 Dec 2016 13:51:58 -0600
Subject: [PATCH 18/28] local changes for hipnccl

Change-Id: I05a1f0381ce2914a800f573342cc954eb5ff82d9
---
 include/hip/hcc_detail/hip_runtime_api.h | 44 ++++++++++++++++++++++--
 src/hip_device.cpp                       | 14 ++++++--
 src/hip_memory.cpp                       | 15 ++++++++
 3 files changed, 68 insertions(+), 5 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 83c60f3e35..3793cfdf8c 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -57,6 +57,30 @@ typedef struct ihipDevice_t *hipDevice_t;
 
 typedef struct ihipStream_t *hipStream_t;
 
+//TODO: IPC implementation
+
+#define hipIpcMemLazyEnablePeerAccess 0
+struct ihipIpcMemHandle_t;
+typedef struct ihipIpcMemHandle_t *hipIpcMemHandle_t;
+struct ihipIpcEventHandle_t;
+typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;
+
+typedef std::nullptr_t nullptr_t ;
+
+__device__ double  
+__longlong_as_double(long long int x)
+{
+  return (double)x;
+}
+__device__ long long int    
+__double_as_longlong(double x)
+{
+  return (long long int)x; 
+}
+
+
+//END TODO
+
 typedef struct ihipModule_t *hipModule_t;
 
 typedef struct ihipFunction_t *hipFunction_t;
@@ -1605,11 +1629,23 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device);
  * @brief Returns a PCI Bus Id string for the device.
  * @param [out] pciBusId
  * @param [in] len
+ * @param [hipDevice_t] device
+ *
+ * @returns #hipSuccess, #hipErrorInavlidDevice
+ */
+// hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
+
+
+/**
+ * @brief Returns a PCI Bus Id string for the device, overloaded to take int device ID.
+ * @param [out] pciBusId
+ * @param [in] len
  * @param [in] device
  *
  * @returns #hipSuccess, #hipErrorInavlidDevice
  */
-hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,int device);
+
 
 /**
  * @brief Returns a handle to a compute device.
@@ -1791,7 +1827,11 @@ hipError_t hipProfilerStop();
  * @}
  */
 
-
+//TODO: implement IPC apis
+hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
+hipError_t hipIpcCloseMemHandle(void *devPtr);
+hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
+hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
 
 
 #ifdef __cplusplus
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 868dd0b05d..38021ad525 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -337,13 +337,21 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return ihipLogStatus(e);
 }
 
-hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
+// hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
+// {
+//     HIP_INIT_API(pciBusId, len, device);
+//     hipError_t e = hipSuccess;
+//     int deviceId= device->_deviceId;
+//     e = ihipDeviceGetAttribute((int*)pciBusId, hipDeviceAttributePciBusId, deviceId);
+//     return ihipLogStatus(e);
+// }
+
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,int device)
 {
     HIP_INIT_API(pciBusId, len, device);
     hipError_t e = hipSuccess;
-    int deviceId= device->_deviceId;
     int tempPciBusId = 0;
-    e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, deviceId);
+    e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, device);
     if( e == hipSuccess) {
         std::string tempPciStr = std::to_string(tempPciBusId);
         memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 4b91228032..b565c5a770 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1034,3 +1034,18 @@ hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDevi
     hipStatus = hipErrorInvalidDevicePointer;
     return hipStatus;
 }
+
+
+//TODO: IPC implementaiton:
+hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
+    return hipSuccess;
+}
+hipError_t hipIpcCloseMemHandle(void *devPtr){
+    return hipSuccess;
+}
+hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
+    return hipSuccess;
+}
+hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
+    return hipSuccess;
+}

From 17b98d59b885306b24ba0aa6f746240c722df6e0 Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Tue, 6 Dec 2016 14:09:53 -0600
Subject: [PATCH 19/28] IPC prototyps and part of the implementation included

Change-Id: Id88c7f155d23ec63f57a6ef05098fba43f8af336
---
 include/hip/hcc_detail/hip_runtime.h      |   3 +-
 include/hip/hcc_detail/hip_runtime_api.h  | 114 ++++++++++++++++++----
 include/hip/hcc_detail/hip_vector_types.h |   4 +
 src/device_util.cpp                       |  18 +++-
 src/hip_hcc.h                             |  27 ++++-
 src/hip_memory.cpp                        |  29 ++++--
 6 files changed, 159 insertions(+), 36 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 55a0485365..78accc0c5b 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -771,7 +771,8 @@ extern "C" __device__ void __threadfence(void);
  *
  * @warning __threadfence_system is a stub and map to no-op.
  */
-__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));
+//__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided with workaround configuration, see hip_kernel_language.md for details")));
+__device__ void  __threadfence_system(void) ;
 
 __device__ unsigned __hip_ds_bpermute(int index, unsigned src);
 __device__ float __hip_ds_bpermutef(int index, float src);
diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 3793cfdf8c..1af8108441 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -60,24 +60,13 @@ typedef struct ihipStream_t *hipStream_t;
 //TODO: IPC implementation
 
 #define hipIpcMemLazyEnablePeerAccess 0
-struct ihipIpcMemHandle_t;
-typedef struct ihipIpcMemHandle_t *hipIpcMemHandle_t;
+
+typedef struct ihipIpcMemHandle *hipIpcMemHandle_t;
+
+//TODO: IPC event handle currently unsupported
 struct ihipIpcEventHandle_t;
 typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;
 
-typedef std::nullptr_t nullptr_t ;
-
-__device__ double  
-__longlong_as_double(long long int x)
-{
-  return (double)x;
-}
-__device__ long long int    
-__double_as_longlong(double x)
-{
-  return (long long int)x; 
-}
-
 
 //END TODO
 
@@ -1828,10 +1817,97 @@ hipError_t hipProfilerStop();
  */
 
 //TODO: implement IPC apis
-hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
-hipError_t hipIpcCloseMemHandle(void *devPtr);
-hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
-hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
+
+/**
+ * @brief Gets an interprocess memory handle for an existing device memory
+ *          allocation
+ *
+ * Takes a pointer to the base of an existing device memory allocation created 
+ * with hipMalloc and exports it for use in another process. This is a 
+ * lightweight operation and may be called multiple times on an allocation
+ * without adverse effects. 
+ *
+ * If a region of memory is freed with hipFree and a subsequent call
+ * to hipMalloc returns memory with the same device address,
+ * hipIpcGetMemHandle will return a unique handle for the
+ * new memory. 
+ *
+ * @param handle - Pointer to user allocated hipIpcMemHandle to return
+ *                    the handle in.
+ * @param devPtr - Base pointer to previously allocated device memory 
+ *
+ * @returns
+ * hipSuccess,
+ * hipErrorInvalidResourceHandle,
+ * hipErrorMemoryAllocation,
+ * hipErrorMapBufferObjectFailed,
+ *
+ */
+extern __host__ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *devPtr);
+
+/**
+ * @brief Opens an interprocess memory handle exported from another process
+ *          and returns a device pointer usable in the local process.
+ *
+ * Maps memory exported from another process with hipIpcGetMemHandle into
+ * the current device address space. For contexts on different devices 
+ * hipIpcOpenMemHandle can attempt to enable peer access between the
+ * devices as if the user called hipDeviceEnablePeerAccess. This behavior is 
+ * controlled by the hipIpcMemLazyEnablePeerAccess flag. 
+ * hipDeviceCanAccessPeer can determine if a mapping is possible.
+ *
+ * Contexts that may open hipIpcMemHandles are restricted in the following way.
+ * hipIpcMemHandles from each device in a given process may only be opened 
+ * by one context per device per other process.
+ *
+ * Memory returned from hipIpcOpenMemHandle must be freed with
+ * hipIpcCloseMemHandle.
+ *
+ * Calling hipFree on an exported memory region before calling
+ * hipIpcCloseMemHandle in the importing context will result in undefined
+ * behavior.
+ * 
+ * @param devPtr - Returned device pointer
+ * @param handle - hipIpcMemHandle to open
+ * @param flags  - Flags for this operation. Must be specified as hipIpcMemLazyEnablePeerAccess
+ *
+ * @returns
+ * hipSuccess,
+ * hipErrorMapBufferObjectFailed,
+ * hipErrorInvalidResourceHandle,
+ * hipErrorTooManyPeers
+ *
+ * @note No guarantees are made about the address returned in @p *devPtr.  
+ * In particular, multiple processes may not receive the same address for the same @p handle.
+ *
+ */
+extern __host__ hipError_t hipIpcOpenMemHandle(void **devPtr, hipIpcMemHandle_t handle, unsigned int flags);
+
+/**
+ * @brief Close memory mapped with hipIpcOpenMemHandle
+ * 
+ * Unmaps memory returnd by hipIpcOpenMemHandle. The original allocation
+ * in the exporting process as well as imported mappings in other processes
+ * will be unaffected.
+ *
+ * Any resources used to enable peer access will be freed if this is the
+ * last mapping using them.
+ *
+ * @param devPtr - Device pointer returned by hipIpcOpenMemHandle
+ * 
+ * @returns
+ * hipSuccess,
+ * hipErrorMapBufferObjectFailed,
+ * hipErrorInvalidResourceHandle,
+ *
+ */
+extern __host__ hipError_t hipIpcCloseMemHandle(void *devPtr);
+
+
+// hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
+// hipError_t hipIpcCloseMemHandle(void *devPtr);
+// // hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
+// hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
 
 
 #ifdef __cplusplus
diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h
index 932e271527..ffe15a27a4 100644
--- a/include/hip/hcc_detail/hip_vector_types.h
+++ b/include/hip/hcc_detail/hip_vector_types.h
@@ -343,6 +343,10 @@ __HIP_DEVICE__ double2 make_double2(double, double );
 __HIP_DEVICE__ double3 make_double3(double, double, double );
 __HIP_DEVICE__ double4 make_double4(double, double, double, double );
 
+extern __HIP_DEVICE__ double  __longlong_as_double(long long int x);
+extern __HIP_DEVICE__ long long int __double_as_longlong(double x);
+
+
 /*
 ///---
 // Inline functions for creating vector types from basic types
diff --git a/src/device_util.cpp b/src/device_util.cpp
index 7b881751ae..7efb12d2d0 100644
--- a/src/device_util.cpp
+++ b/src/device_util.cpp
@@ -14,8 +14,7 @@ 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,
+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.
 */
@@ -2587,7 +2586,18 @@ __HIP_DEVICE__ double4 make_double4(double x, double y, double z, double w)
     return d4;
 }
 
-__device__ void  __threadfence_system(void){
+
+__HIP_DEVICE__ double  __longlong_as_double(long long int x)
+{
+  return static_cast<double>(x);
+}
+
+__HIP_DEVICE__ long long __double_as_longlong(double x)
+{
+  return static_cast<long long>(x);
+}
+
+__HIP_DEVICE__ void  __threadfence_system(void){
     // no-op
 }
 
@@ -3380,3 +3390,5 @@ __host__ double norm4d(double a, double b, double c, double d)
 {
     return std::sqrt(a*a + b*b + c*c + d*d);
 }
+
+
diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index c4b092aa5f..4d7b0eeb0d 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -25,6 +25,7 @@ THE SOFTWARE.
 
 #include <hc.hpp>
 #include <hsa/hsa.h>
+#include "hsa/hsa_ext_amd.h"
 #include "hip_util.h"
 
 
@@ -367,6 +368,26 @@ struct LockedBase {
     MUTEX_TYPE  _mutex;
 };
 
+/**
+ * HIP IPC Handle Size
+ */
+#define HIP_IPC_HANDLE_SIZE 64
+struct __HIP_DEVICE__ ihipIpcMemHandle
+{
+    volatile hsa_amd_ipc_memory_t handle; ///< ipc memory handle on ROCr
+    char reserved[HIP_IPC_HANDLE_SIZE];
+};
+
+
+class ihipModule_t{
+public:
+  hsa_executable_t executable;
+  hsa_code_object_t object;
+  std::string fileName;
+  void *ptr;
+  size_t size;
+};
+
 class ihipFunction_t{
 public:
     ihipFunction_t(const char *name) {
@@ -507,9 +528,9 @@ private:
 
     // The unsigned return is hipMemcpyKind
     unsigned resolveMemcpyDirection(bool srcInDeviceMem, bool dstInDeviceMem);
-    void resolveHcMemcpyDirection(unsigned hipMemKind, 
-                                  const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo, 
-                                  hc::hcCommandKind *hcCopyDir, 
+    void resolveHcMemcpyDirection(unsigned hipMemKind,
+                                  const hc::AmPointerInfo *dstPtrInfo, const hc::AmPointerInfo *srcPtrInfo,
+                                  hc::hcCommandKind *hcCopyDir,
                                   ihipCtx_t **copyDevice,
                                   bool *forceUnpinnedCopy);
 
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index b565c5a770..9a09c5ff70 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1037,15 +1037,24 @@ hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDevi
 
 
 //TODO: IPC implementaiton:
-hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
-    return hipSuccess;
-}
-hipError_t hipIpcCloseMemHandle(void *devPtr){
-    return hipSuccess;
-}
-hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
-    return hipSuccess;
-}
 hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
-    return hipSuccess;
+    // HIP_INIT_API ( devPtr, handle.handle , flags);
+    hipError_t hipStatus = hipSuccess;
+    return hipStatus;
 }
+
+hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
+    HIP_INIT_API ( handle, devPtr);
+    hipError_t hipStatus = hipSuccess;
+    return hipStatus;
+}
+
+hipError_t hipIpcCloseMemHandle(void *devPtr){
+    HIP_INIT_API ( devPtr );
+    hipError_t hipStatus = hipSuccess;
+    return hipStatus;
+}
+
+// hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
+//     return hipSuccess;
+// }

From 01f688587e845a4a7561e262755fcbf134237c71 Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Tue, 6 Dec 2016 14:21:03 -0600
Subject: [PATCH 20/28] change hipgetPCIID to take int as third parameter

Change-Id: I4429b36756a6d868a769abd783bf28a55147c0d0
---
 src/hip_device.cpp | 12 +-----------
 1 file changed, 1 insertion(+), 11 deletions(-)

diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 38021ad525..093277b941 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -337,16 +337,7 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return ihipLogStatus(e);
 }
 
-// hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
-// {
-//     HIP_INIT_API(pciBusId, len, device);
-//     hipError_t e = hipSuccess;
-//     int deviceId= device->_deviceId;
-//     e = ihipDeviceGetAttribute((int*)pciBusId, hipDeviceAttributePciBusId, deviceId);
-//     return ihipLogStatus(e);
-// }
-
-hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,int device)
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
 {
     HIP_INIT_API(pciBusId, len, device);
     hipError_t e = hipSuccess;
@@ -358,7 +349,6 @@ hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,int device)
     }
     return ihipLogStatus(e);
 }
-
 hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
 {
     HIP_INIT_API(bytes, device);

From 6fcfab25528d343ed74af35d5d9a0679853f1938 Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Tue, 6 Dec 2016 19:10:17 -0600
Subject: [PATCH 21/28] Fix issue of ihipModule_t double defined

Change-Id: I508677e595776fd573a2f224691116d01288dc78
---
 src/hip_hcc.h | 10 ----------
 1 file changed, 10 deletions(-)

diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index 4d7b0eeb0d..cd099f026c 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -378,16 +378,6 @@ struct __HIP_DEVICE__ ihipIpcMemHandle
     char reserved[HIP_IPC_HANDLE_SIZE];
 };
 
-
-class ihipModule_t{
-public:
-  hsa_executable_t executable;
-  hsa_code_object_t object;
-  std::string fileName;
-  void *ptr;
-  size_t size;
-};
-
 class ihipFunction_t{
 public:
     ihipFunction_t(const char *name) {

From 266b27ac832d6a44c5ef9a12b96c65d759b09d04 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Wed, 7 Dec 2016 12:12:40 +0530
Subject: [PATCH 22/28] hipDeviceGetPCIBusId int version changes for CUDA
 runtime API

Change-Id: I4d3b995f1d1ac83415ca84808a074e5c8cd72f3c
---
 include/hip/hcc_detail/hip_runtime_api.h  |  2 +-
 include/hip/nvcc_detail/hip_runtime_api.h |  7 ++++++-
 src/hip_device.cpp                        | 14 ++++++++++++++
 3 files changed, 21 insertions(+), 2 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 1af8108441..3ab43ea64d 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -1622,7 +1622,7 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device);
  *
  * @returns #hipSuccess, #hipErrorInavlidDevice
  */
-// hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
 
 
 /**
diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h
index 1a24615475..5f5931cf1b 100644
--- a/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/include/hip/nvcc_detail/hip_runtime_api.h
@@ -770,9 +770,14 @@ inline static hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return hipCUResultTohipError(cuDeviceGetName(name,len,device));
 }
 
+inline static hipError_t hipDeviceGetPCIBusId(char* pciBusId,int len,int device)
+{
+    return hipCUDAErrorTohipError(cudaDeviceGetPCIBusId(pciBusId,len,device));
+}
+
 inline static hipError_t hipDeviceGetPCIBusId(char* pciBusId,int len,hipDevice_t device)
 {
-    return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device));
+    return hipCUResultTohipError(cuDeviceGetPCIBusId(pciBusId,len,device));
 }
 
 inline static hipError_t hipDeviceGetByPCIBusId(int* device, const int *pciBusId)
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 093277b941..75266efd4a 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -337,6 +337,20 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return ihipLogStatus(e);
 }
 
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
+{
+    HIP_INIT_API(pciBusId, len, device);
+    hipError_t e = hipSuccess;
+    int deviceId= device->_deviceId;
+    int tempPciBusId = 0;
+    e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, deviceId);
+    if( e == hipSuccess) {
+        std::string tempPciStr = std::to_string(tempPciBusId);
+        memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );
+    }
+    return ihipLogStatus(e);
+}
+
 hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
 {
     HIP_INIT_API(pciBusId, len, device);

From a53d35fd6c0770f369c38cbaa42b64dedaf68b90 Mon Sep 17 00:00:00 2001
From: pensun <Peng.Sun@amd.com>
Date: Wed, 7 Dec 2016 15:31:23 -0600
Subject: [PATCH 23/28] HIP IPC implementation on ROCr IPC APIs

Change-Id: I1ca9d520f5d0b1b56694211471b81eb7c6c23d16
---
 include/hip/hcc_detail/hip_runtime_api.h | 31 +++++++-------
 include/hip/hip_runtime_api.h            |  1 +
 src/hip_hcc.h                            |  6 ++-
 src/hip_memory.cpp                       | 52 +++++++++++++++++++++---
 4 files changed, 67 insertions(+), 23 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 3ab43ea64d..87b4df3275 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -61,7 +61,7 @@ typedef struct ihipStream_t *hipStream_t;
 
 #define hipIpcMemLazyEnablePeerAccess 0
 
-typedef struct ihipIpcMemHandle *hipIpcMemHandle_t;
+typedef struct ihipIpcMemHandle_t *hipIpcMemHandle_t;
 
 //TODO: IPC event handle currently unsupported
 struct ihipIpcEventHandle_t;
@@ -1822,19 +1822,19 @@ hipError_t hipProfilerStop();
  * @brief Gets an interprocess memory handle for an existing device memory
  *          allocation
  *
- * Takes a pointer to the base of an existing device memory allocation created 
- * with hipMalloc and exports it for use in another process. This is a 
+ * Takes a pointer to the base of an existing device memory allocation created
+ * with hipMalloc and exports it for use in another process. This is a
  * lightweight operation and may be called multiple times on an allocation
- * without adverse effects. 
+ * without adverse effects.
  *
  * If a region of memory is freed with hipFree and a subsequent call
  * to hipMalloc returns memory with the same device address,
  * hipIpcGetMemHandle will return a unique handle for the
- * new memory. 
+ * new memory.
  *
  * @param handle - Pointer to user allocated hipIpcMemHandle to return
  *                    the handle in.
- * @param devPtr - Base pointer to previously allocated device memory 
+ * @param devPtr - Base pointer to previously allocated device memory
  *
  * @returns
  * hipSuccess,
@@ -1850,14 +1850,14 @@ extern __host__ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *d
  *          and returns a device pointer usable in the local process.
  *
  * Maps memory exported from another process with hipIpcGetMemHandle into
- * the current device address space. For contexts on different devices 
+ * the current device address space. For contexts on different devices
  * hipIpcOpenMemHandle can attempt to enable peer access between the
- * devices as if the user called hipDeviceEnablePeerAccess. This behavior is 
- * controlled by the hipIpcMemLazyEnablePeerAccess flag. 
+ * devices as if the user called hipDeviceEnablePeerAccess. This behavior is
+ * controlled by the hipIpcMemLazyEnablePeerAccess flag.
  * hipDeviceCanAccessPeer can determine if a mapping is possible.
  *
  * Contexts that may open hipIpcMemHandles are restricted in the following way.
- * hipIpcMemHandles from each device in a given process may only be opened 
+ * hipIpcMemHandles from each device in a given process may only be opened
  * by one context per device per other process.
  *
  * Memory returned from hipIpcOpenMemHandle must be freed with
@@ -1866,7 +1866,7 @@ extern __host__ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *d
  * Calling hipFree on an exported memory region before calling
  * hipIpcCloseMemHandle in the importing context will result in undefined
  * behavior.
- * 
+ *
  * @param devPtr - Returned device pointer
  * @param handle - hipIpcMemHandle to open
  * @param flags  - Flags for this operation. Must be specified as hipIpcMemLazyEnablePeerAccess
@@ -1877,15 +1877,16 @@ extern __host__ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *d
  * hipErrorInvalidResourceHandle,
  * hipErrorTooManyPeers
  *
- * @note No guarantees are made about the address returned in @p *devPtr.  
+ * @note No guarantees are made about the address returned in @p *devPtr.
  * In particular, multiple processes may not receive the same address for the same @p handle.
  *
  */
-extern __host__ hipError_t hipIpcOpenMemHandle(void **devPtr, hipIpcMemHandle_t handle, unsigned int flags);
+extern __host__ hipError_t hipIpcOpenMemHandle(void **devPtr,
+        hipIpcMemHandle_t handle, unsigned int flags);
 
 /**
  * @brief Close memory mapped with hipIpcOpenMemHandle
- * 
+ *
  * Unmaps memory returnd by hipIpcOpenMemHandle. The original allocation
  * in the exporting process as well as imported mappings in other processes
  * will be unaffected.
@@ -1894,7 +1895,7 @@ extern __host__ hipError_t hipIpcOpenMemHandle(void **devPtr, hipIpcMemHandle_t
  * last mapping using them.
  *
  * @param devPtr - Device pointer returned by hipIpcOpenMemHandle
- * 
+ *
  * @returns
  * hipSuccess,
  * hipErrorMapBufferObjectFailed,
diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h
index 5a8dd44e61..a45a1ee27e 100644
--- a/include/hip/hip_runtime_api.h
+++ b/include/hip/hip_runtime_api.h
@@ -210,6 +210,7 @@ typedef enum hipError_t {
     hipErrorRuntimeOther            = 1053,                   ///< HSA runtime call other than memory returned error.  Typically not seen in production systems.
     hipErrorHostMemoryAlreadyRegistered = 1061, ///< Produced when trying to lock a page-locked memory.
     hipErrorHostMemoryNotRegistered = 1062,   ///< Produced when trying to unlock a non-page-locked memory.
+    hipErrorMapBufferObjectFailed = 1071,   ///< Produced when the IPC memory attach failed from ROCr.
     hipErrorTbd                             ///< Marker that more error codes are needed.
 } hipError_t;
 
diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index cd099f026c..82290bc489 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -372,10 +372,12 @@ struct LockedBase {
  * HIP IPC Handle Size
  */
 #define HIP_IPC_HANDLE_SIZE 64
-struct __HIP_DEVICE__ ihipIpcMemHandle
+class ihipIpcMemHandle_t
 {
-    volatile hsa_amd_ipc_memory_t handle; ///< ipc memory handle on ROCr
+public:
+    hsa_amd_ipc_memory_t ipc_handle; ///< ipc memory handle on ROCr
     char reserved[HIP_IPC_HANDLE_SIZE];
+    size_t psize;
 };
 
 class ihipFunction_t{
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 9a09c5ff70..f2ab6d19a0 100644
--- a/src/hip_memory.cpp
+++ b/src/hip_memory.cpp
@@ -1031,27 +1031,67 @@ hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDevi
         *pbase = amPointerInfo._devicePointer;
         *psize = amPointerInfo._sizeBytes;
     }
-    hipStatus = hipErrorInvalidDevicePointer;
+    else
+        hipStatus = hipErrorInvalidDevicePointer;
     return hipStatus;
 }
 
 
 //TODO: IPC implementaiton:
-hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
-    // HIP_INIT_API ( devPtr, handle.handle , flags);
-    hipError_t hipStatus = hipSuccess;
-    return hipStatus;
-}
 
 hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
     HIP_INIT_API ( handle, devPtr);
     hipError_t hipStatus = hipSuccess;
+    // Get the size of allocated pointer
+    size_t psize;
+    hc::accelerator acc;
+    hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 );
+    am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , devPtr );
+    if (status == AM_SUCCESS) {
+        psize = (size_t)amPointerInfo._sizeBytes;
+    }
+    else
+        hipStatus = hipErrorInvalidResourceHandle;
+
+    // Save the size of the pointer to hipIpcMemHandle
+    (*handle)->psize = psize;
+
+    // Create HSA ipc memory
+    hsa_status_t hsa_status =
+        hsa_amd_ipc_memory_create(devPtr, psize, &(*handle)->ipc_handle);
+    if(hsa_status!= HSA_STATUS_SUCCESS)
+        hipStatus = hipErrorMemoryAllocation;
+
+    return hipStatus;
+}
+
+hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags){
+// HIP_INIT_API ( devPtr, handle.handle , flags);
+    hipError_t hipStatus = hipSuccess;
+
+    // Get the current device agent.
+    hc::accelerator acc;
+    hsa_agent_t *agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+    if(!agent)
+        return hipErrorInvalidResourceHandle;
+
+    //Attach ipc memory
+    hsa_status_t hsa_status =
+        hsa_amd_ipc_memory_attach(&handle->ipc_handle, handle->psize, 1, agent, devPtr);
+    if(hsa_status != HSA_STATUS_SUCCESS)
+        hipStatus = hipErrorMapBufferObjectFailed;
+
     return hipStatus;
 }
 
 hipError_t hipIpcCloseMemHandle(void *devPtr){
     HIP_INIT_API ( devPtr );
     hipError_t hipStatus = hipSuccess;
+
+    hsa_status_t hsa_status =
+        hsa_amd_ipc_memory_detach(devPtr);
+    if(hsa_status != HSA_STATUS_SUCCESS)
+        return hipErrorInvalidResourceHandle;
     return hipStatus;
 }
 

From 6e6b5180983999bd0625eff749c7e231910ffb8e Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 8 Dec 2016 12:50:25 +0530
Subject: [PATCH 24/28] hcc_detail/hip_runtime_api.h: Fix IPC API signature

Change-Id: I0be0f09c62f231620341141bd66183c3338be56a
---
 include/hip/hcc_detail/hip_runtime_api.h | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 87b4df3275..9745ca43d0 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -1843,7 +1843,7 @@ hipError_t hipProfilerStop();
  * hipErrorMapBufferObjectFailed,
  *
  */
-extern __host__ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *devPtr);
+hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *devPtr);
 
 /**
  * @brief Opens an interprocess memory handle exported from another process
@@ -1881,7 +1881,7 @@ extern __host__ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t *handle, void *d
  * In particular, multiple processes may not receive the same address for the same @p handle.
  *
  */
-extern __host__ hipError_t hipIpcOpenMemHandle(void **devPtr,
+hipError_t hipIpcOpenMemHandle(void **devPtr,
         hipIpcMemHandle_t handle, unsigned int flags);
 
 /**
@@ -1902,7 +1902,7 @@ extern __host__ hipError_t hipIpcOpenMemHandle(void **devPtr,
  * hipErrorInvalidResourceHandle,
  *
  */
-extern __host__ hipError_t hipIpcCloseMemHandle(void *devPtr);
+hipError_t hipIpcCloseMemHandle(void *devPtr);
 
 
 // hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);

From d35c8128a85c3fe259b9ae1e2795b3fec9fe9748 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Thu, 8 Dec 2016 14:35:58 +0530
Subject: [PATCH 25/28] Fixed build error due to GetPCIBusId overloaded
 function

Change-Id: I626446f2c72c8143f08c95367bc1c528abeaf69d
---
 include/hip/hcc_detail/hip_runtime_api.h | 22 +++++++++++-----------
 src/hip_device.cpp                       |  2 ++
 2 files changed, 13 insertions(+), 11 deletions(-)

diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index 9745ca43d0..52e59ed17a 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -1614,17 +1614,6 @@ hipError_t hipDeviceComputeCapability(int *major,int *minor,hipDevice_t device);
  */
 hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device);
 
-/**
- * @brief Returns a PCI Bus Id string for the device.
- * @param [out] pciBusId
- * @param [in] len
- * @param [hipDevice_t] device
- *
- * @returns #hipSuccess, #hipErrorInavlidDevice
- */
-hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
-
-
 /**
  * @brief Returns a PCI Bus Id string for the device, overloaded to take int device ID.
  * @param [out] pciBusId
@@ -1915,6 +1904,17 @@ hipError_t hipIpcCloseMemHandle(void *devPtr);
 } /* extern "c" */
 #endif
 
+#ifdef __cplusplus
+/**
+ * @brief Returns a PCI Bus Id string for the device.
+ * @param [out] pciBusId
+ * @param [in] len
+ * @param [hipDevice_t] device
+ *
+ * @returns #hipSuccess, #hipErrorInavlidDevice
+ */
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device);
+#endif
 
 /**
  *-------------------------------------------------------------------------------------------------
diff --git a/src/hip_device.cpp b/src/hip_device.cpp
index 75266efd4a..bf45125b60 100644
--- a/src/hip_device.cpp
+++ b/src/hip_device.cpp
@@ -337,6 +337,7 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return ihipLogStatus(e);
 }
 
+#ifdef __cplusplus
 hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
 {
     HIP_INIT_API(pciBusId, len, device);
@@ -350,6 +351,7 @@ hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device)
     }
     return ihipLogStatus(e);
 }
+#endif
 
 hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
 {

From 23bbe6e4670b452b04bef9d3f5aa34eb59f700d7 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Thu, 8 Dec 2016 20:28:43 +0300
Subject: [PATCH 26/28] Fix the limitation of supported input files.

Actually .cu and .cuda was supported.

+ All the file names are allowed, including file names without extension.

[IMPORTANT]
To hipify CUDA input file, which name is not *.cu, please add option "-x cuda" after tool's options ending marker "--", for instance:

./hipify-clang NeuralNet -- -x cuda

This option will go to clang itself, not the hipify tool.
---
 hipify-clang/src/Cuda2Hip.cpp | 15 +++++++--------
 1 file changed, 7 insertions(+), 8 deletions(-)

diff --git a/hipify-clang/src/Cuda2Hip.cpp b/hipify-clang/src/Cuda2Hip.cpp
index 87a69e8cb9..2c934aebbc 100644
--- a/hipify-clang/src/Cuda2Hip.cpp
+++ b/hipify-clang/src/Cuda2Hip.cpp
@@ -2110,19 +2110,18 @@ int main(int argc, const char **argv) {
   if (dst.empty()) {
     dst = fileSources[0];
     if (!Inplace) {
-      size_t pos = dst.rfind(".cu");
-      if (pos != std::string::npos) {
-        dst = dst.substr(0, pos) + ".hip.cu";
+      size_t pos = dst.rfind(".");
+      if (pos != std::string::npos && pos+1 < dst.size()) {
+        dst = dst.substr(0, pos) + ".hip." + dst.substr(pos+1, dst.size()-pos-1);
       } else {
-        llvm::errs() << "Input .cu file was not specified.\n";
-        return 1;
+        dst += ".hip.cu";
       }
     }
   } else {
     if (Inplace) {
-      llvm::errs() << "Conflict: both -o and -inplace options are specified.";
+      llvm::errs() << "Conflict: both -o and -inplace options are specified.\n";
     }
-    dst += ".cu";
+    dst += ".hip";
   }
   // copy source file since tooling makes changes "inplace"
   std::ifstream source(fileSources[0], std::ios::binary);
@@ -2172,7 +2171,7 @@ int main(int argc, const char **argv) {
   Result = Rewrite.overwriteChangedFiles();
 
   if (!Inplace) {
-    size_t pos = dst.rfind(".cu");
+    size_t pos = dst.rfind(".");
     if (pos != std::string::npos) {
       rename(dst.c_str(), dst.substr(0, pos).c_str());
     }

From 2374153c10dc14ae0eff0e97ca7e560686401150 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Thu, 8 Dec 2016 22:45:10 +0300
Subject: [PATCH 27/28] [HIPIFY] -no-output support.

Actually output file is created by clang itself, but isn't updated and is deleted after processing.
In cooperation with -print-stat -no-output (or single -n) is used for examine the source CUDA code.

Conflicting options:
-inplce
-o
---
 hipify-clang/src/Cuda2Hip.cpp | 22 ++++++++++++++++++----
 1 file changed, 18 insertions(+), 4 deletions(-)

diff --git a/hipify-clang/src/Cuda2Hip.cpp b/hipify-clang/src/Cuda2Hip.cpp
index 2c934aebbc..a124fb7c64 100644
--- a/hipify-clang/src/Cuda2Hip.cpp
+++ b/hipify-clang/src/Cuda2Hip.cpp
@@ -2107,6 +2107,16 @@ int main(int argc, const char **argv) {
   if (N) {
     NoOutput = PrintStats = true;
   }
+  if (NoOutput) {
+    if (Inplace) {
+      llvm::errs() << "Conflict: both -no-output and -inplace options are specified.\n";
+      return 1;
+    }
+    if (!dst.empty()) {
+      llvm::errs() << "Conflict: both -no-output and -o options are specified.\n";
+      return 1;
+    }
+  }
   if (dst.empty()) {
     dst = fileSources[0];
     if (!Inplace) {
@@ -2120,6 +2130,7 @@ int main(int argc, const char **argv) {
   } else {
     if (Inplace) {
       llvm::errs() << "Conflict: both -o and -inplace options are specified.\n";
+      return 1;
     }
     dst += ".hip";
   }
@@ -2167,15 +2178,18 @@ int main(int argc, const char **argv) {
   if (!Tool.applyAllReplacements(Rewrite)) {
     DEBUG(dbgs() << "Skipped some replacements.\n");
   }
-
-  Result = Rewrite.overwriteChangedFiles();
-
-  if (!Inplace) {
+  if (!NoOutput) {
+    Result = Rewrite.overwriteChangedFiles();
+  }
+  if (!Inplace && !NoOutput) {
     size_t pos = dst.rfind(".");
     if (pos != std::string::npos) {
       rename(dst.c_str(), dst.substr(0, pos).c_str());
     }
   }
+  if (NoOutput) {
+    remove(dst.c_str());
+  }
   if (PrintStats) {
     printStats(fileSources[0], PPCallbacks, Callback);
   }

From 428a1bc79f592c71542a32c07edea8620fa50226 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Thu, 8 Dec 2016 23:14:19 +0300
Subject: [PATCH 28/28] [HIPIFY] -no-backup option is added.

Is useful for release hipifying in place.
---
 hipify-clang/src/Cuda2Hip.cpp | 19 +++++++++++++------
 1 file changed, 13 insertions(+), 6 deletions(-)

diff --git a/hipify-clang/src/Cuda2Hip.cpp b/hipify-clang/src/Cuda2Hip.cpp
index a124fb7c64..c66a2b1b87 100644
--- a/hipify-clang/src/Cuda2Hip.cpp
+++ b/hipify-clang/src/Cuda2Hip.cpp
@@ -2004,6 +2004,11 @@ static cl::opt<bool> Inplace("inplace",
        cl::value_desc("inplace"),
        cl::cat(ToolTemplateCategory));
 
+static cl::opt<bool> NoBackup("no-backup",
+       cl::desc("Don't create a backup file for the hipified source"),
+       cl::value_desc("no-backup"),
+       cl::cat(ToolTemplateCategory));
+
 static cl::opt<bool> NoOutput("no-output",
        cl::desc("Don't write any translated output to stdout"),
        cl::value_desc("no-output"),
@@ -2134,12 +2139,14 @@ int main(int argc, const char **argv) {
     }
     dst += ".hip";
   }
-  // copy source file since tooling makes changes "inplace"
-  std::ifstream source(fileSources[0], std::ios::binary);
-  std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
-  dest << source.rdbuf();
-  source.close();
-  dest.close();
+  // backup source file since tooling may change "inplace"
+  if (!NoBackup || !Inplace) {
+    std::ifstream source(fileSources[0], std::ios::binary);
+    std::ofstream dest(Inplace ? dst + ".prehip" : dst, std::ios::binary);
+    dest << source.rdbuf();
+    source.close();
+    dest.close();
+  }
 
   RefactoringTool Tool(OptionsParser.getCompilations(), dst);
   ast_matchers::MatchFinder Finder;