NVIDIA / NVIDIA/cccl

Add additional protections against undefined uses of CUDA extended lambdas

Open
#1,004 2 comments 1 reaction 0 assignees View on GitHub
good first issue libcu++
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.