llvm / llvm/llvm-project

[CUDA] Assumptions on `__nvvm_read_ptx_sreg_meow()` functions

Open
#181,340 0 comments 0 reactions 0 assignees View on GitHub
cuda
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

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.