[Mlir-commits] [mlir] [MLIR][NVVM] Fix assertion failure for insufficient parsing validation of nvvm dialect PureSpecialRangeableRegisterOp (PR #163434)

[Guray Ozen]

grypp wrote:

Not directly related to this PR, but I was wondering — can we always add and verify the `range` attribute in a meaningful way? I implemented this downstream, maybe we should have this nicely in NVVM dialect.

For example, our upper limits have been fixed since CUDA 2.0:

```
threadIdx.x -> 0 - 1024
threadIdx.y -> 0 - 1024
threadIdx.z -> 0 - 64

blockIdx.x  -> 0 - 2^32 - 1
blockIdx.y  -> 0 - 65535
blockIdx.z  -> 0 - 65535

laneid      -> 0 - 32
warpsize    -> 32
```

We could also use `#target` SM information if we expect these numbers to change in the future.


https://github.com/llvm/llvm-project/pull/163434