Skip to content

Commit b1c2267

Browse files
AminsedAminsed
committed
libcudacxx: diagnose extended lambdas in invocability traits
1 parent f328360 commit b1c2267

File tree

4 files changed

+172
-17
lines changed

4 files changed

+172
-17
lines changed

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

Lines changed: 79 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -33,13 +33,39 @@
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+
47+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
48+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
49+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
50+
template <class _Fn>
51+
_CCCL_INLINE_VISIBILITY constexpr bool __disallow_extended_lambda_invocability_v =
52+
# if defined(__nv_is_extended_device_lambda_closure_type)
53+
__nv_is_extended_device_lambda_closure_type(remove_cvref_t<_Fn>) ||
54+
# else
55+
false ||
56+
# endif
57+
# if defined(__nv_is_extended_host_device_lambda_closure_type)
58+
__nv_is_extended_host_device_lambda_closure_type(remove_cvref_t<_Fn>);
59+
# else
60+
false;
61+
# endif
62+
#else
63+
template <class>
64+
_CCCL_INLINE_VISIBILITY constexpr bool __disallow_extended_lambda_invocability_v = false;
65+
#endif
66+
67+
} // namespace __detail
68+
4369
struct __any
4470
{
4571
_CCCL_API inline __any(...);
@@ -214,17 +240,35 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT invoke_result //
214240
215241
template <class _Fn, class... _Args>
216242
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_invocable : bool_constant<__is_invocable<_Fn, _Args...>>
217-
{};
243+
{
244+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
245+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
246+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
247+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
248+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
249+
"that requires querying its invocability in host code. Use a named function object or "
250+
"cuda::proclaim_return_type instead.");
251+
#endif
252+
};
218253
219254
template <class _Ret, class _Fn, class... _Args>
220255
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_invocable_r : bool_constant<__is_invocable_r<_Ret, _Fn, _Args...>>
221-
{};
256+
{
257+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
258+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
259+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
260+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
261+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
262+
"that requires querying its invocability in host code. Use a named function object or "
263+
"cuda::proclaim_return_type instead.");
264+
#endif
265+
};
222266
223267
template <class _Fn, class... _Args>
224-
inline constexpr bool is_invocable_v = __is_invocable<_Fn, _Args...>;
268+
inline constexpr bool is_invocable_v = is_invocable<_Fn, _Args...>::value;
225269
226270
template <class _Ret, class _Fn, class... _Args>
227-
inline constexpr bool is_invocable_r_v = __is_invocable_r<_Ret, _Fn, _Args...>;
271+
inline constexpr bool is_invocable_r_v = is_invocable_r<_Ret, _Fn, _Args...>::value;
228272
229273
// is_nothrow_invocable
230274
@@ -243,22 +287,40 @@ template <class _Ret, class _Fp, class... _Args>
243287
inline constexpr bool __nothrow_invocable_r_imp<true, true, _Ret, _Fp, _Args...> =
244288
noexcept(::cuda::std::__invoke(::cuda::std::declval<_Fp>(), ::cuda::std::declval<_Args>()...));
245289
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-
254290
template <class _Fn, class... _Args>
255-
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_nothrow_invocable : bool_constant<is_nothrow_invocable_v<_Fn, _Args...>>
256-
{};
291+
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_nothrow_invocable
292+
: bool_constant<__nothrow_invocable_r_imp<__is_invocable<_Fn, _Args...>, true, void, _Fn, _Args...>>
293+
{
294+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
295+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
296+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
297+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
298+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
299+
"that requires querying its invocability in host code. Use a named function object or "
300+
"cuda::proclaim_return_type instead.");
301+
#endif
302+
};
257303
258304
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-
{};
305+
struct _CCCL_TYPE_VISIBILITY_DEFAULT is_nothrow_invocable_r
306+
: bool_constant<
307+
__nothrow_invocable_r_imp<__is_invocable_r<_Ret, _Fn, _Args...>, is_void_v<_Ret>, _Ret, _Fn, _Args...>>
308+
{
309+
#if _CCCL_CUDA_COMPILER(NVCC) && defined(__CUDACC_EXTENDED_LAMBDA__) && !_CCCL_DEVICE_COMPILATION() \
310+
&& (defined(__nv_is_extended_device_lambda_closure_type) \
311+
|| defined(__nv_is_extended_host_device_lambda_closure_type))
312+
static_assert(!__detail::__disallow_extended_lambda_invocability_v<_Fn>,
313+
"Attempt to use an extended __device__ or __host__ __device__ lambda in a context "
314+
"that requires querying its invocability in host code. Use a named function object or "
315+
"cuda::proclaim_return_type instead.");
316+
#endif
317+
};
318+
319+
template <class _Fp, class... _Args>
320+
inline constexpr bool is_nothrow_invocable_v = is_nothrow_invocable<_Fp, _Args...>::value;
321+
322+
template <class _Ret, class _Fp, class... _Args>
323+
inline constexpr bool is_nothrow_invocable_r_v = is_nothrow_invocable_r<_Ret, _Fp, _Args...>::value;
262324
263325
// Not going directly through __invoke_result_t because we want the additional device lambda checks in invoke_result
264326
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)