From 6acd8d4cf2a19d9e2fd7c4c91dd5305df2eab2bd 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/clr commit: 13a191774147c6ae6e28023108504cce81ce353a]
---
 .../clr/hipamd/include/hip/hcc_detail/hip_atomic.h     | 10 +++++++---
 1 file changed, 7 insertions(+), 3 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h
index 4af1794ba0..1304cf74a9 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_atomic.h
+++ b/projects/clr/hipamd/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.