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

[Justin Lebar via All-commits]

  Branch: refs/heads/nvvm-undef
  Home:   https://github.com/llvm/llvm-project
  Commit: 7b0ddb98042c69487014b00673e29a8353f5e28f
      https://github.com/llvm/llvm-project/commit/7b0ddb98042c69487014b00673e29a8353f5e28f
  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