From 192839c71bb2e907000580e2f249422b77ec8f57 Mon Sep 17 00:00:00 2001
From: Alex Voicu <alexandru.voicu@amd.com>
Date: Sat, 2 Jun 2018 17:46:37 +0100
Subject: [PATCH] Add missing __device__ for forward declares.

[ROCm/hip commit: 59adb5e52a015a09332f698e5eb82b98789fdb0e]
---
 projects/hip/include/hip/hcc_detail/hip_atomic.h | 10 +++++++---
 1 file changed, 7 insertions(+), 3 deletions(-)

diff --git a/projects/hip/include/hip/hcc_detail/hip_atomic.h b/projects/hip/include/hip/hcc_detail/hip_atomic.h
index 4af1794ba0..1304cf74a9 100644
--- a/projects/hip/include/hip/hcc_detail/hip_atomic.h
+++ b/projects/hip/include/hip/hcc_detail/hip_atomic.h
@@ -176,7 +176,9 @@ __device__
 inline
 unsigned int atomicInc(unsigned int* address, unsigned int val)
 {
-    extern unsigned int __builtin_amdgcn_atomic_inc(
+    __device__
+    extern 
+    unsigned int __builtin_amdgcn_atomic_inc(
         unsigned int*,
         unsigned int,
         unsigned int,
@@ -191,7 +193,9 @@ __device__
 inline
 unsigned int atomicDec(unsigned int* address, unsigned int val)
 {
-    extern unsigned int __builtin_amdgcn_atomic_dec(
+    __device__
+    extern 
+    unsigned int __builtin_amdgcn_atomic_dec(
         unsigned int*,
         unsigned int,
         unsigned int,
@@ -262,4 +266,4 @@ unsigned long long atomicXor(
     return __atomic_fetch_xor(address, val, __ATOMIC_RELAXED);
 }
 
-// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block.
\ No newline at end of file
+// TODO: add scoped atomics i.e. atomic{*}_system && atomic{*}_block.