Move some internal headers from "include/hip/" to src.

Change-Id: I7041bd5c803d9318979f4a7c1d658445c614691e

hipamd/docs/markdown/hip_kernel_language.md

@@ -231,10 +231,10 @@ typedef struct dim3 {
 ```
 
 ## Memory-Fence Instructions
-HIP support for __threadfence(), __threadfence_block() and __threadfence_system() is under development.
-The stubs for the threadfence routines are defined in hcc_details/hip_runtime.h.
-Applications that use these threadfence features should disable both of the L1 and L2 caches by:
-"export HSA_DISABLE_CACHE=1"
+HIP supports __threadfence() and  __threadfence_block().
+
+Applications that use threadfence_system can disable the L1 and L2 caches on the GPU by:
+"export HSA_DISABLE_CACHE=1".  See the hip_porting_guide.md#threadfence_system for more information.
 
 ## Synchronization Functions
 The __syncthreads() built-in function is supported in HIP. The __syncthreads_count(int), __syncthreads_and(int) and __syncthreads_or(int) functions are under development.  
@@ -602,7 +602,8 @@ The printf function is under development.
 
 ## Device-Side Dynamic Global Memory Allocation
 
-Device-side dynamic global memory allocation is not supported.
+Device-side dynamic global memory allocation is under development.  HIP now includes a preliminary
+implementation of malloc and free that can be called from device functions.
 
 ## `__launch_bounds__`
 GPU multiprocessors have a fixed pool of resources (primarily registers and shared memory) that are shared among the active warps. Using more resources can increase the kernel’s IPC, but it reduces the resources available for other warps and limits the number of warps that can run simultaneously. Thus, GPUs exhibit a complex relationship between resource usage and performance. `__launch_bounds__` allows the application to provide usage hints that influence the resources (primarily registers) employed by the generated code. It’s a function attribute that must be attached to a `__global__` function:

hipamd/src/hip_context.cpp

@@ -26,8 +26,8 @@ THE SOFTWARE.
 #include <stack>
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 // Stack of contexts
 thread_local std::stack<ihipCtx_t *>  tls_ctxStack;

hipamd/src/hip_device.cpp

@@ -21,8 +21,8 @@ THE SOFTWARE.
 */
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 //-------------------------------------------------------------------------------------------------
 //Devices

hipamd/src/hip_error.cpp

@@ -21,8 +21,8 @@ THE SOFTWARE.
 */
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------

hipamd/src/hip_event.cpp

@@ -21,8 +21,8 @@ THE SOFTWARE.
 */
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------

hipamd/src/hip_hcc.cpp

@@ -44,8 +44,8 @@ THE SOFTWARE.
 #include "libhsakmt/hsakmt.h"
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 
 

hipamd/src/hip_hcc.h

@@ -25,7 +25,7 @@ THE SOFTWARE.
 
 #include <hc.hpp>
 #include <hsa/hsa.h>
-#include "hip/hcc_detail/hip_util.h"
+#include "hip_util.h"
 
 
 #if defined(__HCC__) && (__hcc_workweek__ < 16354)

hipamd/src/hip_memory.cpp

@@ -25,8 +25,8 @@ THE SOFTWARE.
 #include "hsa/hsa_ext_amd.h"
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------

hipamd/src/hip_module.cpp

@@ -30,8 +30,8 @@ THE SOFTWARE.
 #include "hsa/amd_hsa_kernel_code.h"
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 //TODO Use Pool APIs from HCC to get memory regions.
 

hipamd/src/hip_peer.cpp

@@ -23,8 +23,8 @@ THE SOFTWARE.
 #include <hc_am.hpp>
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 
 // Peer access functions.

hipamd/src/hip_stream.cpp

@@ -21,8 +21,8 @@ THE SOFTWARE.
 */
 
 #include "hip/hip_runtime.h"
-#include "hip/hcc_detail/hip_hcc.h"
-#include "hip/hcc_detail/trace_helper.h"
+#include "hip_hcc.h"
+#include "trace_helper.h"
 
 
 //-------------------------------------------------------------------------------------------------