Skip to content

Commit 150d8f6

Browse files
AminsedAminsed
committed
libcudacxx: diagnose extended lambdas in invocability traits
1 parent 48ff637 commit 150d8f6

File tree

4 files changed

+170
-17
lines changed

4 files changed

+170
-17
lines changed

‎libcudacxx/include/cuda/std/__functional/invoke.h‎

Lines changed: 77 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -33,13 +33,37 @@
3333
#include <cuda/std/__type_traits/is_same.h>
3434
#include <cuda/std/__type_traits/is_void.h>
3535
#include <cuda/std/__type_traits/nat.h>
36+
#include <cuda/std/__type_traits/remove_cvref.h>
3637
#include <cuda/std/__utility/declval.h>
3738
#include <cuda/std/__utility/forward.h>
3839

3940
#include <cuda/std/__cccl/prologue.h>
4041

4142
_CCCL_BEGIN_NAMESPACE_CUDA_STD
4243

44+
namespace __detail
45+
{
46+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
47+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
48+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
49+
template <class _Fn>
50+
_CCCL_INLINE_VISIBILITY constexpr bool __disallow_extended_lambda_invocability_v =
51+
# if defined(__nv_is_extended_device_lambda_closure_type)
52+
__nv_is_extended_device_lambda_closure_type(remove_cvref_t<_Fn>) ||
53+
# else
54+
false ||
55+
# endif
56+
# if defined(__nv_is_extended_host_device_lambda_closure_type)
57+
__nv_is_extended_host_device_lambda_closure_type(remove_cvref_t<_Fn>);
58+
# else
59+
false;
60+
# endif
61+
#else
62+
template <class>
63+
_CCCL_INLINE_VISIBILITY constexpr bool __disallow_extended_lambda_invocability_v = false;
64+
#endif
65+
} // namespace __detail
66+
4367
struct __any
4468
{
4569
_CCCL_API inline __any(...);
@@ -214,17 +238,35 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT invoke_result //
214238
215239
template <class _Fn, class... _Args>
216240
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_invocable : bool_constant<__is_invocable<_Fn, _Args...>>
217-
{};
241+
{
242+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
243+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
244+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
245+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
246+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
247+
"that requires querying its invocability in host code. Use a named function object or "
248+
"cuda::proclaim_return_type instead.");
249+
#endif
250+
};
218251
219252
template <class _Ret, class _Fn, class... _Args>
220253
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_invocable_r : bool_constant<__is_invocable_r<_Ret, _Fn, _Args...>>
221-
{};
254+
{
255+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
256+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
257+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
258+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
259+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
260+
"that requires querying its invocability in host code. Use a named function object or "
261+
"cuda::proclaim_return_type instead.");
262+
#endif
263+
};
222264
223265
template <class _Fn, class... _Args>
224-
inline constexpr bool is_invocable_v = __is_invocable<_Fn, _Args...>;
266+
inline constexpr bool is_invocable_v = is_invocable<_Fn, _Args...>::value;
225267
226268
template <class _Ret, class _Fn, class... _Args>
227-
inline constexpr bool is_invocable_r_v = __is_invocable_r<_Ret, _Fn, _Args...>;
269+
inline constexpr bool is_invocable_r_v = is_invocable_r<_Ret, _Fn, _Args...>::value;
228270
229271
// is_nothrow_invocable
230272
@@ -243,22 +285,40 @@ template <class _Ret, class _Fp, class... _Args>
243285
inline constexpr bool __nothrow_invocable_r_imp<true, true, _Ret, _Fp, _Args...> =
244286
noexcept(::cuda::std::__invoke(::cuda::std::declval<_Fp>(), ::cuda::std::declval<_Args>()...));
245287
246-
template <class _Fp, class... _Args>
247-
inline constexpr bool is_nothrow_invocable_v =
248-
__nothrow_invocable_r_imp<__is_invocable<_Fp, _Args...>, true, void, _Fp, _Args...>;
249-
250-
template <class _Ret, class _Fp, class... _Args>
251-
inline constexpr bool is_nothrow_invocable_r_v =
252-
__nothrow_invocable_r_imp<__is_invocable_r<_Ret, _Fp, _Args...>, is_void_v<_Ret>, _Ret, _Fp, _Args...>;
253-
254288
template <class _Fn, class... _Args>
255-
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_nothrow_invocable : bool_constant<is_nothrow_invocable_v<_Fn, _Args...>>
256-
{};
289+
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_nothrow_invocable
290+
: bool_constant<__nothrow_invocable_r_imp<__is_invocable<_Fn, _Args...>, true, void, _Fn, _Args...>>
291+
{
292+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
293+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
294+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
295+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
296+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
297+
"that requires querying its invocability in host code. Use a named function object or "
298+
"cuda::proclaim_return_type instead.");
299+
#endif
300+
};
257301
258302
template <class _Ret, class _Fn, class... _Args>
259-
struct _CCCL_TYPE_VISIBILITY_DEFAULT
260-
is_nothrow_invocable_r : bool_constant<is_nothrow_invocable_r_v<_Ret, _Fn, _Args...>>
261-
{};
303+
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_nothrow_invocable_r
304+
: bool_constant<
305+
__nothrow_invocable_r_imp<__is_invocable_r<_Ret, _Fn, _Args...>, is_void_v<_Ret>, _Ret, _Fn, _Args...>>
306+
{
307+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
308+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
309+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
310+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
311+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
312+
"that requires querying its invocability in host code. Use a named function object or "
313+
"cuda::proclaim_return_type instead.");
314+
#endif
315+
};
316+
317+
template <class _Fp, class... _Args>
318+
inline constexpr bool is_nothrow_invocable_v = is_nothrow_invocable<_Fp, _Args...>::value;
319+
320+
template <class _Ret, class _Fp, class... _Args>
321+
inline constexpr bool is_nothrow_invocable_r_v = is_nothrow_invocable_r<_Ret, _Fp, _Args...>::value;
262322
263323
// Not going directly through __invoke_result_t because we want the additional device lambda checks in invoke_result
264324
template <class _Fn, class... _Args>

‎libcudacxx/test/libcudacxx/std/utilities/meta/meta.rel/is_invocable.extended_device_lambda.fail.cpp‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
// type_traits
10+
11+
// is_invocable should diagnose extended lambdas in host code.
12+
13+
// UNSUPPORTED: clang && (!nvcc)
14+
15+
#include <cuda/std/type_traits>
16+
17+
#include "test_macros.h"
18+
19+
template <class Fn>
20+
void instantiate()
21+
{
22+
(void) cuda::std::is_invocable_v<Fn>;
23+
}
24+
25+
int main(int, char**)
26+
{
27+
#if TEST_CUDA_COMPILER(NVCC) || TEST_COMPILER(NVRTC)
28+
instantiate<decltype([] __device__() {})>();
29+
#endif
30+
return 0;
31+
}

‎libcudacxx/test/libcudacxx/std/utilities/meta/meta.rel/is_invocable.extended_host_device_lambda.fail.cpp‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
// type_traits
10+
11+
// is_invocable should diagnose extended lambdas in host code.
12+
13+
// UNSUPPORTED: clang && (!nvcc)
14+
15+
#include <cuda/std/type_traits>
16+
17+
#include "test_macros.h"
18+
19+
template <class Fn>
20+
void instantiate()
21+
{
22+
(void) cuda::std::is_invocable_v<Fn>;
23+
}
24+
25+
int main(int, char**)
26+
{
27+
#if TEST_CUDA_COMPILER(NVCC) || TEST_COMPILER(NVRTC)
28+
instantiate<decltype([] __host__ __device__() {})>();
29+
#endif
30+
return 0;
31+
}

‎libcudacxx/test/libcudacxx/std/utilities/meta/meta.rel/is_nothrow_invocable.extended_device_lambda.fail.cpp‎

Lines changed: 31 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
//===----------------------------------------------------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
// type_traits
10+
11+
// is_nothrow_invocable should diagnose extended lambdas in host code.
12+
13+
// UNSUPPORTED: clang && (!nvcc)
14+
15+
#include <cuda/std/type_traits>
16+
17+
#include "test_macros.h"
18+
19+
template <class Fn>
20+
void instantiate()
21+
{
22+
(void) cuda::std::is_nothrow_invocable_v<Fn>;
23+
}
24+
25+
int main(int, char**)
26+
{
27+
#if TEST_CUDA_COMPILER(NVCC) || TEST_COMPILER(NVRTC)
28+
instantiate<decltype([] __device__() {})>();
29+
#endif
30+
return 0;
31+
}

0 commit comments

Comments
 (0)