hliao created this revision.
hliao added reviewers: tra, yaxunl.
hliao requested review of this revision.
Herald added a project: clang.
Herald added a subscriber: cfe-commits.
- MSVC has different `<complex>` implementation which calls into functions
declared in `<ymath.h>`. Instead of through builtins or inline functions,
`<ymath.h>` functions are provided through external libraries. To get
`<complex>` compiled with HIP on MSVC, we need to + Wrap `<ymath.h>` to force
its functions to be `__host__ __device__`. + Provide its function definitions
for the device compilation.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D93638
Files:
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,36 @@
+/*===---- complex - CUDA wrapper for <ymath.h> -----------------------------===
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#ifndef __CLANG_CUDA_WRAPPERS_COMPLEX
+#define __CLANG_CUDA_WRAPPERS_COMPLEX
+
+// 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,18 @@
#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__)
+__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; }
+#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/cuda_wrappers/ymath.h
===================================================================
--- /dev/null
+++ clang/lib/Headers/cuda_wrappers/ymath.h
@@ -0,0 +1,36 @@
+/*===---- complex - CUDA wrapper for <ymath.h> -----------------------------===
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to
+ * deal in the Software without restriction, including without limitation the
+ * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
+ * sell copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
+ * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
+ * IN THE SOFTWARE.
+ *
+ *===-----------------------------------------------------------------------===
+ */
+
+#ifndef __CLANG_CUDA_WRAPPERS_COMPLEX
+#define __CLANG_CUDA_WRAPPERS_COMPLEX
+
+// 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,18 @@
#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__)
+__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; }
+#endif // defined(_MSC_VER) && defined(__HIP_DEVICE_COMPILE__)
+
#endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
#define __CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__ 1
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
