diff --git a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
index 81f9fe4761..fccbcfbfdc 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16_math_fwd.h
@@ -28,49 +28,49 @@ THE SOFTWARE.
 
 extern "C"
 {
-    __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
-    _Float16 __ocml_cos_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
-    __attribute__((const))
+    __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
+    __device__ _Float16 __ocml_cos_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
+    __device__ __attribute__((const))
     _Float16 __ocml_fma_f16(_Float16, _Float16, _Float16);
-    __attribute__((const)) int __ocml_isinf_f16(_Float16);
-    __attribute__((const)) int __ocml_isnan_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
-    __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
-    __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
-    _Float16 __ocml_sin_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
-    __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
+    __device__ __attribute__((const)) int __ocml_isinf_f16(_Float16);
+    __device__ __attribute__((const)) int __ocml_isnan_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
+    __device__ __attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
+    __device__ _Float16 __ocml_sin_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
+    __device__ __attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
 
     typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
     typedef short __2i16 __attribute__((ext_vector_type(2)));
 
-    __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
-    __2f16 __ocml_cos_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
-    __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
-    __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
-    __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
-    inline
+    __device__ __attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
+    __device__ __2f16 __ocml_cos_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
+    __device__ __attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
+    __device__ __attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
+    __device__ __attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
+    __device__ inline
     __2f16 __llvm_amdgcn_rcp_2f16(__2f16 x) // Not currently exposed by ROCDL.
     {
         return __2f16{__llvm_amdgcn_rcp_f16(x.x), __llvm_amdgcn_rcp_f16(x.y)};
     }
-    __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
-    __2f16 __ocml_sin_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
-    __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
+    __device__ __2f16 __ocml_sin_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
+    __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
 }
\ No newline at end of file