From 2bd034a0d94d99cebfcc7ffb5fc2a82f3f3634ff 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 [ROCm/hip commit: b7641cf8353fb4b510c66bf7f1edae3f946a9a9f] --- projects/hip/include/hip/hcc_detail/hip_runtime.h | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 9737ecb19f..1a6b0f7dda 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/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