yaxunl created this revision. yaxunl added a reviewer: tra. Herald added a subscriber: mgorny. yaxunl requested review of this revision.
This patch adds wrapper headers for `<functional>` and a few others which is required to support `<functional>`. The basic idea is to make template functions defined in these headers host device by pragmas. Since this only works for libc++. The code is conditioned for libc++ only. For libstdc++ it is NFC. https://reviews.llvm.org/D102507 Files: clang/docs/HIPSupport.rst clang/lib/Headers/CMakeLists.txt clang/lib/Headers/__clang_hip_runtime_wrapper.h clang/lib/Headers/cuda_wrappers/__tuple clang/lib/Headers/cuda_wrappers/array clang/lib/Headers/cuda_wrappers/functional clang/lib/Headers/cuda_wrappers/tuple clang/lib/Headers/cuda_wrappers/type_traits clang/lib/Headers/cuda_wrappers/utility
Index: clang/lib/Headers/cuda_wrappers/utility =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/utility @@ -0,0 +1,24 @@ +/*===---- utility - CUDA/HIP wrapper for <utility> -------------------------=== + * + * 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_UTILITY +#define __CLANG_CUDA_WRAPPERS_UTILITY + +#if __HIP_USE_LIBCPP +#pragma push_macro("_LIBCPP_NO_EXCEPTIONS") +#define _LIBCPP_NO_EXCEPTIONS +#pragma clang force_cuda_host_device begin +#include_next <utility> +#pragma clang force_cuda_host_device end +#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS") +#else +#include_next <utility> +#endif // __HIP_USE_LIBCPP + +#endif // __CLANG_CUDA_WRAPPERS_UTILITY Index: clang/lib/Headers/cuda_wrappers/type_traits =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/type_traits @@ -0,0 +1,24 @@ +/*===---- type_traits - CUDA/HIP wrapper for <type_traits> -----------------=== + * + * 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_TYPE_TRAITS +#define __CLANG_CUDA_WRAPPERS_TYPE_TRAITS + +#if __HIP_USE_LIBCPP +#pragma push_macro("_LIBCPP_NO_EXCEPTIONS") +#define _LIBCPP_NO_EXCEPTIONS +#pragma clang force_cuda_host_device begin +#include_next <type_traits> +#pragma clang force_cuda_host_device end +#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS") +#else +#include_next <type_traits> +#endif // __HIP_USE_LIBCPP + +#endif // __CLANG_CUDA_WRAPPERS_TYPE_TRAITS Index: clang/lib/Headers/cuda_wrappers/tuple =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/tuple @@ -0,0 +1,24 @@ +/*===---- tuple - CUDA/HIP wrapper for <tuple> -----------------------------=== + * + * 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_TUPLE +#define __CLANG_CUDA_WRAPPERS_TUPLE + +#if __HIP_USE_LIBCPP +#pragma push_macro("_LIBCPP_NO_EXCEPTIONS") +#define _LIBCPP_NO_EXCEPTIONS +#pragma clang force_cuda_host_device begin +#include_next <tuple> +#pragma clang force_cuda_host_device end +#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS") +#else +#include_next <tuple> +#endif // __HIP_USE_LIBCPP + +#endif // __CLANG_CUDA_WRAPPERS_TUPLE Index: clang/lib/Headers/cuda_wrappers/functional =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/functional @@ -0,0 +1,34 @@ +/*===---- functional - CUDA/HIP wrapper for <functional> -------------------=== + * + * 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_FUNCTIONAL +#define __CLANG_CUDA_WRAPPERS_FUNCTIONAL + +#if __HIP_USE_LIBCPP +#pragma push_macro("_LIBCPP_NO_EXCEPTIONS") +#define _LIBCPP_NO_EXCEPTIONS + +extern __device__ void abort() __attribute__ ((__noreturn__)); +namespace std { +namespace __1 { + inline __device__ void abort() { + return ::abort(); + } +} +} + +#pragma clang force_cuda_host_device begin +#include_next <functional> +#pragma clang force_cuda_host_device end +#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS") +#else +#include_next <functional> +#endif // __HIP_USE_LIBCPP + +#endif // __CLANG_CUDA_WRAPPERS_FUNCTIONAL Index: clang/lib/Headers/cuda_wrappers/array =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/array @@ -0,0 +1,24 @@ +/*===---- array - CUDA/HIP wrapper for <array> -----------------------------=== + * + * 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_ARRAY +#define __CLANG_CUDA_WRAPPERS_ARRAY + +#if __HIP_USE_LIBCPP +#pragma push_macro("_LIBCPP_NO_EXCEPTIONS") +#define _LIBCPP_NO_EXCEPTIONS +#pragma clang force_cuda_host_device begin +#include_next <array> +#pragma clang force_cuda_host_device end +#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS") +#else +#include_next <array> +#endif // __HIP_USE_LIBCPP + +#endif // __CLANG_CUDA_WRAPPERS_ARRAY Index: clang/lib/Headers/cuda_wrappers/__tuple =================================================================== --- /dev/null +++ clang/lib/Headers/cuda_wrappers/__tuple @@ -0,0 +1,24 @@ +/*===---- __tuple - CUDA/HIP wrapper for <__tuple> -------------------------=== + * + * 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___TUPLE +#define __CLANG_CUDA_WRAPPERS___TUPLE + +#if __HIP_USE_LIBCPP +#pragma push_macro("_LIBCPP_NO_EXCEPTIONS") +#define _LIBCPP_NO_EXCEPTIONS +#pragma clang force_cuda_host_device begin +#include_next <__tuple> +#pragma clang force_cuda_host_device end +#pragma pop_macro("_LIBCPP_NO_EXCEPTIONS") +#else +#include_next <__tuple> +#endif // __HIP_USE_LIBCPP + +#endif // __CLANG_CUDA_WRAPPERS___TUPLE 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 @@ -18,6 +18,10 @@ #if __HIP__ +#if __has_include(<__libcpp_version>) +#define __HIP_USE_LIBCPP 1 +#endif // __has_include(<__libcpp_version>) + #if !defined(__HIPCC_RTC__) #include <cmath> #include <cstdlib> Index: clang/lib/Headers/CMakeLists.txt =================================================================== --- clang/lib/Headers/CMakeLists.txt +++ clang/lib/Headers/CMakeLists.txt @@ -140,8 +140,14 @@ set(cuda_wrapper_files cuda_wrappers/algorithm + cuda_wrappers/array cuda_wrappers/complex + cuda_wrappers/functional cuda_wrappers/new + cuda_wrappers/type_traits + cuda_wrappers/tuple + cuda_wrappers/utility + cuda_wrappers/__tuple ) set(ppc_wrapper_files Index: clang/docs/HIPSupport.rst =================================================================== --- /dev/null +++ clang/docs/HIPSupport.rst @@ -0,0 +1,23 @@ +============ +HIP Support +============ + +.. contents:: + :local: + +Introduction +============ + +This document describes HIP support in clang. More details are provided in +`external document <https://github.com/ROCm-Developer-Tools/HIP>`_\ , +which are going to be added to clang documentation in the future. + +Standard Library Support +======================== + +<std::functional> +----------------- + +Clang supports calling std::functioinal functors in HIP device code. However +this is limited to `-stdlib=libc++`. +
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
