From 0e0b028846c6eb47c4ee467fdf497253411481b2 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Wed, 30 May 2018 16:33:18 -0400 Subject: [PATCH] Fix __syncthreads for hip-clang --- hipamd/include/hip/hcc_detail/hip_runtime.h | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h index 9737ecb19f..1a6b0f7dda 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -471,12 +471,12 @@ static inline __device__ void printf(const char* format, All... all) {} #endif #endif - - -#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) #endif //__HCC_OR_HIP_CLANG__ #ifdef __HCC__ + +#define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) + #define HIP_KERNEL_NAME(...) (__VA_ARGS__) #define HIP_SYMBOL(X) #X @@ -656,6 +656,8 @@ __DEVICE__ void inline __assert_fail(const char * __assertion, __device_trap(); } +extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); + #pragma push_macro("__DEVICE__") #include