hliao updated this revision to Diff 313211. hliao added a comment. Herald added a subscriber: mgorny.
Fix the cmake to distribute that header wrapper. Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D93638/new/ https://reviews.llvm.org/D93638 Files: clang/lib/Headers/CMakeLists.txt clang/lib/Headers/__clang_hip_runtime_wrapper.h clang/lib/Headers/cuda_wrappers/ymath.h Index: clang/lib/Headers/cuda_wrappers/ymath.h =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/ymath.h @@ -0,0 +1,22 @@ +/*===---- ymath.h - HIP wrapper for <ymath.h> ------------------------------=== + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. + * See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + *===-----------------------------------------------------------------------=== + */ + +#ifndef __CLANG_CUDA_WRAPPERS_YMATH_H +#define __CLANG_CUDA_WRAPPERS_YMATH_H + +// Wrapper around <ymath.h> that forces its functions to be __host__ +// __device__. + +#pragma clang force_cuda_host_device begin + +#include_next <ymath.h> + +#pragma clang force_cuda_host_device end + +#endif // include guard Index: clang/lib/Headers/__clang_hip_runtime_wrapper.h =================================================================== --- clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -61,6 +61,28 @@ #include <algorithm> #include <complex> #include <new> + +// Define math functions from <ymath.h> on MSVC for the device compilation +// only to avoid conflicts with MSVC runtime in the host compilation. +#if defined(_MSC_VER) && defined(__HIP_DEVICE_COMPILE__) +#if defined(__cplusplus) +extern "C" { +#endif // defined(__cplusplus) +#pragma push_macro("__HOST_DEVICE__") +#define __HOST_DEVICE__ \ + static __host__ __device__ inline __attribute__((always_inline)) +__HOST_DEVICE__ double _Cosh(double x, double y) { return cosh(x) * y; } +__HOST_DEVICE__ float _FCosh(float x, float y) { return coshf(x) * y; } +__HOST_DEVICE__ short _Dtest(double *p) { return fpclassify(*p); } +__HOST_DEVICE__ short _FDtest(float *p) { return fpclassify(*p); } +__HOST_DEVICE__ double _Sinh(double x, double y) { return sinh(x) * y; } +__HOST_DEVICE__ float _FSinh(float x, float y) { return sinhf(x) * y; } +#pragma pop_macro("__HOST_DEVICE__") +#if defined(__cplusplus) +} +#endif // defined(__cplusplus) +#endif // defined(_MSC_VER) && defined(__HIP_DEVICE_COMPILE__) + #endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1 Index: clang/lib/Headers/CMakeLists.txt =================================================================== --- clang/lib/Headers/CMakeLists.txt +++ clang/lib/Headers/CMakeLists.txt @@ -142,6 +142,7 @@ cuda_wrappers/algorithm cuda_wrappers/complex cuda_wrappers/new + cuda_wrappers/ymath.h ) set(ppc_wrapper_files
Index: clang/lib/Headers/cuda_wrappers/ymath.h =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/ymath.h @@ -0,0 +1,22 @@ +/*===---- ymath.h - HIP wrapper for <ymath.h> ------------------------------=== + * + * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. + * See https://llvm.org/LICENSE.txt for license information. + * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + * + *===-----------------------------------------------------------------------=== + */ + +#ifndef __CLANG_CUDA_WRAPPERS_YMATH_H +#define __CLANG_CUDA_WRAPPERS_YMATH_H + +// Wrapper around <ymath.h> that forces its functions to be __host__ +// __device__. + +#pragma clang force_cuda_host_device begin + +#include_next <ymath.h> + +#pragma clang force_cuda_host_device end + +#endif // include guard Index: clang/lib/Headers/__clang_hip_runtime_wrapper.h =================================================================== --- clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -61,6 +61,28 @@ #include <algorithm> #include <complex> #include <new> + +// Define math functions from <ymath.h> on MSVC for the device compilation +// only to avoid conflicts with MSVC runtime in the host compilation. +#if defined(_MSC_VER) && defined(__HIP_DEVICE_COMPILE__) +#if defined(__cplusplus) +extern "C" { +#endif // defined(__cplusplus) +#pragma push_macro("__HOST_DEVICE__") +#define __HOST_DEVICE__ \ + static __host__ __device__ inline __attribute__((always_inline)) +__HOST_DEVICE__ double _Cosh(double x, double y) { return cosh(x) * y; } +__HOST_DEVICE__ float _FCosh(float x, float y) { return coshf(x) * y; } +__HOST_DEVICE__ short _Dtest(double *p) { return fpclassify(*p); } +__HOST_DEVICE__ short _FDtest(float *p) { return fpclassify(*p); } +__HOST_DEVICE__ double _Sinh(double x, double y) { return sinh(x) * y; } +__HOST_DEVICE__ float _FSinh(float x, float y) { return sinhf(x) * y; } +#pragma pop_macro("__HOST_DEVICE__") +#if defined(__cplusplus) +} +#endif // defined(__cplusplus) +#endif // defined(_MSC_VER) && defined(__HIP_DEVICE_COMPILE__) + #endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__ #define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1 Index: clang/lib/Headers/CMakeLists.txt =================================================================== --- clang/lib/Headers/CMakeLists.txt +++ clang/lib/Headers/CMakeLists.txt @@ -142,6 +142,7 @@ cuda_wrappers/algorithm cuda_wrappers/complex cuda_wrappers/new + cuda_wrappers/ymath.h ) set(ppc_wrapper_files
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
