[all-commits] [llvm/llvm-project] 302578: Mark threadIdx.x and friends as noundef.

[Justin Lebar via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 302578d9016299b5c1c22d580de4ccbc68f7a120
      https://github.com/llvm/llvm-project/commit/302578d9016299b5c1c22d580de4ccbc68f7a120
  Author: Justin Lebar <justin.lebar at gmail.com>
  Date:   2023-04-05 (Wed, 05 Apr 2023)

  Changed paths:
    M llvm/include/llvm/IR/IntrinsicsNVVM.td

  Log Message:
  -----------
  Mark threadIdx.x and friends as noundef.

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.

Differential Revision: https://reviews.llvm.org/D147589