NVIDIA / NVIDIA/cccl

[EPIC]: Improve usability of architecture specific features in libcudacxx

Open
#1,083 1 comment 1 reaction 1 assignee Claimed by @wmaxey View on GitHub
Dominant language
C++
Stars
2.5k
Forks
487
Avg merge
2d 7h
Merged PRs (30d)
296

Description

### Is this a duplicate?

- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Area

libcu++

### Is your feature request related to a problem? Please describe.

As a CUDA developer using libcu++, I want to be able to use architecture dependent features of libcudacxx in my CUDA application. For any given libcudacxx header and feature, I need to be able to do the following:

```
#include

__global__ void kernel(...){
NV_DISPATCH_TARGET(
NV_IS_EXACTLY_SM_60, ( do_sm60_thing(); ),
NV_PROVIDES_SM_70, ( do_sm70_thing(); ),
NV_PROVIDES_SM_90, ( do_sm90_thing(); )
)
}
```

I need to be able to compile this file with any set of architectures (`-gencode arch=compute_XX,code=sm_XX`) and for it to be able to compile and link successfully so long as I am always careful to use an architecture dependent feature in an appropriately guarded code path, whether using `NV_IF_TARGET` or `__CUDA_ARCH__`.

However, this does not work universally today. For example, the following fails to compile when compiled with `-gencode arch=compute_52,code=sm_52 -gencode arch=compute_70,code=sm_70`

```
#include
#include

__global__ void kernel(){
NV_IF_TARGET(
NV_PROVIDES_SM_70,
cuda::atomic i;
)
}
```
https://godbolt.org/z/ddMaW65Ej

This is because the `cuda/atomic` header will unconditionally error any time it is included in a TU that compiles for an architecture less than `sm60`, even if the feature is never used in code paths for the unsupported architecture.

A similar problem exists with `cuda/barrier`: https://godbolt.org/z/aEjsMT5YK

### Describe the solution you'd like

I should be able to do the following with all libcu++ headers and features:

```cpp
#include

__global__ void kernel(...){
NV_DISPATCH_TARGET(
NV_IS_EXACTLY_SM_60, ( do_sm60_thing(); ),
NV_PROVIDES_SM_70, ( do_sm70_thing(); ),
NV_PROVIDES_SM_90, ( do_sm90_thing(); )
)
}
```

### Tasks
- [ ] https://github.com/NVIDIA/cccl/issues/1084
- [ ] Enable using relevant parts of `` on sm_52
- [ ] Enable including `` on sm_52

### Describe alternatives you've considered

If libcu++ doesn't do this, then I am forced to use lower level things like `atomicAdd()` or inline PTX.

### Additional context
Related issues:
https://github.com/NVIDIA/cccl/issues/997
https://github.com/NVIDIA/cccl/issues/1082
https://github.com/NVIDIA/cccl/issues/624

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.