diff --git a/hipamd/docs/markdown/hip_kernel_language.md b/hipamd/docs/markdown/hip_kernel_language.md
index 4b8762025e..3b13cde08a 100644
--- a/hipamd/docs/markdown/hip_kernel_language.md
+++ b/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:
diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp
index 4f45f049f6..a72ae9b58d 100644
--- a/hipamd/src/hip_context.cpp
+++ b/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;
diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp
index 6e36b935c8..fda5724d7f 100644
--- a/hipamd/src/hip_device.cpp
+++ b/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
diff --git a/hipamd/src/hip_error.cpp b/hipamd/src/hip_error.cpp
index 72e7ab0084..60c45cc1f7 100644
--- a/hipamd/src/hip_error.cpp
+++ b/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"
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
diff --git a/hipamd/src/hip_event.cpp b/hipamd/src/hip_event.cpp
index 441918d6c4..5494fc32f2 100644
--- a/hipamd/src/hip_event.cpp
+++ b/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"
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp
index 8bf7f1091c..4fcf427a3c 100644
--- a/hipamd/src/hip_hcc.cpp
+++ b/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"
 
 
 
diff --git a/hipamd/include/hip/hcc_detail/hip_hcc.h b/hipamd/src/hip_hcc.h
similarity index 99%
rename from hipamd/include/hip/hcc_detail/hip_hcc.h
rename to hipamd/src/hip_hcc.h
index 9e7499e4ac..d8045d4cb7 100644
--- a/hipamd/include/hip/hcc_detail/hip_hcc.h
+++ b/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)
diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index 8b030799fb..9e0473f02c 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/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"
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp
index c8368abeec..f7ac35c77b 100644
--- a/hipamd/src/hip_module.cpp
+++ b/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.
 
diff --git a/hipamd/src/hip_peer.cpp b/hipamd/src/hip_peer.cpp
index 4f3227de82..e66a0d2971 100644
--- a/hipamd/src/hip_peer.cpp
+++ b/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.
diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp
index 9bb615ebf7..3b1d6af038 100644
--- a/hipamd/src/hip_stream.cpp
+++ b/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"
 
 
 //-------------------------------------------------------------------------------------------------
diff --git a/hipamd/include/hip/hcc_detail/hip_util.h b/hipamd/src/hip_util.h
similarity index 100%
rename from hipamd/include/hip/hcc_detail/hip_util.h
rename to hipamd/src/hip_util.h
diff --git a/hipamd/include/hip/hcc_detail/trace_helper.h b/hipamd/src/trace_helper.h
similarity index 100%
rename from hipamd/include/hip/hcc_detail/trace_helper.h
rename to hipamd/src/trace_helper.h