Add documentation on threadfence_system workaround guidelines.

Change-Id: I9636a3808798f3dabe992285ce5652187cee6eb8

docs/markdown/hip_kernel_language.md

@@ -233,7 +233,11 @@ typedef struct dim3 {
 ## Memory-Fence Instructions
 HIP supports __threadfence() and  __threadfence_block().
 
-Support for threadfence_system() is under development.
+HIP provides workaround for threadfence_system() under HCC path.
+To enable the workaround, HIP should be built with environment variable HIP_COHERENT_HOST_ALLOC enabled.
+In addition,the kernels that use __threadfence_system() should be modified as follows:
+- The kernel should only operate on finegrained system memory; which should be allocated with hipHostMalloc().
+- Remove all memcpy for those allocated finegrained system memory regions.
 
 ## 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.