diff --git a/docs/markdown/hip_kernel_language.md b/docs/markdown/hip_kernel_language.md index 3b13cde08a..d629063f73 100644 --- a/docs/markdown/hip_kernel_language.md +++ b/docs/markdown/hip_kernel_language.md @@ -233,8 +233,7 @@ typedef struct dim3 { ## Memory-Fence Instructions 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. +Support for threadfence_system() is under development. ## 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. diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 65bf42516f..73fc022362 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -556,7 +556,7 @@ extern "C" __device__ void __threadfence(void); * * @param void * - * @warning __threadfence_system is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches. + * @warning __threadfence_system is a stub and map to no-op. */ __device__ void __threadfence_system(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));