From 24c621db5bbd056f495ca7f809b2ac940015960f Mon Sep 17 00:00:00 2001 From: pensun Date: Wed, 9 Nov 2016 20:54:19 -0600 Subject: [PATCH] Add documentation on threadfence_system workaround guidelines. Change-Id: I9636a3808798f3dabe992285ce5652187cee6eb8 --- docs/markdown/hip_kernel_language.md | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index d629063f73..f84868987c 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/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.