Add additional protections against undefined uses of CUDA extended lambdas
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
[CUDA Extended Lambdas](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda) (lambdas with `__device__` or `__host__ __device__`) annotations are a useful convenience, but they come with [several restrictions](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda-restrictions) that can cause silent or confusing failures in user's code.
For many of these restrictions, libcu++ is powerless to do anything to help users, but where possible, libcu++ should make an effort to detect invalid uses of extended lambdas and emit a more helpful diagnostic.
Some of the most pertinent restrictions stem from the fact that nvcc replaces extended lambdas with a placeholder type whose `operator()` is not equivalent to the lambda definition
> 14. As described above, the CUDA compiler replaces a __device__ extended lambda defined in a host function with a placeholder type defined in namespace scope. This placeholder type does not define a operator() function equivalent to the original lambda declaration. An attempt to determine the return type or parameter types of the operator() function may therefore work incorrectly in host code, as the code processed by the host compiler will be semantically different than the input code processed by the CUDA compiler. However, it is OK to introspect the return type or parameter types of the operator() function within device code. Note that this restriction does not apply to __host__ __device__ extended lambdas.
> 17. As described previously, the CUDA compiler replaces an extended __device__ or __host__ __device__ lambda expression with an instance of a placeholder type in the code sent to the host compiler. This placeholder type may define C++ special member functions (e.g. constructor, destructor). As a result, some standard C++ type traits may return different results for the closure type of the extended lambda, in the CUDA frontend compiler versus the host compiler.
nvcc provides the `__nv_is_extended_device_lambda_closure_type(X)` and `__nv_is_extended_host_device_lambda_closure_type(X)` built-in traits to detect a `__device__` or `__host__ __device__` lambda at compile time. This enables us to detect and emit compile time diagnostics for invalid uses with libcu++ constructs.
For example, one of the restrictions on extended lambdas is that you cannot query their return type in host code, so in https://github.com/NVIDIA/libcudacxx/pull/284 we updated `cuda::std::invoke_result_t` to emit a compile time error when used in host code.
As mentioned in restriction 17, there are several other type traits where we should do similar changes as was done for `cuda::std::invoke_result` (note that unlike with `invoke_result`, the following traits should be guarded for _both_ `__device__` and `__host__ __device__` lambdas).
### Tasks
- [ ] https://github.com/NVIDIA/cccl/issues/1984
I also suspect there are changes we can/should make to things in `` like `cuda::std::invoke`, but that will require additional investigation.
Contributor guide
Research direction
Start with the cuda::std::invoke_result_t changes from pull request 284 and the linked task in issue 1984, then inspect the related type traits and the entry points mentioned here. Determine which host and device uses need diagnostics, and consider the work complete when invalid extended-lambda uses are consistently rejected with compile-time errors.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 30/100