AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

added threadfence feature for hcc

Browse Source 
1. Added feature for __threadfence and __threadfence_block
2. Added feature for using LLVM IR files directly while compilation
3. Added test for threadfence and threadfence_block

Change-Id: Ib7e5d89b4cca1a135952b317e5809cd05b56a3c9
This commit is contained in:
Aditya Atluri
2016-10-10 15:29:50 -05:00
parent 7c943ef20a
commit dc1042ce6d
5 changed files with 94 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/hip/hcc_detail/hip_runtime.h
+4 -2
View File
@@ -533,7 +533,8 @@ __device__ float __dsqrt_rz(double x);
*
* @warning __threadfence_block is a stub and map to no-op.
*/
__device__ void __threadfence_block(void);
// __device__ void __threadfence_block(void);
extern "C" __device__ void __threadfence_block(void);
/**
* @brief threadfence makes wirtes visible to other threads running on same GPU.
@@ -544,7 +545,8 @@ __device__ void __threadfence_block(void);
*
* @warning __threadfence is a stub and map to no-op, application should set "export HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches.
*/
__device__ void __threadfence(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
// __device__ void __threadfence(void) __attribute__((deprecated("Provided for compile-time compatibility, not yet functional")));
extern "C" __device__ void __threadfence(void);
/**
* @brief threadfence_system makes writes to pinned system memory visible on host CPU.