Comments (2)
bdice
commented on June 12, 2024
I think this is a minimal reproducer: https://godbolt.org/z/oaMhrcoPv
#include <cuda/functional>
int main()
{
auto f = cuda::proclaim_return_type<bool>([] __device__() { return false; });
auto g = cuda::proclaim_return_type<bool>([f] __device__() { return f(); });
return 0;
}/opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/std/detail/libcxx/include/__functional/invoke.h(402): error: calling a __device__ function("main::[lambda() (instance 1)]::operator ()() const") from a __host__ __device__ function("__invoke") is not allowed
{ return static_cast<_Fp&&>(__f)(static_cast<_Args&&>(__args)...); }
^
detected during:
instantiation of "decltype((<expression>)) cuda::std::__4::__invoke(_Fp &&, _Args &&...) [with _Fp=const lambda []()->bool &, _Args=<>]" at line 126 of /opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/functional
instantiation of "_Ret cuda::__4::__detail::__return_type_wrapper<_Ret, _DecayFn>::operator()(_As &&...) const & noexcept [with _Ret=bool, _DecayFn=lambda []()->bool, _As=<>]" at line 6 of <source>
/opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/std/detail/libcxx/include/__functional/invoke.h(402): error: calling a __device__ function("main::[lambda() (instance 1)]::operator ()() const") from a __host__ __device__ function("__invoke") is not allowed
{ return static_cast<_Fp&&>(__f)(static_cast<_Args&&>(__args)...); }
^
detected during:
instantiation of "decltype((<expression>)) cuda::std::__4::__invoke(_Fp &&, _Args &&...) [with _Fp=const lambda []()->bool &, _Args=<>]" at line 126 of /opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/functional
instantiation of "_Ret cuda::__4::__detail::__return_type_wrapper<_Ret, _DecayFn>::operator()(_As &&...) const & noexcept [with _Ret=bool, _DecayFn=lambda []()->bool, _As=<>]" at line 6 of <source>
2 errors detected in the compilation of "<source>".
ASM generation compiler returned: 2
/opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/std/detail/libcxx/include/__functional/invoke.h(402): error: calling a __device__ function("main::[lambda() (instance 1)]::operator ()() const") from a __host__ __device__ function("__invoke") is not allowed
{ return static_cast<_Fp&&>(__f)(static_cast<_Args&&>(__args)...); }
^
detected during:
instantiation of "decltype((<expression>)) cuda::std::__4::__invoke(_Fp &&, _Args &&...) [with _Fp=const lambda []()->bool &, _Args=<>]" at line 126 of /opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/functional
instantiation of "_Ret cuda::__4::__detail::__return_type_wrapper<_Ret, _DecayFn>::operator()(_As &&...) const & noexcept [with _Ret=bool, _DecayFn=lambda []()->bool, _As=<>]" at line 6 of <source>
/opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/std/detail/libcxx/include/__functional/invoke.h(402): error: calling a __device__ function("main::[lambda() (instance 1)]::operator ()() const") from a __host__ __device__ function("__invoke") is not allowed
{ return static_cast<_Fp&&>(__f)(static_cast<_Args&&>(__args)...); }
^
detected during:
instantiation of "decltype((<expression>)) cuda::std::__4::__invoke(_Fp &&, _Args &&...) [with _Fp=const lambda []()->bool &, _Args=<>]" at line 126 of /opt/compiler-explorer/libs/libcudacxx/trunk/include/cuda/functional
instantiation of "_Ret cuda::__4::__detail::__return_type_wrapper<_Ret, _DecayFn>::operator()(_As &&...) const & noexcept [with _Ret=bool, _DecayFn=lambda []()->bool, _As=<>]" at line 6 of <source>
2 errors detected in the compilation of "<source>".
Execution build compiler returned: 2
from libcudacxx.
bdice
commented on June 12, 2024
I'm using f in a context where it is both called by g and used as a predicate function for a thrust call by itself. The f lambda needs to have a "return type proclamation" so it can be used by thrust, but that prevents it from being captured and used in the g lambda.
If I remove the cuda::proclaim_return_type from around f and wrap it separately, the program compiles: https://godbolt.org/z/oonnWqhxq
#include <cuda/functional>
int main()
{
auto f = [] __device__() { return false; };
auto f_ret = cuda::proclaim_return_type<bool>(f);
auto g = cuda::proclaim_return_type<bool>([f] __device__() { return f(); });
return 0;
}from libcudacxx.
Related Issues (20)
- [QUESTION] Can we remove the `isnan` check in the `complex<T>` multiplication operator? HOT 6
- [Question] understanding documentation of cuda::atomic::atomic_thread_fence HOT 2
- CUDA Tuple Structured Binding Declaration Broken HOT 4
- libcu++ should document policy for mixing a custom version with the CUDA Toolkit HOT 2
- `error: "cuda::std::cuda" is ambiguous` when including <cuda/std> HOT 6
- Issue with cuda::atomic_ref::notify_one and wait HOT 1
- Feature request: fp32 multiplication with zero optimization specifier HOT 5
- Copyright in `nv` files HOT 2
- A device variable cannot be marked constexpr (cuda::std::swap / CTK 11.2) HOT 2
- Add `cuda::std::expected`
- Placement new creates performance issues
- libcu++ should not define __host__/__device__/__forceinline__ (or any other runtime macros) HOT 2
- atomic fence questions HOT 2
- Hide if target bits from NVCC/NVRTC
- Function objects in `<functional>` need to be annotated with `nv_exec_check_disable` HOT 1
- Does `cuda::std::complex` have opt-in for simplified complex math HOT 3
- libcu++ should build with gcc-6
- Passing `cuda::proclaim_return_type` to `thrust::make_transform_iterator` gives compilation error HOT 1
- Support fixed-width floating point types
Recommend Projects
-
React
A declarative, efficient, and flexible JavaScript library for building user interfaces.
-
Vue.js
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
-
Typescript
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
-
TensorFlow
An Open Source Machine Learning Framework for Everyone
-
Django
The Web framework for perfectionists with deadlines.
-
Laravel
A PHP framework for web artisans
-
D3
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
-
Recommend Topics
-
javascript
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
-
web
Some thing interesting about web. New door for the world.
-
server
A server is a program made to process requests and deliver data to clients.
-
Machine learning
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
-
Visualization
Some thing interesting about visualization, use data art
-
Game
Some thing interesting about game, make everyone happy.
Recommend Org
-
Facebook
We are working to build community through open source technology. NB: members must have two-factor auth.
-
Microsoft
Open source projects and samples from Microsoft.
-
Google
Google ❤️ Open Source for everyone.
-
Alibaba
Alibaba Open Source for everyone
-
D3
Data-Driven Documents codes.
-
Tencent
China tencent open source team.
from libcudacxx.