Redesign libcudacxx architecture specific testing
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
# Summary
As described in https://github.com/NVIDIA/cccl/issues/1083, as a user of libcudacxx headers and features, I want to be able to use libcudacxx architecture specific features in a TU compiled for multiple architectures (`-gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70`) so long as I am careful to guard the code paths in which I use those features, e.g.,
```
#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(); )
)
}
```
This is not the case today, and this is indirectly due in part to how libcudacxx tests architecture specific features.
libcudacxx uses `lit` for compiling and running its tests, and it uses the `UNSUPPORTED:` keyword to indicate environments that the test doesn't support, including GPU architectures. For example:
https://github.com/NVIDIA/libcudacxx/blob/206d8f9179deda6006795865c9c61cbd24b5e6cc/.upstream-tests/test/cuda/memcpy_async_16.pass.cpp#L11-L20
The `// UNSUPPORTED: pre-sm-70` means that this test file will be **skipped entirely** if compiling for `();
return 0;
}
```
we want this:
```
#include "memcpy_async.h"
int main(int argc, char ** argv)
{
NV_IF_TARGET(
NV_PROVIDES_SM_70, test_select_source();
)
return 0;
}
```
It would be nice to avoid needing to actually rewrite tests that are currently using `// UNSUPPORTED: pre-sm-70` and could somehow implicitly inject the appropriate `NV_IF_TARGET` logic based on the information in the `UNSUPPORTED:` key. Perhaps by injecting a different `fake_main` that bakes the appropriate `NV_IF_TARGET` logic in.
### Tasks
- [ ] https://github.com/NVIDIA/cccl/issues/1188
- [ ] Refactor existing architecture specific tests using new design
Contributor guide
Assessment
This issue has not been assessed yet.