| Issue |
181340
|
| Summary |
[CUDA] Assumptions on `__nvvm_read_ptx_sreg_meow()` functions
|
| Labels |
new issue
|
| Assignees |
|
| Reporter |
davebayer
|
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 <unsigned block_dim_x>
__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!
![]()
_______________________________________________
llvm-bugs mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
