[CUDA] Assumptions on `__nvvm_read_ptx_sreg_meow()` functions
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
When working on `cuda::launch` in the CCCL project, we often have compile time information about the dimensions of block/cluster/grid. When using `nvcc`, we are able to do:
```cpp
template
__global__ void kernel()
{
__builtin_assume(blockDim.x == block_dim_x);
if constexpr (block_dim_x == 1)
{
__builtin_assume(threadIdx.x == 0);
}
// ...
}
```
Which can replace the builtin variable reads with values and produce a better optimized code as a result.
When I try to do same thing with `clang-cuda`, I get a warning about the assumption having side effects and being ignored. I thought this is caused by the way Clang implements these builtin variables, but I get the same behaviour when using the corresponding `__nvvm_read_ptx_sreg_meow()` builtin functions instead.
Is there a way to make this work, or one can just not make assumptions on builtin functions?
Thank you!
Contributor guide
Research direction
Start by reproducing the warning with clang-cuda using __builtin_assume and the corresponding __nvvm_read_ptx_sreg_meow() builtins, as described in the issue. Trace how clang-cuda handles these builtin calls and determine whether assumptions are expected to be accepted; done means establishing a confirmed supported behavior or a clearly documented limitation.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100