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 09c4d813d0..7d2675d0ba 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} ";
         }
     }
 
@@ -192,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  = ();
@@ -260,6 +262,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') {
@@ -273,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 -
@@ -322,16 +334,12 @@ 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) {
-    $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/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/hipify-clang/src/Cuda2Hip.cpp b/hipify-clang/src/Cuda2Hip.cpp
index 87a69e8cb9..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"),
@@ -2107,29 +2112,41 @@ 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) {
-      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";
+      return 1;
     }
-    dst += ".cu";
+    dst += ".hip";
+  }
+  // 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();
   }
-  // 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();
 
   RefactoringTool Tool(OptionsParser.getCompilations(), dst);
   ast_matchers::MatchFinder Finder;
@@ -2168,15 +2185,18 @@ int main(int argc, const char **argv) {
   if (!Tool.applyAllReplacements(Rewrite)) {
     DEBUG(dbgs() << "Skipped some replacements.\n");
   }
-
-  Result = Rewrite.overwriteChangedFiles();
-
-  if (!Inplace) {
-    size_t pos = dst.rfind(".cu");
+  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);
   }
diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index b1877ed0b3..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);
@@ -828,7 +829,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 +839,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/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h
index ee703c4eec..52e59ed17a 100644
--- a/include/hip/hcc_detail/hip_runtime_api.h
+++ b/include/hip/hcc_detail/hip_runtime_api.h
@@ -57,6 +57,19 @@ typedef struct ihipDevice_t *hipDevice_t;
 
 typedef struct ihipStream_t *hipStream_t;
 
+//TODO: IPC implementation
+
+#define hipIpcMemLazyEnablePeerAccess 0
+
+typedef struct ihipIpcMemHandle_t *hipIpcMemHandle_t;
+
+//TODO: IPC event handle currently unsupported
+struct ihipIpcEventHandle_t;
+typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;
+
+
+//END TODO
+
 typedef struct ihipModule_t *hipModule_t;
 
 typedef struct ihipFunction_t *hipFunction_t;
@@ -1281,6 +1294,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
@@ -1590,14 +1615,15 @@ 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.
+ * @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 (int *pciBusId,int len,hipDevice_t device);
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,int device);
+
 
 /**
  * @brief Returns a handle to a compute device.
@@ -1779,13 +1805,116 @@ hipError_t hipProfilerStop();
  * @}
  */
 
+//TODO: implement IPC apis
 
+/**
+ * @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,
+ *
+ */
+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.
+ *
+ */
+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,
+ *
+ */
+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
 } /* 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/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/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/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h
index 625448094b..5f5931cf1b 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));
@@ -765,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(int *pciBusId,int len,hipDevice_t device)
+inline static hipError_t hipDeviceGetPCIBusId(char* pciBusId,int len,int device)
 {
-    return hipCUResultTohipError(cuDeviceGetPCIBusId((char*)pciBusId,len,device));
+    return hipCUDAErrorTohipError(cudaDeviceGetPCIBusId(pciBusId,len,device));
+}
+
+inline static hipError_t hipDeviceGetPCIBusId(char* pciBusId,int len,hipDevice_t device)
+{
+    return hipCUResultTohipError(cuDeviceGetPCIBusId(pciBusId,len,device));
 }
 
 inline static hipError_t hipDeviceGetByPCIBusId(int* device, const int *pciBusId)
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)
 
diff --git a/src/device_util.cpp b/src/device_util.cpp
index 203e7a7826..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.
 */
@@ -23,6 +22,7 @@ THE SOFTWARE.
 #include <hc.hpp>
 #include <grid_launch.h>
 #include <hc_math.hpp>
+#include "device_util.h"
 
 #include "hip/hip_runtime.h"
 
@@ -33,17 +33,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];
 __attribute__((address_space(1))) uint32_t gpuFlags[NUM_PAGES];
 
@@ -2597,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
 }
 
@@ -3390,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/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 4f5729791e..bf45125b60 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);
@@ -306,7 +305,7 @@ hipError_t hipSetDeviceFlags( unsigned int flags)
            }
        }
 
-       unsigned supportedFlags = hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax; 
+       unsigned supportedFlags = hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax;
 
        if (flags & (~supportedFlags)) {
           e = hipErrorInvalidValue;
@@ -338,15 +337,34 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device)
     return ihipLogStatus(e);
 }
 
-hipError_t hipDeviceGetPCIBusId (int *pciBusId,int len,hipDevice_t device)
+#ifdef __cplusplus
+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);
+    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);
 }
+#endif
 
+hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device)
+{
+    HIP_INIT_API(pciBusId, len, device);
+    hipError_t e = hipSuccess;
+    int tempPciBusId = 0;
+    e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, device);
+    if( e == hipSuccess) {
+        std::string tempPciStr = std::to_string(tempPciBusId);
+        memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() );
+    }
+    return ihipLogStatus(e);
+}
 hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device)
 {
     HIP_INIT_API(bytes, device);
diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp
index 36790271de..84794b02fb 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;
@@ -78,15 +81,17 @@ 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;
 
+int HIP_COHERENT_HOST_ALLOC = 0;
 
 
 
 
 
-#define HIP_USE_PRODUCT_NAME 0
+
+#define HIP_USE_PRODUCT_NAME 1
 //#define DISABLE_COPY_EXT 1
 
 
@@ -105,7 +110,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 +138,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 +216,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 +287,7 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
         } else if (HIP_WAIT_MODE == 2) {
             waitMode = hc::hcWaitModeActive;
         }
-            
+
         crit->_av.wait(waitMode);
     }
 
@@ -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.
@@ -385,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);
@@ -410,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) |
@@ -455,7 +465,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 +483,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 +823,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 +1191,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?
             }
@@ -1243,7 +1253,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,20 +1270,26 @@ 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.");
     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"); 
+
+    // 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 +1422,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 +1540,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();
 }
 
@@ -1716,10 +1741,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 +1774,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 +1814,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 +1871,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);
@@ -1883,8 +1908,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..82290bc489 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"
 
 
@@ -32,7 +33,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
 //
 
 
@@ -45,6 +46,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 +58,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 +160,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.
@@ -369,17 +368,18 @@ struct LockedBase {
     MUTEX_TYPE  _mutex;
 };
 
-
-class ihipModule_t{
+/**
+ * HIP IPC Handle Size
+ */
+#define HIP_IPC_HANDLE_SIZE 64
+class ihipIpcMemHandle_t
+{
 public:
-  hsa_executable_t executable;
-  hsa_code_object_t object;
-  std::string fileName;
-  void *ptr;
-  size_t size;
+    hsa_amd_ipc_memory_t ipc_handle; ///< ipc memory handle on ROCr
+    char reserved[HIP_IPC_HANDLE_SIZE];
+    size_t psize;
 };
 
-
 class ihipFunction_t{
 public:
     ihipFunction_t(const char *name) {
@@ -401,6 +401,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>
@@ -455,7 +478,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);
@@ -497,9 +520,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_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 }
diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp
index 089d1a626e..f2ab6d19a0 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>
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
@@ -156,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)
 {
@@ -192,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)){
@@ -852,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;
@@ -905,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) {
@@ -1037,3 +1020,81 @@ 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;
+    }
+    else
+        hipStatus = hipErrorInvalidDevicePointer;
+    return hipStatus;
+}
+
+
+//TODO: IPC implementaiton:
+
+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;
+}
+
+// hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle){
+//     return hipSuccess;
+// }
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index f7ac35c77b..badeac2dcd 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;
@@ -164,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)
@@ -172,6 +176,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 +202,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 +278,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 +316,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 +331,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,10 +365,26 @@ 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
 
 
-        ihipPostLaunchKernel(hStream, lp);
+        ihipPostLaunchKernel(f->_kernelName, hStream, lp);
 
     }
 
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/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;
+}
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
  */