Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_robust_constant

include/hip/hcc_detail/device_functions.h

@@ -33,6 +33,15 @@ THE SOFTWARE.
 #include <hip/hip_vector_types.h>
 #include <hip/hcc_detail/device_library_decls.h>
 #include <hip/hcc_detail/llvm_intrinsics.h>
+
+#if __HIP_CLANG_ONLY__
+#if __HIP_VDI__
+extern "C" __device__ int printf(const char *fmt, ...);
+#else
+static inline __device__ void printf(const char* format, All... all) {}
+#endif
+#endif
+
 /*
 Integer Intrinsics
 */
@@ -548,7 +557,7 @@ long __shfl_xor(long var, int lane_mask, int width = warpSize)
     return tmp1;
     #else
     static_assert(sizeof(long) == sizeof(int), "");
-    return static_cast<long>(__shfl_down(static_cast<int>(var), lane_delta, width));
+    return static_cast<long>(__shfl_xor(static_cast<int>(var), lane_mask, width));
     #endif
 }
 __device__

include/hip/hcc_detail/device_library_decls.h

@@ -72,6 +72,7 @@ extern "C" __device__ __attribute__((const)) uint __ockl_multi_grid_thread_rank(
 extern "C" __device__ __attribute__((const)) int __ockl_multi_grid_is_valid(void);
 extern "C" __device__ __attribute__((convergent)) void __ockl_multi_grid_sync(void);
 
+extern "C" __device__ void __ockl_global_atomic_add_f32(__attribute__((address_space(1))) float*, float);
 
 // Introduce local address space
 #define __local __attribute__((address_space(3)))

include/hip/hcc_detail/hip_atomic.h

@@ -73,6 +73,14 @@ float atomicAdd(float* address, float val)
 
     return __uint_as_float(r);
 }
+
+__device__
+inline
+void atomicAddNoRet(float* address, float val)
+{
+    __ockl_global_atomic_add_f32((__attribute__((address_space(1))) float*)address, val);
+}
+
 __device__
 inline
 double atomicAdd(double* address, double val)

include/hip/hcc_detail/hip_runtime.h

@@ -311,15 +311,19 @@ extern "C" __device__ void* __hip_free(void* ptr);
 static inline __device__ void* malloc(size_t size) { return __hip_malloc(size); }
 static inline __device__ void* free(void* ptr) { return __hip_free(ptr); }
 
-#if defined(__HCC_ACCELERATOR__) && defined(HC_FEATURE_PRINTF)
+// Declare printf only for the HCC compiler. hip-clang is handled in
+// device_functions.h
+#if __HCC_ACCELERATOR__
+#if HC_FEATURE_PRINTF
 template <typename... All>
 static inline __device__ void printf(const char* format, All... all) {
     hc::printf(format, all...);
 }
-#elif defined(__HCC_ACCELERATOR__) || __HIP__
+#else
 template <typename... All>
 static inline __device__ void printf(const char* format, All... all) {}
-#endif
+#endif // HC_FEATURE_PRINTF
+#endif // __HCC_ACCELERATOR__
 
 #endif //__HCC_OR_HIP_CLANG__