diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 71901da4a4..a26e848985 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -128,8 +128,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
     set(CMAKE_C_COMPILER   "${HCC_HOME}/bin/hcc")
 
     # Set HIP_HCC so we know this is HIP compile, some files are shared with HCC (staging_buffer).
-    set(CMAKE_CXX_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_CXX_FLAGS}")
-    set(CMAKE_C_FLAGS   " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_C_FLAGS}")
+    set(CMAKE_CXX_FLAGS " -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_CXX_FLAGS}")
+    set(CMAKE_C_FLAGS   " -fPIC -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC ${CMAKE_C_FLAGS}")
 
     set(SOURCE_FILES src/device_util.cpp
                      src/hip_hcc.cpp
@@ -149,6 +149,10 @@ if(HIP_PLATFORM STREQUAL "hcc")
         #add_library(hip_hcc STATIC ${SOURCE_FILES})
         add_library(hip_hcc OBJECT ${SOURCE_FILES})
     endif()
+
+    # Generate .hipconfig
+    string(TIMESTAMP _timestamp)
+    file(WRITE "${PROJECT_BINARY_DIR}/.hipconfig" "# Auto-generated by cmake on ${_timestamp} local time\nHCC_HOME=${HCC_HOME}\nHCC_VERSION=${HCC_VERSION}\n")
 endif()
 
 # Build doxygen documentation
@@ -166,6 +170,9 @@ if(HIP_PLATFORM STREQUAL "hcc")
         #install(TARGETS hip_hcc DESTINATION lib)
         install(DIRECTORY ${PROJECT_BINARY_DIR}/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
     endif()
+
+    # Install .hipconfig
+    install(FILES ${PROJECT_BINARY_DIR}/.hipconfig DESTINATION lib)
 endif()
 
 # Install src, bin, include if necessary
diff --git a/projects/hip/INSTALL.md b/projects/hip/INSTALL.md
index 0b05a11f6b..4139cb2010 100644
--- a/projects/hip/INSTALL.md
+++ b/projects/hip/INSTALL.md
@@ -1,18 +1,17 @@
-<!-- START doctoc generated TOC please keep comment here to allow auto update -->
-<!-- DON'T EDIT THIS SECTION, INSTEAD RE-RUN doctoc TO UPDATE -->
-**Installation** 
+## Table of Contents
 
-- [Installing pre-built packages:](#installing-pre-built-packages)
-  - [Prerequisites](#prerequisites)
-  - [AMD (hcc)](#amd-hcc)
-  - [NVIDIA (nvcc)](#nvidia-nvcc)
-  - [Verify your installation](#verify-your-installation)
+<!-- toc -->
+
+- [Installing pre-built packages](#installing-pre-built-packages)
+  * [Prerequisites](#prerequisites)
+  * [AMD-hcc](#amd-hcc)
+  * [NVIDIA-nvcc](#nvidia-nvcc)
+  * [Verify your installation](#verify-your-installation)
 - [Building HIP from source](#building-hip-from-source)
-  - [HCC Options](#hcc-options)
-    - [Using HIP with the AMD Native-GCN compiler.](#using-hip-with-the-amd-native-gcn-compiler)
-    - [Compiling CodeXL markers for HIP Functions](#compiling-codexl-markers-for-hip-functions)
+  * [HCC Options](#hcc-options)
+    + [Using HIP with the AMD Native-GCN compiler.](#using-hip-with-the-amd-native-gcn-compiler)
 
-<!-- END doctoc generated TOC please keep comment here to allow auto update -->
+<!-- tocstop -->
 
 # Installing pre-built packages
 
@@ -62,7 +61,7 @@ HIP source code is available and the project can be built from source on the HCC
 2. Download HIP source code (from the [GitHub repot](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP).)
 3. Build and install HIP (This is the simple version assuming default paths ; see below for additional options.)
 ```
-cd HIP-privatestaging
+cd HIP
 mkdir build
 cd build
 cmake .. 
@@ -78,7 +77,7 @@ make install
 Here's a richer command-line that overrides the default paths:
 
 ```shell
-cd HIP-privatestaging
+cd HIP
 mkdir build
 cd build
 cmake -DHSA_PATH=/path/to/hsa -DHCC_HOME=/path/to/hcc -DCMAKE_INSTALL_PREFIX=/where/to/install/hip -DCMAKE_BUILD_TYPE=Release ..
@@ -105,42 +104,3 @@ Alternatively, this sections describes how to build it from source:
 export HCC_HOME=/path/to/native/hcc
 ```
 
-
-### Compiling CodeXL markers for HIP Functions
-HIP can generate markers at function begin/end which are displayed on the CodeXL timeline view.  To do this, you need to install CodeXL, tell HIP
-where the CodeXL install directory lives, and enable HIP to generate the markers:
-
-1. Install CodeXL
-See [CodeXL Download](http://developer.amd.com/tools-and-sdks/opencl-zone/codexl/?webSyncID=9d9c2cb9-3d73-5e65-268a-c7b06428e5e0&sessionGUID=29beacd0-d654-ddc6-a3e2-b9e6c0b0cc77) for the installation file.
-Also this [blog](http://gpuopen.com/getting-up-to-speed-with-the-codexl-gpu-profiler-and-radeon-open-compute/) provides more information and tips for using CodeXL.  In addition to installing the CodeXL profiling 
-and visualization tools, CodeXL also comes with an SDK that allow applications to add markers to the timeline viewer.  We'll be linking HIP against this library.
-
-2. Set CODEXL_PATH
-```shell
-# set to your code-xl installation location:
-export CODEXL_PATH=/opt/AMD/CodeXL
-```
-
-3. Enable in source code.
-In src/hip_hcc.cpp, enable the define 
-```c
-#define COMPILE_TRACE_MARKER 1
-```
-
-
-Then recompile the target application, run with profiler enabled to generate ATP file or trace log.
-```shell
-# Use profiler to generate timeline view:
-$CODEXL_PATH/CodeXLGpuProfiler -A  -o  ./myHipApp  
-...
-Session output path: /home/me/HIP-privatestaging/tests/b1/mytrace.atp
-```
-
-You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable.  This can be useful for tracing application flow.  Also can be combined with the more detailed debug information provided
-by the HIP_DB switch.  For example:
-```shell
-# Trace to stderr showing begin/end of each function (with arguments) + intermediate debug trace during the execution of each function.
-HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp  
-```
-
-Note this trace mode uses colors.  "less -r" can handle raw control characters and will display the debug output in proper colors.
diff --git a/projects/hip/README.md b/projects/hip/README.md
index 4b69485ba3..f9f653f2ba 100644
--- a/projects/hip/README.md
+++ b/projects/hip/README.md
@@ -21,10 +21,69 @@ New projects can be developed directly in the portable HIP C++ language and can
 - [clang-hipify](clang-hipify/README.md)
 - [Developer/CONTRIBUTING Info](CONTRIBUTING.md)
 - [Release Notes](RELEASE.md)
+
 ## How do I get set up?
 
 See the [Installation](INSTALL.md) notes.
 
+## Simple Example
+The HIP API includes functions such as hipMalloc, hipMemcpy, and hipFree.  
+Programmers familiar with CUDA will also be able to quickly learn and start coding with the HIP API.
+Compute kernels are launched with the "hipLaunchKernel" macro call.    Here is simple example showing a
+snippet of HIP API code:
+
+```cpp
+hipMalloc(&A_d, Nbytes));
+hipMalloc(&C_d, Nbytes));
+
+hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice);
+
+const unsigned blocks = 512;
+const unsigned threadsPerBlock = 256;
+hipLaunchKernel(vector_square,   /* compute kernel*/
+                dim3(blocks), dim3(threadsPerBlock), 0/*dynamic shared*/, 0/*stream*/,     /* launch config*/
+                C_d, A_d, N);  /* arguments to the compute kernel */
+
+hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost); 
+```
+
+
+The HIP kernel language defines builtins for determining grid and block coordinates, math functions, short vectors, 
+atomics, and timer functions. It also specifies additional defines and keywords for function types, address spaces, and 
+optimization controls.  (See the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) for a full description).
+Here's an example of defining a simple 'vector_square' kernel.  
+
+
+
+```cpp
+template <typename T>
+__global__ void
+vector_square(T *C_d, const T *A_d, size_t N)
+{
+    size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
+    size_t stride = hipBlockDim_x * hipGridDim_x ;
+
+    for (size_t i=offset; i<N; i+=stride) {
+        C_d[i] = A_d[i] * A_d[i];
+    }
+}
+```
+
+The HIP Runtime API code and compute kernel definition can exist in the same source file - HIP takes care of generating host and device code appropriately.
+
+## HIP Portability and Compiler Technology
+HIP C++ code can be compiled with either :
+- On the Nvidia CUDA platform, HIP provides header file which translate from the HIP runtime APIs to CUDA runtime APIs.  The header file contains mostly inlined 
+  functions and thus has very low overhead - developers coding in HIP should expect the same perforamnce as coding in native CUDA.  The code is then 
+  compiled with nvcc, the standard C++ compiler provided with the CUDA SDK.  Developers can use any tools supported by the CUDA SDK including the CUDA
+  profiler and debugger.
+- On the AMD ROCm platform, HIP provides a header and runtime library built on top of hcc compiler.  The HIP runtime implements HIP streams, events, and memory APIs, 
+  and is a object library that is linked with the application.  The source code for all headers and the library implementation is available on GitHub.  
+  HIP developers on ROCm can use AMD's CodeXL for debugging and profiling.
+
+Thus HIP source code can be compiled to run on either platform.  Platform-specific features can be isolated to a specific platform using conditional compilation.  Thus HIP
+provides source portability to either platform.   HIP provides the _hipcc_ compiler driver which will call the appropriate toolchain depending on the desired platform.
+
 
 ## Examples and Getting Started:
 
diff --git a/projects/hip/RELEASE.md b/projects/hip/RELEASE.md
index f60dafd4b5..da689f42ea 100644
--- a/projects/hip/RELEASE.md
+++ b/projects/hip/RELEASE.md
@@ -2,9 +2,26 @@
 
 We have attempted to document known bugs and limitations - in particular the [HIP Kernel Language](docs/markdown/hip_kernel_language.md) document uses the phrase "Under Development", and the [HIP Runtime API bug list](http://gpuopen-professionalcompute-tools.github.io/HIP/bug.html) lists known bugs. 
 
-===================================================================================================
 Upcoming:
-- hipLaunchKernel supports one-dimensional grid and/or block dims, without explicit cast to dim3 type.
+- Stability: Enforce perioidic host synchronization to reclaim resources if the application has launched a large
+  number of commands (>1K) without synchronizing.  
+- Register keyword now silently ignored on HCC (previously would emit warning).
+- Doc updates: Add some more frequently asked questions to FAQ, fix TOC in some files, review.
+
+===================================================================================================
+
+## Revision History:
+
+===================================================================================================
+Release:0.92.00
+- hipLaunchKernel supports one-dimensional grid and/or block dims, without explicit cast to dim3 type (actually in 0.90.00)
+- fp16 software support
+- Support for Hawaii dGPUs using environment variable ROCM_TARGET=hawaii
+- Support hipArray
+- Improved profiler support
+- Documentation updates
+- Improvements to clang-hipify
+
 
 ===================================================================================================
 Release:0.90.00
@@ -16,9 +33,6 @@ Date: 2016.06.29
 - Improve error code reporting on nvcc.
 - Add hipPeekAtError for nvcc.
 
-===================================================================================================
-
-## Revision History:
 
 ===================================================================================================
 Release:0.86.00
diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc
index ce888fea8f..1141e1c08d 100755
--- a/projects/hip/bin/hipcc
+++ b/projects/hip/bin/hipcc
@@ -31,6 +31,25 @@ $verbose = 0 unless defined $verbose;
 $HIP_PATH=$ENV{'HIP_PATH'};
 $HIP_PATH=dirname (dirname $0) unless defined $HIP_PATH;        # use parent directory of hipcc
 
+#---
+# Read .hipconfig
+my %hipConfig = ();
+$hipConfig{'VALID'}=0;
+if (open (CONFIG, "$HIP_PATH/lib/.hipconfig")) {
+    while (<CONFIG>) {
+        my $config_line=$_;
+        chop ($config_line);
+        $config_line =~ s/^\s*//;
+        $config_line =~ s/\s*$//;
+        if (($config_line !~ /^#/) && ($config_line ne "")) {
+            my ($name, $value) = split (/=/, $config_line);
+            $hipConfig{$name} = $value;
+            $hipConfig{'VALID'}=1;
+        }
+    }
+    close(CONFIG);
+}
+
 #---
 #HIP_PLATFORM controls whether to use NVCC or HCC for compilation:
 $HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform`;
@@ -52,14 +71,13 @@ if ($HIP_PLATFORM eq "hcc") {
 
     $HCC_HOME=$ENV{'HCC_HOME'};
     $HCC_HOME="/opt/rocm/hcc" unless defined $HCC_HOME;
+    $HCC_VERSION=`${HCC_HOME}/bin/hcc --version | cut -d" " -f9 | tr -d "\n"`;
 
     $ROCM_PATH=$ENV{'ROCM_PATH'};
     $ROCM_PATH="/opt/rocm" unless defined $ROCM_PATH;
 
     $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'};
-    if ($HIP_ATP_MARKER) {
-        $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger";
-    }
+    $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger";
 
     $ROCM_TARGET=$ENV{'ROCM_TARGET'};
     $ROCM_TARGET="fiji" unless defined $ROCM_TARGET;
@@ -73,6 +91,7 @@ if ($HIP_PLATFORM eq "hcc") {
 
     $HIPCXXFLAGS .= " -I$HIP_PATH/include/hcc_detail/cuda";
     $HIPCXXFLAGS .= " -I$HSA_PATH/include";
+    $HIPCXXFLAGS .= " -Wno-deprecated-register";
     $HIPLDFLAGS =  "-hc -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive";
 
     # Suppress linker warnings in case HCC distribution contains OpenCL/SPIR symbols
@@ -103,11 +122,11 @@ if ($HIP_PLATFORM eq "hcc") {
         if (-e $marker_inc_path) {
             $HIPCXXFLAGS .= " -I$marker_inc_path";
         }
+    }
 
-        $marker_lib_path = "$marker_path/bin/x86_64";
-        if (-e $marker_lib_path) {
-            $HIPLDFLAGS .= " -L$marker_lib_path -lCXLActivityLogger -Wl,--rpath=$marker_lib_path";
-        }
+    $marker_lib_path = "$marker_path/bin/x86_64";
+    if (-e $marker_lib_path) {
+        $HIPLDFLAGS .= " -L$marker_lib_path -lCXLActivityLogger -Wl,--rpath=$marker_lib_path";
     }
 
     # Add C++ libs for GCC.
@@ -268,5 +287,8 @@ if ($printHipVersion) {
     print $HIP_VERSION, "\n";
 }
 if ($runCmd) {
+    if ($hipConfig{'VALID'} and $HIP_PLATFORM eq "hcc" and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) {
+        print ("HIP was built using $hipConfig{'HCC_VERSION'}, but you are using $HCC_VERSION. Please rebuild HIP.\n") && die ();
+    }
     system ("$CMD") and die ();
 }
diff --git a/projects/hip/clang-hipify/README.md b/projects/hip/clang-hipify/README.md
index 6ea9e4a7a7..f95eed1fc8 100644
--- a/projects/hip/clang-hipify/README.md
+++ b/projects/hip/clang-hipify/README.md
@@ -1,3 +1,14 @@
+## Table of Contents
+
+<!-- toc -->
+
+- [Using hipify-clang](#using-hipify-clang)
+  * [Build and install](#build-and-install)
+  * [Running and using hipify-clang](#running-and-using-hipify-clang)
+    + [Disclaimer](#disclaimer)
+
+<!-- tocstop -->
+
 ## Using hipify-clang
 
 `hipify-clang` is a clang-based tool which can automate the translation of CUDA source code into portable HIP C++.
@@ -43,4 +54,4 @@ The information contained herein is for informational purposes only, and is subj
 
 AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies.
 
-Copyright (c) 2014-2016 Advanced Micro Devices, Inc. All rights reserved.
\ No newline at end of file
+Copyright (c) 2014-2016 Advanced Micro Devices, Inc. All rights reserved.
diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
index fc4f7c3fe5..189f172c95 100644
--- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
+++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md
@@ -25,7 +25,7 @@
 | `cudaIpcOpenEventHandle`                                  |                               | Opens an interprocess event handle for use in the current process.                                                             |
 | `cudaIpcOpenMemHandle`                                    |                               | Opens an interprocess memory handle exported from another process and returns a device pointer usable in the local process.    |
 | `cudaSetDevice`                                           | `hipSetDevice`                | Set device to be used for GPU executions.                                                                                      |
-| `cudaSetDeviceFlags`                                      |                               | Sets flags to be used for device executions.                                                                                   |
+| `cudaSetDeviceFlags`                                      | `hipSetDeviceFlags`           | Sets flags to be used for device executions.                                                                                   |
 | `cudaSetValidDevices`                                     |                               | Set a list of devices that can be used for CUDA.                                                                               |`
 
 **2. Error Handling**
diff --git a/projects/hip/docs/markdown/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md
index e71dcdfb36..0f2a43fe26 100644
--- a/projects/hip/docs/markdown/hip_faq.md
+++ b/projects/hip/docs/markdown/hip_faq.md
@@ -1,12 +1,14 @@
 # FAQ
 
-- [What APIs and features does HIP support ?](#Q1)
-- [What is not supported?](#Q2)
-  - [Run-time features](#run-time-features)
-  - [Kernel language features](#kernel-language-features)
-- [Is HIP a drop-in replacement for CUDA?](#Q3)
-- [What version of CUDA is supported?](#Q4)
-- [What libraries does HIP support?](#Q5)
+<!-- toc -->
+
+- [What APIs and features does HIP support?](#what-apis-and-features-does-hip-support)
+- [What is not supported?](#what-is-not-supported)
+  * [Run-time features](#run-time-features)
+  * [Kernel language features](#kernel-language-features)
+- [Is HIP a drop-in replacement for CUDA?](#is-hip-a-drop-in-replacement-for-cuda)
+- [What specific version of CUDA does HIP support?](#what-specific-version-of-cuda-does-hip-support)
+- [What libraries does HIP support?](#what-libraries-does-hip-support)
 - [How does HIP compare with OpenCL?](#how-does-hip-compare-with-opencl)
 - [What hardware does HIP support?](#what-hardware-does-hip-support)
 - [Does Hipify automatically convert all source code?](#does-hipify-automatically-convert-all-source-code)
@@ -17,11 +19,16 @@
 - [Can I develop HIP code on an AMD HCC platform?](#can-i-develop-hip-code-on-an-amd-hcc-platform)
 - [Can a HIP binary run on both AMD and Nvidia platforms?](#can-a-hip-binary-run-on-both-amd-and-nvidia-platforms)
 - [What's the difference between HIP and hc?](#whats-the-difference-between-hip-and-hc)
+- [On HCC, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ?](#on-hcc-can-i-link-hip-code-with-host-code-compiled-with-another-compiler-such-as-gcc-icc-or-clang-)
 - [HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?](#hip-detected-my-platform-hcc-vs-nvcc-incorrectly---what-should-i-do)
+- [Can I install both CUDA SDK and HCC on same machine?](#can-i-install-both-cuda-sdk-and-hcc-on-same-machine)
+- [How do I trace HIP application flow?](#how-do-i-trace-hip-application-flow)
+  * [Using CodeXL markers for HIP Functions](#using-codexl-markers-for-hip-functions)
+  * [Using HIP_TRACE_API](#using-hip_trace_api)
 
+<!-- tocstop -->
 
-
-### What APIs and features does HIP support?<a name="Q1"></a>
+### What APIs and features does HIP support?
 HIP provides the following:
 - Devices (hipSetDevice(), hipGetDeviceProperties(), etc.)
 - Memory management (hipMalloc(), hipMemcpy(), hipFree(), etc.)
@@ -34,7 +41,7 @@ HIP provides the following:
 
 The HIP API documentation describes each API and its limitations, if any, compared with the equivalent CUDA API.
 
-### What is not supported?<a name="Q2"></a>
+### What is not supported?
 #### Run-time features
 - Textures 
 - MemcpyToSymbol functions
@@ -51,19 +58,19 @@ The HIP API documentation describes each API and its limitations, if any, compar
 - PTX assembly (CUDA 4.0)
 - Several kernel features are under development.  See the [HIP Kernel Language](hip_kernel_language.md) for more information.
 
-### Is HIP a drop-in replacement for CUDA?<a name="Q3"></a>
+### Is HIP a drop-in replacement for CUDA?
 No. HIP provides porting tools which do most of the work do convert CUDA code into portable C++ code that uses the HIP APIs.
 Most developers will port their code from CUDA to HIP and then maintain the HIP version. 
 HIP code provides the same performance as coding in native CUDA, plus the benefit that the code can also run on AMD platforms.
 
-### What version of CUDA is supported?<a name="Q4"></a>
+### What specific version of CUDA does HIP support?
 HIP APIs and features do not map to a specific CUDA version.  HIP provides a strong subset of functionality provided in CUDA, and the hipify tools can 
 scan code to identify any unsupported CUDA functions - this is very useful for identifying the specific features required by a given application.
 
 However, we can provide a rough summary of the features included in each CUDA SDK and the support level in HIP:
 
 - CUDA 4.0 and earlier :  
-    - HIP supports CUDA 4.0 except for the limitations described [above](#Q2).
+    - HIP supports CUDA 4.0 except for the limitations described above.
 - CUDA 5.0 : 
     - Dynamic Parallelism (not supported) 
     - cuIpc functions (under development).
@@ -81,7 +88,7 @@ However, we can provide a rough summary of the features included in each CUDA SD
 - CUDA 8.0
     - No new language features.
 
-### What libraries does HIP support? <a name="Q5"></a>
+### What libraries does HIP support?
 HIP includes growing support for the 4 key math libraries using hcBlas, hcFft, hcrng, and hcsparse).
 These offer pointer-based memory interfaces (as opposed to opaque buffers) and can be easily interfaces with other HCC code.  Developers should use conditional compliation if portability to nvcc systems is desired - using calls to cu* routines on one path and hc* routines on the other.  
 
@@ -146,8 +153,13 @@ HIP is a portable C++ language that supports a strong subset of the CUDA run-tim
 A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time, C++ kernel-launch APIs (parallel_for_each), C++ kernel language, and several memory-management options, including pointers, arrays and array_view (with implicit data synchronization). It's intended to be a leading indicator of the ISO C++ standard.
 
 
-### HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?
+### On HCC, can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ?
+Yes!  HIP/HCC generates the object code which conforms to the GCC ABI, and also links with libstdc++.  This means you can compile host code with the compiler of your choice and link this
+with GPU code compiler with HIP.  Larger projects often contain a mixture of accelerator code (initially written in CUDA with nvcc) plus host code (compiled with gcc, icc, or clang).   These projects
+can convert the accelerator code to HIP, compile that code with hipcc, and link with object code from the preferred compiler.
 
+
+### HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?
 HIP will set the platform to HCC if it sees that the AMD graphics driver is installed and has detected an AMD GPU.
 Sometimes this isn't what you want - you can force HIP to recognize the platform by setting HIP_PLATFORM to hcc (or nvcc)
 ```
@@ -157,3 +169,42 @@ export HIP_PLATFORM=hcc
 One symptom of this problem is the message "error: 'unknown error'(11) at square.hipref.cpp:56".  This can occur if you have a CUDA installation on an AMD platform, and HIP incorrectly detects the platform as nvcc.  HIP may be able to compile the application using the nvcc tool-chain, but will generate this error at runtime since the platform does not have a CUDA device. The fix is to set HIP_PLATFORM=hcc and rebuild the issue. 
 
 If you see issues related to incorrect platform detection, please file an issue with the GitHub issue tracker so we can improve HIP's platform detection logic.
+
+### Can I install both CUDA SDK and HCC on same machine?
+Yes. You can use HIP_PLATFORM to choose which path hipcc targets.  This configuration can be useful when using HIP to develop an application which is portable to both AMD and NVIDIA.
+
+### How do I trace HIP application flow?
+#### Using CodeXL markers for HIP Functions
+HIP can generate markers at function being/end which are displayed on the CodeXL timeline view.
+To do this, you need to install ROCm-Profiler and enable HIP to generate the markers:
+
+1. Install ROCm-Profiler
+Installing HIP from the [rocm](http://gpuopen.com/getting-started-with-boltzmann-components-platforms-installation/) pre-built packages, installs the ROCm-Profiler as well.
+Alternatively, you can build ROCm-Profiler using the instructions [here](https://github.com/RadeonOpenCompute/ROCm-Profiler#building-the-rocm-profiler).
+
+2. Build HIP with ATP markers enabled
+HIP pre-built packages are enabled with ATP marker support by default.
+To enable ATP marker support when building HIP from source, use the option ```-DCOMPILE_HIP_ATP_MARKER=1``` during the cmake configure step.
+
+3. Set HIP_ATP_MARKER
+```shell
+export HIP_ATP_MARKER=1
+```
+
+4. Recompile the target application
+
+5. Run with profiler enabled to generate ATP file.
+```shell
+# Use profile to generate timeline view:
+/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>
+```
+
+#### Using HIP_TRACE_API
+You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided
+by the HIP_DB switch. For example:
+```shell
+# Trace to stderr showing being/end of each function (with arguments) + intermediate debug trace during the execution of each function.
+HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp
+```
+
+Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors.
diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md
index ec2530e46e..88f0237706 100644
--- a/projects/hip/docs/markdown/hip_kernel_language.md
+++ b/projects/hip/docs/markdown/hip_kernel_language.md
@@ -1,38 +1,40 @@
-**Table of Contents**  
+## Table of Contents
+
+<!-- toc -->
 
 - [Introduction](#introduction)
 - [Function-Type Qualifiers](#function-type-qualifiers)
-  - [`__device__`](#__device__)
-  - [`__global__`](#__global__)
-  - [`__host__`](#__host__)
+  * [`__device__`](#__device__)
+  * [`__global__`](#__global__)
+  * [`__host__`](#__host__)
 - [Calling `__global__` Functions](#calling-__global__-functions)
 - [Kernel-Launch Example](#kernel-launch-example)
 - [Variable-Type Qualifiers](#variable-type-qualifiers)
-  - [`__constant__`](#__constant__)
-  - [`__shared__`](#__shared__)
-  - [`__managed__`](#__managed__)
-  - [`__restrict__`](#__restrict__)
+  * [`__constant__`](#__constant__)
+  * [`__shared__`](#__shared__)
+  * [`__managed__`](#__managed__)
+  * [`__restrict__`](#__restrict__)
 - [Built-In Variables](#built-in-variables)
-  - [Coordinate Built-Ins](#coordinate-built-ins)
-  - [warpSize](#warpsize)
+  * [Coordinate Built-Ins](#coordinate-built-ins)
+  * [warpSize](#warpsize)
 - [Vector Types](#vector-types)
-  - [Short Vector Types](#short-vector-types)
-  - [dim3](#dim3)
+  * [Short Vector Types](#short-vector-types)
+  * [dim3](#dim3)
 - [Memory-Fence Instructions](#memory-fence-instructions)
 - [Synchronization Functions](#synchronization-functions)
 - [Math Functions](#math-functions)
-  - [Single Precision Mathematical Functions](#single-precision-mathematical-functions)
-  - [Double Precision Mathematical Functions](#double-precision-mathematical-functions)
-  - [Integer Intrinsics](#integer-intrinsics)
-  - [Floating-point Intrinsics](#floating-point-intrinsics)
+  * [Single Precision Mathematical Functions](#single-precision-mathematical-functions)
+  * [Double Precision Mathematical Functions](#double-precision-mathematical-functions)
+  * [Integer Intrinsics](#integer-intrinsics)
+  * [Floating-point Intrinsics](#floating-point-intrinsics)
 - [Texture Functions](#texture-functions)
 - [Surface Functions](#surface-functions)
 - [Timer Functions](#timer-functions)
 - [Atomic Functions](#atomic-functions)
-  - [Caveats and Features Under-Development:](#caveats-and-features-under-development)
+  * [Caveats and Features Under-Development:](#caveats-and-features-under-development)
 - [Warp Cross-Lane Functions](#warp-cross-lane-functions)
-  - [Warp Vote and Ballot Functions](#warp-vote-and-ballot-functions)
-  - [Warp Shuffle Functions](#warp-shuffle-functions)
+  * [Warp Vote and Ballot Functions](#warp-vote-and-ballot-functions)
+  * [Warp Shuffle Functions](#warp-shuffle-functions)
 - [Profiler Counter Function](#profiler-counter-function)
 - [Assert](#assert)
 - [Printf](#printf)
@@ -43,7 +45,7 @@
 - [In-Line Assembly](#in-line-assembly)
 - [C++ Support](#c-support)
 
-
+<!-- tocstop -->
 
 ## Introduction
 
@@ -628,12 +630,12 @@ The compiler ensures that the kernel uses fewer registers than both allowed maxi
 
 HIP/hcc will parse the `launch_bounds` attribute but silently ignores the performance hint. Full support is under development.
 
-The hcc compiler does not support the "--maxregcount" option like nvcc.  Instead, users are encouraged to use the hip_launch_bounds directive since the parameters are more intuitive and portable than
+Unlike nvcc, hcc does not support the "--maxregcount" option.  Instead, users are encouraged to use the hip_launch_bounds directive since the parameters are more intuitive and portable than
 micro-architecture details like registers, and also the directive allows per-kernel control rather than an entire file.  hip_launch_bounds works on both hcc and nvcc targets.
 
 
 ## Register Keyword
-The register keyword affects code generation in neither nvcc nor hcc.  It’s deprecated in standard C++, so hcc will generate a warning. (nvcc silently ignores use of this keyword.) To disable the warning, you can pass the option `-Wno-deprecated-register` to hcc.
+The register keyword is deprecated in C++, and is silently ignored by both nvcc and hcc.  To see warnings, you can pass the option `-Wdeprecated-register` to hcc.
 
 
 ## Pragma Unroll
diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md
index 76cf86aa49..7857e4b983 100644
--- a/projects/hip/docs/markdown/hip_porting_guide.md
+++ b/projects/hip/docs/markdown/hip_porting_guide.md
@@ -4,7 +4,6 @@ the porting of existing CUDA code into the HIP environment.  This section descri
 and provides practical suggestions on how to port CUDA code and work through common issues. 
 
 ## Table of Contents
-=================
 
 <!-- toc -->
 
@@ -41,8 +40,8 @@ and provides practical suggestions on how to port CUDA code and work through com
   * [Debugging hipcc](#debugging-hipcc)
   * [What Does This Error Mean?](#what-does-this-error-mean)
     + [/usr/include/c++/v1/memory:5172:15: error: call to implicitly deleted default constructor of 'std::__1::bad_weak_ptr' throw bad_weak_ptr();](#usrincludecv1memory517215-error-call-to-implicitly-deleted-default-constructor-of-std__1bad_weak_ptr-throw-bad_weak_ptr)
-    + [HIP Environment Variables](#hip-environment-variables)
-    + [Editor Highlighting](#editor-highlighting)
+  * [HIP Environment Variables](#hip-environment-variables)
+  * [Editor Highlighting](#editor-highlighting)
 
 <!-- tocstop -->
 
@@ -461,7 +460,7 @@ hipcc-cmd: /opt/hcc/bin/hcc  -hc -I/opt/hcc/include -stdlib=libc++ -I../../../..
 If you pass a ".cu" file, hcc will attempt to compile it as a Cuda language file. You must tell hcc that it’s in fact a C++ file: use the "-x c++" option.
 
 
-#### HIP Environment Variables
+### HIP Environment Variables
 
 On the HCC path, HIP provides a number of environment variables that control the behavior of HIP.  Some of these are useful for appliction development (for example HIP_VISIBLE_DEVICES, HIP_LAUNCH_BLOCKING),
 some are useful for performance tuning or experimentation (for example HIP_STAGING*), and some are useful for debugging (HIP_DB).  You can see the environment variables supported by HIP as well as
@@ -484,5 +483,5 @@ HIP_DISABLE_HW_COPY_DEP        =  1 : Disable HW dependencies before copy comman
 ```
 
 
-#### Editor Highlighting
+### Editor Highlighting
 See the utils/vim or utils/gedit directories to add handy highlighting to hip files.
diff --git a/projects/hip/docs/markdown/hip_terms.md b/projects/hip/docs/markdown/hip_terms.md
index c1a40e0573..cc32757047 100644
--- a/projects/hip/docs/markdown/hip_terms.md
+++ b/projects/hip/docs/markdown/hip_terms.md
@@ -1,11 +1,3 @@
-<!-- START doctoc generated TOC please keep comment here to allow auto update -->
-<!-- DON'T EDIT THIS SECTION, INSTEAD RE-RUN doctoc TO UPDATE -->
-**Table of Contents**  *generated with [DocToc](https://github.com/thlorenz/doctoc)*
-
-- [Table Comparing Syntax for Different Compute APIs](#table-comparing-syntax-for-different-compute-apis)
-
-<!-- END doctoc generated TOC please keep comment here to allow auto update -->
-
 # Table Comparing Syntax for Different Compute APIs
 
 |Term|CUDA|HIP|HC|C++AMP|OpenCL|
@@ -46,5 +38,3 @@
 2. The indexing functions (starting with `thread-index`) show the terminology for a 1D grid.  Some APIs use reverse order of xyz / 012 indexing for 3D grids.
 3. HC allows tile dimensions to be specified at runtime while C++AMP requires that tile dimensions be specified at compile-time.  Thus hc syntax for tile dims is `t_ext.tile_dim[0]` while C++AMP is t_ext.tile_dim0.
 
-
-
diff --git a/projects/hip/docs/markdown/hip_terms2.md b/projects/hip/docs/markdown/hip_terms2.md
index 6807338b8d..82174405cd 100644
--- a/projects/hip/docs/markdown/hip_terms2.md
+++ b/projects/hip/docs/markdown/hip_terms2.md
@@ -1,12 +1,3 @@
-<!-- START doctoc generated TOC please keep comment here to allow auto update -->
-<!-- DON'T EDIT THIS SECTION, INSTEAD RE-RUN doctoc TO UPDATE -->
-**Table of Contents**  *generated with [DocToc](https://github.com/thlorenz/doctoc)*
-
-- [Terms used in HIP Documentation](#terms-used-in-hip-documentation)
-
-<!-- END doctoc generated TOC please keep comment here to allow auto update -->
-
-
 # Terms used in HIP Documentation
 
 - host,  host cpu : Executes the HIP runtime API and is capable of initiating kernel launches to one or more devices.
@@ -24,6 +15,3 @@ The default device can be set with hipSetDevice.
 - hipify - tool to convert CUDA(R) code to portable C++ code.
 - hipconfig - tool to report various confoguration properties of the target platform.
 
-
-
-
diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index 635b65f384..405e2de67c 100644
--- a/projects/hip/include/hcc_detail/hip_hcc.h
+++ b/projects/hip/include/hcc_detail/hip_hcc.h
@@ -351,7 +351,9 @@ public:
         _last_copy_signal(NULL),
         _signalCursor(0),
         _oldest_live_sig_id(1),
-        _stream_sig_id(0)
+        _stream_sig_id(0),
+        _kernelCnt(0),
+        _signalCnt(0)
     {
         _signalPool.resize(HIP_STREAM_SIGNALS > 0 ? HIP_STREAM_SIGNALS : 1);
     };
@@ -362,7 +364,6 @@ public:
 
     ihipStreamCriticalBase_t<StreamMutex>  * mlock() { LockedBase<MUTEX_TYPE>::lock(); return this;};
 
-
 public:
     // Critical Data:
     ihipCommand_t               _last_command_type;  // type of the last command
@@ -378,8 +379,10 @@ public:
     int                         _signalCursor;
     SIGSEQNUM                   _oldest_live_sig_id; // oldest live seq_id, anything < this can be allocated.
     std::deque<ihipSignal_t>    _signalPool;   // Pool of signals for use by this stream.
-
-
+    uint32_t                    _signalCnt;    // Count of inflight commands using signals from the signal pool.   
+                                               // Each copy may use 1-2 signals depending on command transitions: 
+                                               //   2 are required if a barrier packet is inserted.
+    uint32_t                    _kernelCnt;    // Count of inflight kernels in this stream.  Reset at ::wait().
     SIGSEQNUM                   _stream_sig_id;      // Monotonically increasing unique signal id.
 };
 
@@ -393,7 +396,6 @@ typedef LockedAccessor<ihipStreamCritical_t> LockedAccessor_StreamCrit_t;
 class ihipStream_t {
 public:
 typedef uint64_t SeqNum_t ;
-
     ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags);
     ~ihipStream_t();
 
@@ -420,7 +422,7 @@ typedef uint64_t SeqNum_t ;
 
 
     // Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function.
-    SIGSEQNUM            lastCopySeqId (LockedAccessor_StreamCrit_t &crit) { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; };
+    SIGSEQNUM            lastCopySeqId (LockedAccessor_StreamCrit_t &crit) const { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; };
     ihipSignal_t *       allocSignal (LockedAccessor_StreamCrit_t &crit);
 
 
@@ -440,14 +442,12 @@ private:
     // Critical Data.  THis MUST be accessed through LockedAccessor_StreamCrit_t
     ihipStreamCritical_t        _criticalData;
 
-    // Array of dependency completion_future.
-    std::vector<hc::completion_future> _depFutures;
-
 private:
-    void                        enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal);
-    void                        enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal);
+    void                        enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal);
     void                        waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal);
 
+
+
     // The unsigned return is hipMemcpyKind
     unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem);
     void setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent);
diff --git a/projects/hip/packaging/hip_hcc.txt b/projects/hip/packaging/hip_hcc.txt
index 661d77e530..556bade072 100644
--- a/projects/hip/packaging/hip_hcc.txt
+++ b/projects/hip/packaging/hip_hcc.txt
@@ -20,12 +20,12 @@ set(CPACK_PACKAGE_FILE_NAME ${CPACK_PACKAGE_NAME}-${CPACK_PACKAGE_VERSION_MAJOR}
 set(CPACK_GENERATOR "TGZ;DEB;RPM")
 set(CPACK_BINARY_DEB "ON")
 set(CPACK_DEBIAN_PACKAGE_CONTROL_EXTRA "${PROJECT_BINARY_DIR}/postinst;${PROJECT_BINARY_DIR}/prerm")
-set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), hcc_lc (= @HCC_VERSION@), rocm-profiler")
+set(CPACK_DEBIAN_PACKAGE_DEPENDS "hip_base (= ${CPACK_PACKAGE_VERSION}), hcc_lc (= @HCC_VERSION@)")
 set(CPACK_BINARY_RPM "ON")
 set(CPACK_RPM_PACKAGE_ARCHITECTURE "x86_64")
 set(CPACK_RPM_POST_INSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/postinst")
 set(CPACK_RPM_PRE_UNINSTALL_SCRIPT_FILE "${PROJECT_BINARY_DIR}/prerm")
 set(CPACK_RPM_PACKAGE_AUTOREQPROV " no")
-set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, hcc_lc = @HCC_VERSION@, rocm-profiler")
+set(CPACK_RPM_PACKAGE_REQUIRES "hip_base = ${CPACK_PACKAGE_VERSION}, hcc_lc = @HCC_VERSION@")
 set(CPACK_SOURCE_GENERATOR "TGZ")
 include(CPack)
diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp
index f16ae715c1..cfc285427c 100644
--- a/projects/hip/src/hip_device.cpp
+++ b/projects/hip/src/hip_device.cpp
@@ -184,7 +184,7 @@ hipError_t hipDeviceReset(void)
  */
 hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device)
 {
-    HIP_INIT_API(attr, device);
+    HIP_INIT_API(pi, attr, device);
 
     hipError_t e = hipSuccess;
 
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 70890d9374..c6c8691419 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -67,7 +67,7 @@ int HIP_OPTIMAL_MEM_TRANSFER = 0; //ENV Variable to test different memory transf
 int HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING = 0;
 int HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE = 0;
 int HIP_D2H_MEM_TRANSFER_THRESHOLD = 0;
-int HIP_STREAM_SIGNALS = 2;  /* number of signals to allocate at stream creation */
+int HIP_STREAM_SIGNALS = 32;  /* number of signals to allocate at stream creation */
 int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */
 
 
@@ -146,6 +146,9 @@ ihipStream_t::~ihipStream_t()
 
 
 //---
+//TODO - this function is dangerous since it does not propertly account
+//for younger commands which may be depending on the signals we are reclaiming.
+//Will fix when we move to HCC management of copy signals.
 void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum)
 {
     LockedAccessor_StreamCrit_t crit(_criticalData);
@@ -180,17 +183,19 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
     if (! assertQueueEmpty) {
         tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this);
         _av.wait();
-    }
+    } 
+     
     if (crit->_last_copy_signal) {
         tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, lastCopySeqId(crit) );
         this->waitCopy(crit, crit->_last_copy_signal);
     }
 
+    crit->_kernelCnt = 0;
+
     // Reset the stream to "empty" - next command will not set up an inpute dependency on any older signal.
     crit->_last_command_type = ihipCommandCopyH2D;
     crit->_last_copy_signal = NULL;
-
-    _depFutures.clear();
+//    crit->_signalCnt = 0;
 }
 
 
@@ -292,6 +297,8 @@ ihipDevice_t * ihipStream_t::getDevice() const
     return ::getDevice(_device_index);
 };
 
+#define HIP_NUM_SIGNALS_PER_STREAM 32
+
 
 //---
 // Allocate a new signal from the signal pool.
@@ -300,8 +307,17 @@ ihipDevice_t * ihipStream_t::getDevice() const
 ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit)
 {
     int numToScan = crit->_signalPool.size();
+    crit->_signalCnt++;
+    if(crit->_signalCnt == HIP_STREAM_SIGNALS){
+        this->wait(crit);
+        crit->_signalCnt = 0;
+    }
+
+    return &crit->_signalPool[crit->_signalCnt];
+
     do {
         auto thisCursor = crit->_signalCursor;
+
         if (++crit->_signalCursor == crit->_signalPool.size()) {
             crit->_signalCursor = 0;
         }
@@ -338,7 +354,7 @@ ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit)
 
 
 //---
-void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal)
+void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal)
 {
 
     // Obtain the write index for the command queue
@@ -356,44 +372,9 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal)
     //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
     barrier->header = header;
 
-    barrier->dep_signal[0] = depSignal->_hsa_signal;
+    barrier->dep_signal[0].handle = depSignal ? depSignal->_hsa_signal.handle: 0;
 
-    barrier->completion_signal.handle = 0;
-
-    // TODO - check queue overflow, return error:
-    // Increment write index and ring doorbell to dispatch the kernel
-    hsa_queue_store_write_index_relaxed(queue, index+1);
-    hsa_signal_store_relaxed(queue->doorbell_signal, index);
-}
-
-void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal)
-{
-
-    // Obtain the write index for the command queue
-    uint64_t index = hsa_queue_load_write_index_relaxed(queue);
-    const uint32_t queueMask = queue->size - 1;
-
-    // Define the barrier packet to be at the calculated queue index address
-    hsa_barrier_and_packet_t* barrier = &(((hsa_barrier_and_packet_t*)(queue->base_address))[index&queueMask]);
-    memset(barrier, 0, sizeof(hsa_barrier_and_packet_t));
-
-    // setup header
-    uint16_t header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE;
-    header |= 1 << HSA_PACKET_HEADER_BARRIER;
-    //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE;
-    //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE;
-    barrier->header = header;
-
-    barrier->dep_signal[0].handle = 0;
-    barrier->dep_signal[1].handle = 0;
-    barrier->dep_signal[2].handle = 0;
-    barrier->dep_signal[3].handle = 0;
-    barrier->dep_signal[4].handle = 0;
-
-    hsa_signal_t signal;
-    hsa_signal_create(1, 0, NULL, &signal);
-    *depSignal = signal;
-    barrier->completion_signal = signal;
+    barrier->completion_signal.handle = completionSignal ? completionSignal->_hsa_signal.handle : 0;
 
     // TODO - check queue overflow, return error:
     // Increment write index and ring doorbell to dispatch the kernel
@@ -402,6 +383,7 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal)
 }
 
 
+int HIP_NUM_KERNELS_INFLIGHT = 128;
 
 //--
 //When the commands in a stream change types (ie kernel command follows a data command,
@@ -414,6 +396,12 @@ bool ihipStream_t::lockopen_preKernelCommand()
     LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/);
 
     bool addedSync = false;
+
+    if(crit->_kernelCnt > HIP_NUM_KERNELS_INFLIGHT){
+        this->wait(crit);
+       crit->_kernelCnt = 0;
+    }
+    crit->_kernelCnt++;
     // If switching command types, we need to add a barrier packet to synchronize things.
     if (crit->_last_command_type != ihipCommandKernel) {
         if (crit->_last_copy_signal) {
@@ -421,7 +409,7 @@ bool ihipStream_t::lockopen_preKernelCommand()
 
             hsa_queue_t * q =  (hsa_queue_t*)_av.get_hsa_queue();
             if (HIP_DISABLE_HW_KERNEL_DEP == 0) {
-                this->enqueueBarrier(q, crit->_last_copy_signal);
+                this->enqueueBarrier(q, crit->_last_copy_signal, NULL);
                 tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n",
                         this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel], crit->_last_copy_signal->_sig_id)
 
@@ -456,14 +444,12 @@ void ihipStream_t::lockclose_postKernelCommand(hc::completion_future &kernelFutu
 //---
 // Called whenever a copy command is set to the stream.
 // Examines the last command sent to this stream and returns a signal to wait on, if required.
-int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *lastCopy, hsa_signal_t *waitSignal, ihipCommand_t copyType)
+int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *copyCompletionSignal, hsa_signal_t *waitSignal, ihipCommand_t copyType)
 {
     int needSync = 0;
 
     waitSignal->handle = 0;
 
-    //_mutex.lock(); // will be unlocked in postCopyCommand
-
     // If switching command types, we need to add a barrier packet to synchronize things.
     if (FORCE_SAMEDIR_COPY_DEP || (crit->_last_command_type != copyType)) {
 
@@ -472,16 +458,10 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
             tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on prev kernel)\n",
                     this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType]);
             needSync = 1;
-            this->enqueueBarrier(static_cast<hsa_queue_t*>(_av.get_hsa_queue()), waitSignal);
-//            hsa_signal_t *hsaSignal = (static_cast<hsa_signal_t*> (crit->_last_kernel_future.get_native_handle()));
-/*            if (hsaSignal) {
-                // Keep reference to the kernel future in order to keep the
-                // dependent signal alive.
-                _depFutures.push_back(crit->_last_kernel_future);
-                *waitSignal = * hsaSignal;
-            } else {
-                assert(0); // if NULL signal, and we return 1, hsa_amd_memory_copy_async will fail.  Confirm this never happens.
-            }*/
+            ihipSignal_t *depSignal = allocSignal(crit);
+            hsa_signal_store_relaxed(depSignal->_hsa_signal,1);
+            this->enqueueBarrier(static_cast<hsa_queue_t*>(_av.get_hsa_queue()), NULL, depSignal);
+            *waitSignal = depSignal->_hsa_signal;
         } else if (crit->_last_copy_signal) {
             needSync = 1;
             tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on other copy #%lu)\n",
@@ -504,7 +484,7 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
         crit->_last_command_type = copyType;
     }
 
-    crit->_last_copy_signal = lastCopy;
+    crit->_last_copy_signal = copyCompletionSignal;
 
     return needSync;
 }
@@ -1135,6 +1115,8 @@ void ihipInit()
     READ_ENV_I(release, HIP_DISABLE_HW_KERNEL_DEP, 0, "Disable HW dependencies before kernel commands  - instead wait for dependency on host. -1 means ignore these dependencies. (debug mode)");
     READ_ENV_I(release, HIP_DISABLE_HW_COPY_DEP, 0, "Disable HW dependencies before copy commands  - instead wait for dependency on host. -1 means ifnore these dependencies (debug mode)");
 
+    READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Number of kernels per stream ");
+
     if (HIP_OPTIMAL_MEM_TRANSFER && !HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING) {
         HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING= MEMCPY_H2D_DIRECT_VS_STAGING_COPY_THRESHOLD;
         fprintf (stderr, "warning: env var HIP_OPTIMAL_MEM_TRANSFER=0x%x but HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING=0.Using default value for this.\n", HIP_OPTIMAL_MEM_TRANSFER);
@@ -1270,28 +1252,6 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
     }
 }
 
-// HIP uses only 64 kernels. If the performance decrease, add more
-uint32_t kernelCount = 0;
-std::vector<hc::completion_future*> vCF(64);
-
-void incKernelCnt(hc::completion_future *cf){
-    vCF[kernelCount] = cf;
-    kernelCount++;
-}
-
-void decKernelCnt(){
-    if(kernelCount > 63){
-        uint32_t len = kernelCount;
-        for(uint32_t i =0;i<len;i++){
-            if(vCF[i] != NULL){
-                vCF[i]->wait();
-            }
-            delete vCF[i];
-            vCF[i] = NULL;
-            kernelCount--;
-        }
-    }
-}
 
 // TODO - data-up to data-down:
 // Called just before a kernel is launched from hipLaunchKernel.
@@ -1321,7 +1281,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
 //    *av = &stream->_av;
     lp->av = &stream->_av;
     lp->cf = new hc::completion_future;
-    incKernelCnt(lp->cf);
 //    lp->av = static_cast<void*>(av);
 //    lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
     return (stream);
@@ -1351,7 +1310,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri
 //    *av = &stream->_av;
     lp->av = &stream->_av;
     lp->cf = new hc::completion_future;
-    incKernelCnt(lp->cf);
 //    lp->av = static_cast<void*>(av);
 //    lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
     return (stream);
@@ -1382,7 +1340,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri
 //    *av = &stream->_av;
     lp->av = &stream->_av;
     lp->cf = new hc::completion_future;
-    incKernelCnt(lp->cf);
 //    lp->av = static_cast<void*>(av);
 //    lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
     return (stream);
@@ -1413,7 +1370,6 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
 //    *av = &stream->_av;
     lp->av = &stream->_av;
     lp->cf = new hc::completion_future;
-    incKernelCnt(lp->cf);
 //    lp->av = static_cast<void*>(av);
 //    lp->cf = static_cast<void*>(malloc(sizeof(hc::completion_future)));
     return (stream);
@@ -1426,7 +1382,6 @@ void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp)
 {
 //    stream->lockclose_postKernelCommand(cf);
     stream->lockclose_postKernelCommand(*lp.cf);
-    decKernelCnt();
     if (HIP_LAUNCH_BLOCKING) {
         tprintf(DB_SYNC, " stream:%p LAUNCH_BLOCKING for kernel completion\n", stream);
     }
@@ -1619,7 +1574,8 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
                 hc::am_copy(dst, src, sizeBytes);
 #endif
             }
-        }else{
+        } else {
+            // This is H2D copy, and source is pinned host memory : we can copy directly w/o using staging buffer.
             hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
             hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
 
@@ -1657,10 +1613,6 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
                 {
                     device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL);
                 }
-                if(crit->_last_command_type == ihipCommandKernel){
-                    std::cout<<"Destroying depSignal MemcpySync"<<std::endl;
-                    hsa_signal_destroy(depSignal);
-                }
                 // The copy completes before returning so can reset queue to empty:
                 this->wait(crit, true);
 
@@ -1673,7 +1625,8 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
                 hc::am_copy(dst, src, sizeBytes);
 #endif
             }
-        }else{
+        } else {
+            // This is D2H copy, and destination is pinned host memory : we can copy directly w/o using staging buffer.
             hsa_agent_t dstAgent = *(static_cast<hsa_agent_t*>(dstPtrInfo._acc.get_hsa_agent()));
             hsa_agent_t srcAgent = *(static_cast<hsa_agent_t*>(srcPtrInfo._acc.get_hsa_agent()));
 
@@ -1743,9 +1696,6 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const
             throw ihipException(hipErrorInvalidValue);
         }
     }
-    if(crit->_last_command_type == ihipCommandKernel){
-        hsa_signal_destroy(depSignal); 
-    }
 
 }
 
@@ -1820,9 +1770,6 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig
 
             hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal);
 
-            if (crit->_last_command_type == ihipCommandKernel) {
-                hsa_signal_destroy(depSignal);
-            }
 
             if (hsa_status == HSA_STATUS_SUCCESS) {
                 if (HIP_LAUNCH_BLOCKING) {