[PATCH] D147589: Mark threadIdx.x and friends as noundef.

[Justin Lebar via Phabricator via llvm-commits]

jlebar created this revision.
jlebar added a reviewer: tra.
Herald added subscribers: mattd, gchakrabarti.
Herald added a project: All.
jlebar requested review of this revision.
Herald added subscribers: llvm-commits, jholewinski.
Herald added a project: LLVM.

threadIdx.x and similar functions never return undef.

Simple enough to say, but why does it matter?

Consider the following IR that reads threadIdx.x and blockIdx.x.

  %a = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !range !138
  %b = call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !range !139
  %c = shl nuw nsw i32 %a, 6
  %linear_index = or i32 %c, %b
  %linear_index.fr = freeze i32 %linear_index

If %a or %b may be undef, then computeKnownBits will not recurse through
the freeze instruction.  Therefore we will not know anything about the
known bits of linear_index.fr, even though we have range metadata!  Bad
Things fall out of this.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D147589

Files:
  llvm/include/llvm/IR/IntrinsicsNVVM.td

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D147589.511000.patch
Type: text/x-patch
Size: 12920 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20230405/1bf0295e/attachment.bin>