From adfaa47840118b674896e4144f175e0dc7577245 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Mon, 25 Jul 2016 14:53:15 +0530
Subject: [PATCH 01/22] Documentation updates

Change-Id: Ia624d86915c4c96da0ac0242f767135f30ff73c6


[ROCm/hip commit: 90a3e4e1da1c22a4b4d0254d8d5af0e20ab0bbc4]
---
 projects/hip/INSTALL.md                       | 66 ++++--------------
 projects/hip/clang-hipify/README.md           | 13 +++-
 projects/hip/docs/markdown/hip_faq.md         | 69 +++++++++++++++----
 .../hip/docs/markdown/hip_kernel_language.md  | 42 +++++------
 .../hip/docs/markdown/hip_porting_guide.md    |  9 ++-
 projects/hip/docs/markdown/hip_terms.md       | 10 ---
 projects/hip/docs/markdown/hip_terms2.md      | 12 ----
 7 files changed, 106 insertions(+), 115 deletions(-)

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/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/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md
index e71dcdfb36..31844346c0 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 version of CUDA is supported?](#what-version-of-cuda-is-supported)
+- [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)
@@ -18,10 +20,13 @@
 - [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)
 - [HIP detected my platform (hcc vs nvcc) incorrectly - what should I do?](#hip-detected-my-platform-hcc-vs-nvcc-incorrectly---what-should-i-do)
+- [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 +39,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,12 +56,12 @@ 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 version of CUDA is supported?
 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.
 
@@ -81,7 +86,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.  
 
@@ -147,7 +152,6 @@ A C++ dialect, hc is supported by the AMD HCC compiler. It provides C++ run time
 
 
 ### 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 +161,40 @@ 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.
+
+
+### 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 478de7cda6..34778628da 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
 
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.
 
-
-
-

From 69aed2c6ea263cbbf52b2a9e27b55dc62e4dbcc1 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 26 Jul 2016 09:22:59 -0500
Subject: [PATCH 02/22] Added re-fix for memcpy kernel sync 1. The patch uses
 HIP signal pools to sync between copy and kernel commands 2. The
 hsa_signal_create is removed 3. Left the redundant enqueueBarrier method just
 in case

Change-Id: I3dff3e8ee57fff3cd49bec802ff735ed128e5ca1


[ROCm/hip commit: 4bdf26a82efb863d23e5ffe8f62e00050f34374e]
---
 projects/hip/include/hcc_detail/hip_hcc.h |  2 +-
 projects/hip/src/hip_hcc.cpp              | 28 +++++++++--------------
 2 files changed, 12 insertions(+), 18 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index 635b65f384..cf6c705082 100644
--- a/projects/hip/include/hcc_detail/hip_hcc.h
+++ b/projects/hip/include/hcc_detail/hip_hcc.h
@@ -444,7 +444,7 @@ private:
     std::vector<hc::completion_future> _depFutures;
 
 private:
-    void                        enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal);
+    void                        enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal);
     void                        enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal);
     void                        waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal);
 
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 70890d9374..46f0b2674f 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -338,7 +338,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,9 +356,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;
+    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
@@ -390,10 +390,7 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal)
     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 = *depSignal;
 
     // TODO - check queue overflow, return error:
     // Increment write index and ring doorbell to dispatch the kernel
@@ -421,7 +418,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)
 
@@ -472,7 +469,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);
+            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;
 //            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
@@ -1272,7 +1272,7 @@ 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);
+std::vector<hc::completion_future*> vCF(1024);
 
 void incKernelCnt(hc::completion_future *cf){
     vCF[kernelCount] = cf;
@@ -1280,7 +1280,7 @@ void incKernelCnt(hc::completion_future *cf){
 }
 
 void decKernelCnt(){
-    if(kernelCount > 63){
+    if(kernelCount > 1023){
         uint32_t len = kernelCount;
         for(uint32_t i =0;i<len;i++){
             if(vCF[i] != NULL){
@@ -1743,9 +1743,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 +1817,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) {

From b909cbb2a3ad5d3fa40cc837a21c8e790cd10d56 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Mon, 25 Jul 2016 10:15:02 -0500
Subject: [PATCH 03/22] Make HCC ignore register keyword

- (previously would emit a warning)
- Also tweak documentation.

Change-Id: I0f4f00f82f8cc53d420112570f2d7675535e6aea


[ROCm/hip commit: 0a74e00ac88ff7e18e81a0e1bdf34ec4b2ff9303]
---
 projects/hip/RELEASE.md                                      | 5 ++++-
 projects/hip/bin/hipcc                                       | 1 +
 .../markdown/CUDA_Runtime_API_functions_supported_by_HIP.md  | 2 +-
 projects/hip/docs/markdown/hip_faq.md                        | 2 +-
 projects/hip/docs/markdown/hip_kernel_language.md            | 4 ++--
 5 files changed, 9 insertions(+), 5 deletions(-)

diff --git a/projects/hip/RELEASE.md b/projects/hip/RELEASE.md
index f60dafd4b5..6e6fc60a8d 100644
--- a/projects/hip/RELEASE.md
+++ b/projects/hip/RELEASE.md
@@ -5,7 +5,10 @@ We have attempted to document known bugs and limitations - in particular the [HI
 ===================================================================================================
 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 (>10K) without synchronizing.  
+- Doc updates: Add some more frequently asked questions to FAQ, fix TOC in some files, review.
+- Register keyword now silently ignored on HCC (previously would emit warning).
 ===================================================================================================
 Release:0.90.00
 Date: 2016.06.29
diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc
index ce888fea8f..91ace61cec 100755
--- a/projects/hip/bin/hipcc
+++ b/projects/hip/bin/hipcc
@@ -73,6 +73,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
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..1e578842fa 100644
--- a/projects/hip/docs/markdown/hip_faq.md
+++ b/projects/hip/docs/markdown/hip_faq.md
@@ -56,7 +56,7 @@ No. HIP provides porting tools which do most of the work do convert CUDA code in
 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?<a name="Q4"></a>
 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.
 
diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md
index ec2530e46e..5310147ca5 100644
--- a/projects/hip/docs/markdown/hip_kernel_language.md
+++ b/projects/hip/docs/markdown/hip_kernel_language.md
@@ -628,12 +628,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

From 46cb4cdc7b911a89b3c83c33c7bd407a3904a35b Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Mon, 25 Jul 2016 17:45:31 -0500
Subject: [PATCH 04/22] Doc update for README.md - add more intro text, example

Change-Id: I99b8eaacd6460dfdbdbc8ddba3fe589647d877e7


[ROCm/hip commit: 6ea9d2eacecd1788e6b4e10c5b32569d359a57d6]
---
 projects/hip/README.md | 59 ++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 59 insertions(+)

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:
 

From 3027a1f269d5c9bf99c3eb75586e0e1df34f6d0f Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Tue, 26 Jul 2016 11:51:54 -0500
Subject: [PATCH 05/22] Doc update for FAQ and future RELEASE notes

Change-Id: I7e7c32d9a19fdaea0a0e41f1d4fa4652e53640f2


[ROCm/hip commit: 99bc0f03e161c17894eb8fa8189dfd1b1bf0def7]
---
 projects/hip/RELEASE.md               |  6 +++---
 projects/hip/docs/markdown/hip_faq.md | 11 +++++++++++
 2 files changed, 14 insertions(+), 3 deletions(-)

diff --git a/projects/hip/RELEASE.md b/projects/hip/RELEASE.md
index 6e6fc60a8d..2f1a5e8c92 100644
--- a/projects/hip/RELEASE.md
+++ b/projects/hip/RELEASE.md
@@ -2,14 +2,14 @@
 
 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 (>10K) without synchronizing.  
-- Doc updates: Add some more frequently asked questions to FAQ, fix TOC in some files, review.
+  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.
 ===================================================================================================
+
 Release:0.90.00
 Date: 2016.06.29
 - Support dynamic shared memory allocations
diff --git a/projects/hip/docs/markdown/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md
index 1e578842fa..6a3d46f075 100644
--- a/projects/hip/docs/markdown/hip_faq.md
+++ b/projects/hip/docs/markdown/hip_faq.md
@@ -16,8 +16,10 @@
 - [Can I develop HIP code on an Nvidia CUDA platform?](#can-i-develop-hip-code-on-an-nvidia-cuda-platform)
 - [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)
+- [Can I link HIP code with host code compiled with another compiler such as gcc, icc, or clang ?]
 - [What's the difference between HIP and hc?](#whats-the-difference-between-hip-and-hc)
 - [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?](#Q100)
 
 
 
@@ -146,6 +148,12 @@ 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.
 
 
+### [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.
@@ -157,3 +165,6 @@ 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?]<a name="Q100"></a>
+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.

From 06abacce9835f035610957bbd3fededfad5f87dc Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 26 Jul 2016 13:35:35 -0500
Subject: [PATCH 06/22] removed redundant signal destroy

Change-Id: Icf0cd76b2620d34c87cfb6c7a83049087c0a0bc4


[ROCm/hip commit: fa7933eb910e7e753f9542dd15a2383913f6c366]
---
 projects/hip/src/hip_hcc.cpp | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 46f0b2674f..b89fc43582 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -1657,10 +1657,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);
 

From eccab29360b971398869362c297d02036dcdfc16 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 26 Jul 2016 14:03:51 -0500
Subject: [PATCH 07/22] Signal Fix: Changed global signal count to per stream
 signal count 1. The number of kernels that can use signals are increased to
 128 2. The kernel count is now specific to the stream

Change-Id: Ie6d1aa3f437aad8f08c3333fe48bd3f46e551e60


[ROCm/hip commit: 53d7629a85a6ba28e3c9f89e0aae739e6e1d38bd]
---
 projects/hip/include/hcc_detail/hip_hcc.h |  2 +-
 projects/hip/src/hip_hcc.cpp              | 35 +++++------------------
 2 files changed, 8 insertions(+), 29 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index cf6c705082..8637f62457 100644
--- a/projects/hip/include/hcc_detail/hip_hcc.h
+++ b/projects/hip/include/hcc_detail/hip_hcc.h
@@ -393,7 +393,7 @@ typedef LockedAccessor<ihipStreamCritical_t> LockedAccessor_StreamCrit_t;
 class ihipStream_t {
 public:
 typedef uint64_t SeqNum_t ;
-
+    uint32_t kernelCnt;
     ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags);
     ~ihipStream_t();
 
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index b89fc43582..20deb69aef 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -134,6 +134,7 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsig
     _flags(flags),
     _device_index(device_index)
 {
+    kernelCnt = 0;
     tprintf(DB_SYNC, " streamCreate: stream=%p\n", this);
 };
 
@@ -398,7 +399,7 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal)
     hsa_signal_store_relaxed(queue->doorbell_signal, index);
 }
 
-
+#define HIP_NUM_KERNELS_INFLIGHT 128
 
 //--
 //When the commands in a stream change types (ie kernel command follows a data command,
@@ -411,6 +412,11 @@ bool ihipStream_t::lockopen_preKernelCommand()
     LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/);
 
     bool addedSync = false;
+    if(kernelCnt > HIP_NUM_KERNELS_INFLIGHT){
+        this->wait(crit);
+        kernelCnt = 0;
+    }
+    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) {
@@ -1270,28 +1276,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(1024);
-
-void incKernelCnt(hc::completion_future *cf){
-    vCF[kernelCount] = cf;
-    kernelCount++;
-}
-
-void decKernelCnt(){
-    if(kernelCount > 1023){
-        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 +1305,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 +1334,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 +1364,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 +1394,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 +1406,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);
     }

From 6ff74b77805f26918f1e492f41002fffaa641805 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 26 Jul 2016 17:09:27 -0500
Subject: [PATCH 08/22] Signal Fix: Moved kernel count to critical stream 1.
 Added environment variable HIP_NUM_KERNELS_INFLIGHT 2. Moved kernelcount
 variable inside stream critical section

Change-Id: I51d24d0a2a109467209170de117a6d02ba4e308e


[ROCm/hip commit: 0a31b47e2ea7aac21c3bedc856ffca5365e55a54]
---
 projects/hip/include/hcc_detail/hip_hcc.h |  3 +--
 projects/hip/src/hip_hcc.cpp              | 13 ++++++++-----
 2 files changed, 9 insertions(+), 7 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index 8637f62457..07bbe95742 100644
--- a/projects/hip/include/hcc_detail/hip_hcc.h
+++ b/projects/hip/include/hcc_detail/hip_hcc.h
@@ -379,7 +379,7 @@ public:
     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                    _kernelCnt;
     SIGSEQNUM                   _stream_sig_id;      // Monotonically increasing unique signal id.
 };
 
@@ -393,7 +393,6 @@ typedef LockedAccessor<ihipStreamCritical_t> LockedAccessor_StreamCrit_t;
 class ihipStream_t {
 public:
 typedef uint64_t SeqNum_t ;
-    uint32_t kernelCnt;
     ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags);
     ~ihipStream_t();
 
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 20deb69aef..a319234538 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -134,7 +134,6 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsig
     _flags(flags),
     _device_index(device_index)
 {
-    kernelCnt = 0;
     tprintf(DB_SYNC, " streamCreate: stream=%p\n", this);
 };
 
@@ -181,6 +180,7 @@ 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();
+        crit->_kernelCnt = 0;
     }
     if (crit->_last_copy_signal) {
         tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, lastCopySeqId(crit) );
@@ -399,7 +399,7 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal)
     hsa_signal_store_relaxed(queue->doorbell_signal, index);
 }
 
-#define HIP_NUM_KERNELS_INFLIGHT 128
+int HIP_NUM_KERNELS_INFLIGHT = 128;
 
 //--
 //When the commands in a stream change types (ie kernel command follows a data command,
@@ -412,11 +412,12 @@ bool ihipStream_t::lockopen_preKernelCommand()
     LockedAccessor_StreamCrit_t crit(_criticalData, false/*no unlock at destruction*/);
 
     bool addedSync = false;
-    if(kernelCnt > HIP_NUM_KERNELS_INFLIGHT){
+
+    if(crit->_kernelCnt > HIP_NUM_KERNELS_INFLIGHT){
         this->wait(crit);
-        kernelCnt = 0;
+       crit->_kernelCnt = 0;
     }
-    kernelCnt++;
+    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) {
@@ -1141,6 +1142,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);

From 1acf8919ee517d20a7e5af4c1e880125fb1fa496 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 27 Jul 2016 20:30:04 +0530
Subject: [PATCH 09/22] Update release notes for 0.92.00 release

Change-Id: I9ca588cd0d5d752dc6521e76ba943500eb55525f


[ROCm/hip commit: 059a8d51ba096c58a9d737f721865f8059e2f191]
---
 projects/hip/RELEASE.md | 18 +++++++++++++-----
 1 file changed, 13 insertions(+), 5 deletions(-)

diff --git a/projects/hip/RELEASE.md b/projects/hip/RELEASE.md
index f60dafd4b5..80dc240192 100644
--- a/projects/hip/RELEASE.md
+++ b/projects/hip/RELEASE.md
@@ -3,8 +3,19 @@
 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.
+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
+
+
+===================================================================================================
+
+## Revision History:
 
 ===================================================================================================
 Release:0.90.00
@@ -16,9 +27,6 @@ Date: 2016.06.29
 - Improve error code reporting on nvcc.
 - Add hipPeekAtError for nvcc.
 
-===================================================================================================
-
-## Revision History:
 
 ===================================================================================================
 Release:0.86.00

From c2b4a3936c3b2895b63d24458a99b3a19ac2b1fd Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Wed, 27 Jul 2016 13:48:49 -0500
Subject: [PATCH 10/22] Signal Fix: Added signal limit to allocSignal 1. Did
 not change the logic in allocSignal 2. Added guard to wait on signal limit

Change-Id: I78f29097e6a584b3c3d78319dac19869067bd1fe


[ROCm/hip commit: 1859c6e515ac89b963f646e7105a1ef445c666b7]
---
 projects/hip/include/hcc_detail/hip_hcc.h |  6 ++++--
 projects/hip/src/hip_hcc.cpp              | 10 ++++++++++
 2 files changed, 14 insertions(+), 2 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index 07bbe95742..db51a497e9 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);
     };
@@ -378,7 +380,7 @@ 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;
     uint32_t                    _kernelCnt;
     SIGSEQNUM                   _stream_sig_id;      // Monotonically increasing unique signal id.
 };
diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index a319234538..54a16a56d8 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -293,6 +293,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.
@@ -301,8 +303,16 @@ 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_NUM_SIGNALS_PER_STREAM){
+        crit->_signalCnt = 0;
+        this->wait(crit);
+    }
+
     do {
         auto thisCursor = crit->_signalCursor;
+
         if (++crit->_signalCursor == crit->_signalPool.size()) {
             crit->_signalCursor = 0;
         }

From 07c9fb8b38a736c86c97d6afe8619efc877d818f Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Wed, 27 Jul 2016 16:18:14 -0500
Subject: [PATCH 11/22] Fix API string message for hipDeviceGetAttribute

Change-Id: I30f54627630c8ee835506be8c9921742bb68a43a


[ROCm/hip commit: 4cf0f86e99ae29fa24eff99db06c191aa7729df3]
---
 projects/hip/src/hip_device.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

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;
 

From 91368967bb7b76175ce0c13f8e8723113ccbe9b3 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Wed, 27 Jul 2016 18:31:11 -0500
Subject: [PATCH 12/22] Cleanup sync code.

Remove dead depFutures, enqueueBarrier call.
Rename some parms to reflect usage.
Add comments to better explain tricky parts of sync code.

Change-Id: I763296421d9c2b3b58fc8cef5f010b12ab49553c


[ROCm/hip commit: 02dd7a739923f0a55a81ed89d9459586e8b9461a]
---
 projects/hip/include/hcc_detail/hip_hcc.h | 15 +++++-----
 projects/hip/src/hip_hcc.cpp              | 34 ++++++++++-------------
 2 files changed, 22 insertions(+), 27 deletions(-)

diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h
index db51a497e9..405e2de67c 100644
--- a/projects/hip/include/hcc_detail/hip_hcc.h
+++ b/projects/hip/include/hcc_detail/hip_hcc.h
@@ -364,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
@@ -380,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;
-    uint32_t                    _kernelCnt;
+    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.
 };
 
@@ -421,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);
 
 
@@ -441,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, ihipSignal_t *completionSignal);
-    void                        enqueueBarrier(hsa_queue_t* queue, hsa_signal_t *depSignal);
     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/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 54a16a56d8..48d1eebf50 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -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,18 +183,21 @@ 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();
-        crit->_kernelCnt = 0;
+    } else {
+        assert (crit->_kernelCnt == 0);
     }
+     
     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;
 }
 
 
@@ -306,7 +312,6 @@ ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit)
 
     crit->_signalCnt++;
     if(crit->_signalCnt == HIP_NUM_SIGNALS_PER_STREAM){
-        crit->_signalCnt = 0;
         this->wait(crit);
     }
 
@@ -470,14 +475,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)) {
 
@@ -490,15 +493,6 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t
             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;
-//            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.
-            }*/
         } 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",
@@ -521,7 +515,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;
 }
@@ -1611,7 +1605,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()));
 
@@ -1661,7 +1656,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()));
 

From 30447110426bc17b0d7f674abb77071b51e2e8c2 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 28 Jul 2016 22:48:22 -0500
Subject: [PATCH 13/22] Remove dead enqueueBarrier function.

Change-Id: Ib18fe6bd96ce24dbeb342961ddb5721f7d03f2b2


[ROCm/hip commit: 666c227c7d914b6a3562d5a0fe9d6ed423c012af]
---
 projects/hip/src/hip_hcc.cpp | 31 -------------------------------
 1 file changed, 31 deletions(-)

diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 48d1eebf50..9b5657bef1 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -382,37 +382,6 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, i
     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;
-
-    barrier->completion_signal = *depSignal;
-
-    // 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);
-}
 
 int HIP_NUM_KERNELS_INFLIGHT = 128;
 

From 2f53a50718130d35b5bb3a1de192eea2d1d210cf Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Thu, 28 Jul 2016 23:01:35 -0500
Subject: [PATCH 14/22] Signal Fix: The signals in a stream are re-used 1.
 Before, the signal pool is increased depending on the usage 2. After, a
 static number of signals are allocated to the pool Only these are used by hip
 in a stream 3. If the signals required are more than the pool size, the
 stream has to wait to make sure all the signals are available 4. Once they
 are available, the stream can use them 5. Removed HIP_NUM_SIGNALS_PER_STREAM
 because of redundancy with HIP_STREAM_SIGNALS 6. Increased signal count from
 2 to 32. Future Work: Dynamically increase the pool size depending on the
 number of streams allocated by the application. And, null stream should have
 more signals

Change-Id: I6be36e084f26bb04766fabf776c7210aee0f9e91


[ROCm/hip commit: 9062ebcf3af57bc10f9c0a5c1459db8e9fef3576]
---
 projects/hip/src/hip_hcc.cpp | 10 ++++++----
 1 file changed, 6 insertions(+), 4 deletions(-)

diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 9b5657bef1..9af4146891 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 */
 
 
@@ -197,7 +197,7 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty
     // 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;
-    crit->_signalCnt = 0;
+//    crit->_signalCnt = 0;
 }
 
 
@@ -309,12 +309,14 @@ 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_NUM_SIGNALS_PER_STREAM){
+    if(crit->_signalCnt == HIP_STREAM_SIGNALS){
         this->wait(crit);
+        crit->_signalCnt = 0;
     }
 
+    return &crit->_signalPool[crit->_signalCnt];
+
     do {
         auto thisCursor = crit->_signalCursor;
 

From 5ffc0a1e25b737379273047adff61914376fe51b Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Mon, 1 Aug 2016 13:38:47 -0500
Subject: [PATCH 15/22] Remove faulty assert for kernelCnt==0

Change-Id: I8a925c95f48e857c0a31f44561499e90dc6df552


[ROCm/hip commit: f43d02027e38c907c02afb870b282115c9d3c4be]
---
 projects/hip/src/hip_hcc.cpp | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp
index 9af4146891..c6c8691419 100644
--- a/projects/hip/src/hip_hcc.cpp
+++ b/projects/hip/src/hip_hcc.cpp
@@ -183,9 +183,7 @@ 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();
-    } else {
-        assert (crit->_kernelCnt == 0);
-    }
+    } 
      
     if (crit->_last_copy_signal) {
         tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, lastCopySeqId(crit) );

From 8c104aaeebd5c21b8d66a8b346a4c5f74d632e79 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 2 Aug 2016 15:45:46 -0500
Subject: [PATCH 16/22] Added configuration file to be used by hipcc 1. CMake
 will create .hip-config file in bin directory Future Work: Need to make
 changes to hipcc to read the file

Change-Id: Ia7dc48d43787921d5af4ab07d7a5befbcf904465


[ROCm/hip commit: 9c45d9eaedbf7136335eac0884c020aed2694d24]
---
 projects/hip/CMakeLists.txt | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 71901da4a4..85e6ff1c21 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -36,6 +36,8 @@ if(HIP_PLATFORM STREQUAL "hcc")
             set(HCC_HOME $ENV{HCC_HOME} CACHE PATH "Path to which HCC has been installed")
         endif()
     endif()
+    message(STATUS ${HCC_HOME})
+    file(WRITE "./bin/.hip-config" ${HCC_HOME})
     if(IS_ABSOLUTE ${HCC_HOME} AND EXISTS ${HCC_HOME} AND IS_DIRECTORY ${HCC_HOME})
         execute_process(COMMAND ${HCC_HOME}/bin/hcc --version
                         COMMAND cut -d\  -f9

From 7bab8bfb27c77b8eedce018b84bbef05427fa161 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 3 Aug 2016 09:07:58 +0530
Subject: [PATCH 17/22] hip_hcc rpm package does not depend on rocm-profiler

Change-Id: I291c046371e704ce395781d5de3f6430c0dcba7b


[ROCm/hip commit: b2035c98acd3ba4f8a386a43bbbb54a548672892]
---
 projects/hip/packaging/hip_hcc.txt | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

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)

From 960d1ec4cca0adcb24ef32c58036485ccd6403fb Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 3 Aug 2016 09:08:40 +0530
Subject: [PATCH 18/22] Fix hipcc to work correctly when HIP_ATP_MARKER is not
 defined

Change-Id: I7cc525daccf896704e6ccf6d04ed395fda622031


[ROCm/hip commit: e818f2378d1509c07a8f46019e75781b2c13a44e]
---
 projects/hip/bin/hipcc | 12 +++++-------
 1 file changed, 5 insertions(+), 7 deletions(-)

diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc
index ce888fea8f..81d9ec62ec 100755
--- a/projects/hip/bin/hipcc
+++ b/projects/hip/bin/hipcc
@@ -57,9 +57,7 @@ if ($HIP_PLATFORM eq "hcc") {
     $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;
@@ -103,11 +101,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.

From 0229cddf67ae35d1ca230666fa889a5907c90fbb Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 3 Aug 2016 11:25:59 +0530
Subject: [PATCH 19/22] Only create .hipconfig on install

Change-Id: I9c20440023401f5794b33a0165e0918372783a68


[ROCm/hip commit: c38d0466b8c88a5e1db24b5344bcfc9197c30384]
---
 projects/hip/CMakeLists.txt | 9 +++++++--
 1 file changed, 7 insertions(+), 2 deletions(-)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 85e6ff1c21..88b14b0521 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -36,8 +36,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
             set(HCC_HOME $ENV{HCC_HOME} CACHE PATH "Path to which HCC has been installed")
         endif()
     endif()
-    message(STATUS ${HCC_HOME})
-    file(WRITE "./bin/.hip-config" ${HCC_HOME})
     if(IS_ABSOLUTE ${HCC_HOME} AND EXISTS ${HCC_HOME} AND IS_DIRECTORY ${HCC_HOME})
         execute_process(COMMAND ${HCC_HOME}/bin/hcc --version
                         COMMAND cut -d\  -f9
@@ -151,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
@@ -168,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 bin)
 endif()
 
 # Install src, bin, include if necessary

From 8c98963134f896551413fb8a18f07a434ba92cc1 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 3 Aug 2016 11:32:08 +0530
Subject: [PATCH 20/22] hipcc reads .hipconfig

hipcc dies if HCC_VERSION is not same as the one used to build HIP

Change-Id: I180c0108812fe5cb6c5304477557c524a4a61f75


[ROCm/hip commit: 703a287bb0e15b9bf67a8e4c48b7e430ff28b6f9]
---
 projects/hip/bin/hipcc | 22 ++++++++++++++++++++++
 1 file changed, 22 insertions(+)

diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc
index 1758cb5dce..d1f757d0c4 100755
--- a/projects/hip/bin/hipcc
+++ b/projects/hip/bin/hipcc
@@ -52,6 +52,7 @@ 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;
@@ -240,6 +241,24 @@ if ($needHipHcc) {
     }
 }
 
+# Read .hipconfig
+my %hipConfig = ();
+$hipConfig{'VALID'}=0;
+if (open (CONFIG, "$HIP_PATH/bin/.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);
+}
+
 # hipcc currrently requires separate compilation of source files, ie it is not possible to pass
 # CPP files combined with .O files
 # Reason is that NVCC uses the file extension to determine whether to compile in CUDA mode or
@@ -267,5 +286,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 ();
 }

From 41da8068f1e909ad0ab93d49f7115520decf90d3 Mon Sep 17 00:00:00 2001
From: scchan <siuchi.chan@amd.com>
Date: Tue, 2 Aug 2016 15:58:10 -0400
Subject: [PATCH 21/22] compile all hip objects with -fPIC so that they work in
 a shared library

Change-Id: I1f8355d4a81dbd5f408062c317c16a5efc668876


[ROCm/hip commit: 2f6ac8c36ef0b3657fcb252e2c5ee89b6ac2569e]
---
 projects/hip/CMakeLists.txt | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 88b14b0521..08e59ba32b 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

From ea3ad49a8c47570200f75eaa9b3eb1310c0a9115 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Wed, 3 Aug 2016 12:01:48 +0530
Subject: [PATCH 22/22] .hipconfig now lives in lib

Also moved reading .hipconfig to begining of hipcc

Change-Id: Ibb9d32bef64a79f189aa037c36814759c8cc8052


[ROCm/hip commit: bc394505cc624860cfb5ec3581c36523264aa127]
---
 projects/hip/CMakeLists.txt |  2 +-
 projects/hip/bin/hipcc      | 37 +++++++++++++++++++------------------
 2 files changed, 20 insertions(+), 19 deletions(-)

diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt
index 08e59ba32b..a26e848985 100644
--- a/projects/hip/CMakeLists.txt
+++ b/projects/hip/CMakeLists.txt
@@ -172,7 +172,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
     endif()
 
     # Install .hipconfig
-    install(FILES ${PROJECT_BINARY_DIR}/.hipconfig DESTINATION bin)
+    install(FILES ${PROJECT_BINARY_DIR}/.hipconfig DESTINATION lib)
 endif()
 
 # Install src, bin, include if necessary
diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc
index d1f757d0c4..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`;
@@ -241,24 +260,6 @@ if ($needHipHcc) {
     }
 }
 
-# Read .hipconfig
-my %hipConfig = ();
-$hipConfig{'VALID'}=0;
-if (open (CONFIG, "$HIP_PATH/bin/.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);
-}
-
 # hipcc currrently requires separate compilation of source files, ie it is not possible to pass
 # CPP files combined with .O files
 # Reason is that NVCC uses the file extension to determine whether to compile in CUDA mode or