From 49565f9d9f4c738c94f0a8fe55db1887b6637a8f Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Mon, 15 Dec 2025 17:23:33 +0100 Subject: [PATCH] SWDEV-548892 - Always declare used ocml and ockl device libs functions (#2230) Ignore __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__. This should not be relying on declarations from the clang builtin headers. There is no issue declaring the same intrinsics multiple times. This will enable removal of declarations from the clang builtin headers. --- projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h | 2 -- projects/clr/hipamd/include/hip/amd_detail/math_fwd.h | 3 --- 2 files changed, 5 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h b/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h index 6798ef5dbd..cfdfe26803 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_fp16_math_fwd.h @@ -28,7 +28,6 @@ THE SOFTWARE. #if !defined(__HIPCC_RTC__) #include "host_defines.h" #endif -#ifndef __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ extern "C" { __device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16); __device__ _Float16 __ocml_cos_f16(_Float16); @@ -79,7 +78,6 @@ __device__ __attribute__((const)) _Float16 __ocml_cvtrtn_f16_f32(float); __device__ __attribute__((const)) _Float16 __ocml_cvtrtp_f16_f32(float); __device__ __attribute__((const)) _Float16 __ocml_cvtrtz_f16_f32(float); } -#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // TODO: remove these after they get into clang header __clang_hip_libdevice_declares.h' extern "C" { __device__ __attribute__((const)) _Float16 __ocml_fmax_f16(_Float16, _Float16); diff --git a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h index a20594d617..533de15994 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h +++ b/projects/clr/hipamd/include/hip/amd_detail/math_fwd.h @@ -55,7 +55,6 @@ __device__ __attribute__((const)) unsigned int __ockl_udot8(unsigned int, unsign unsigned int, bool); #endif -#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ // BEGIN FLOAT __device__ __attribute__((const)) float __ocml_acos_f32(float); __device__ __attribute__((pure)) float __ocml_acosh_f32(float); @@ -282,8 +281,6 @@ __device__ __attribute__((const)) double __ocml_fma_rtz_f64(double, double, doub // END INTRINSICS // END DOUBLE -#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ - #if defined(__cplusplus) } // extern "C" #endif