This is an archive of the discontinued LLVM Phabricator instance.

Differential D147589

Mark threadIdx.x and friends as noundef.
ClosedPublic

Authored by jlebar on Apr 4 2023, 11:39 PM.

Details

Reviewers
tra
Commits
rG302578d90162: Mark threadIdx.x and friends as noundef.
Summary

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.

Diff Detail

Event Timeline

jlebar created this revision.Apr 4 2023, 11:39 PM
Herald added a project: Restricted Project. · View Herald TranscriptApr 4 2023, 11:39 PM
Herald added subscribers: mattd, gchakrabarti. · View Herald Transcript
jlebar requested review of this revision.Apr 4 2023, 11:39 PM
Herald added a project: Restricted Project. · View Herald TranscriptApr 4 2023, 11:39 PM
Herald added subscribers: llvm-commits, jholewinski. · View Herald Transcript
Harbormaster completed remote builds in B223727: Diff 511000.Apr 5 2023, 12:12 AM
tra accepted this revision.Apr 5 2023, 10:06 AM
This revision is now accepted and ready to land.Apr 5 2023, 10:06 AM
This revision was landed with ongoing or failed builds.Apr 5 2023, 1:44 PM
Closed by commit rG302578d90162: Mark threadIdx.x and friends as noundef. (authored by jlebar). · Explain Why
This revision was automatically updated to reflect the committed changes.
jlebar added a commit: rG302578d90162: Mark threadIdx.x and friends as noundef..

Revision Contents

PathSize
llvm/
include/
llvm/
IR/
84 lines

Diff 511201

llvm/include/llvm/IR/IntrinsicsNVVM.td

Show First 20 Lines • Show All 1,594 Lines • ▼ Show 20 Lines
def int_nvvm_isspacep_shared def int_nvvm_isspacep_shared
: DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty], : DefaultAttrsIntrinsic<[llvm_i1_ty], [llvm_ptr_ty],
[IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>], [IntrNoMem, IntrSpeculatable, NoCapture<ArgIndex<0>>],
"llvm.nvvm.isspacep.shared">, "llvm.nvvm.isspacep.shared">,
ClangBuiltin<"__nvvm_isspacep_shared">; ClangBuiltin<"__nvvm_isspacep_shared">;
// Environment register read // Environment register read
def int_nvvm_read_ptx_sreg_envreg0 def int_nvvm_read_ptx_sreg_envreg0
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg0">, "llvm.nvvm.read.ptx.sreg.envreg0">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg0">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg0">;
def int_nvvm_read_ptx_sreg_envreg1 def int_nvvm_read_ptx_sreg_envreg1
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg1">, "llvm.nvvm.read.ptx.sreg.envreg1">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg1">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg1">;
def int_nvvm_read_ptx_sreg_envreg2 def int_nvvm_read_ptx_sreg_envreg2
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg2">, "llvm.nvvm.read.ptx.sreg.envreg2">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg2">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg2">;
def int_nvvm_read_ptx_sreg_envreg3 def int_nvvm_read_ptx_sreg_envreg3
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg3">, "llvm.nvvm.read.ptx.sreg.envreg3">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg3">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg3">;
def int_nvvm_read_ptx_sreg_envreg4 def int_nvvm_read_ptx_sreg_envreg4
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg4">, "llvm.nvvm.read.ptx.sreg.envreg4">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg4">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg4">;
def int_nvvm_read_ptx_sreg_envreg5 def int_nvvm_read_ptx_sreg_envreg5
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg5">, "llvm.nvvm.read.ptx.sreg.envreg5">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg5">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg5">;
def int_nvvm_read_ptx_sreg_envreg6 def int_nvvm_read_ptx_sreg_envreg6
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg6">, "llvm.nvvm.read.ptx.sreg.envreg6">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg6">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg6">;
def int_nvvm_read_ptx_sreg_envreg7 def int_nvvm_read_ptx_sreg_envreg7
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg7">, "llvm.nvvm.read.ptx.sreg.envreg7">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg7">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg7">;
def int_nvvm_read_ptx_sreg_envreg8 def int_nvvm_read_ptx_sreg_envreg8
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg8">, "llvm.nvvm.read.ptx.sreg.envreg8">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg8">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg8">;
def int_nvvm_read_ptx_sreg_envreg9 def int_nvvm_read_ptx_sreg_envreg9
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg9">, "llvm.nvvm.read.ptx.sreg.envreg9">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg9">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg9">;
def int_nvvm_read_ptx_sreg_envreg10 def int_nvvm_read_ptx_sreg_envreg10
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg10">, "llvm.nvvm.read.ptx.sreg.envreg10">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg10">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg10">;
def int_nvvm_read_ptx_sreg_envreg11 def int_nvvm_read_ptx_sreg_envreg11
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg11">, "llvm.nvvm.read.ptx.sreg.envreg11">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg11">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg11">;
def int_nvvm_read_ptx_sreg_envreg12 def int_nvvm_read_ptx_sreg_envreg12
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg12">, "llvm.nvvm.read.ptx.sreg.envreg12">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg12">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg12">;
def int_nvvm_read_ptx_sreg_envreg13 def int_nvvm_read_ptx_sreg_envreg13
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg13">, "llvm.nvvm.read.ptx.sreg.envreg13">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg13">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg13">;
def int_nvvm_read_ptx_sreg_envreg14 def int_nvvm_read_ptx_sreg_envreg14
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg14">, "llvm.nvvm.read.ptx.sreg.envreg14">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg14">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg14">;
def int_nvvm_read_ptx_sreg_envreg15 def int_nvvm_read_ptx_sreg_envreg15
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg15">, "llvm.nvvm.read.ptx.sreg.envreg15">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg15">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg15">;
def int_nvvm_read_ptx_sreg_envreg16 def int_nvvm_read_ptx_sreg_envreg16
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg16">, "llvm.nvvm.read.ptx.sreg.envreg16">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg16">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg16">;
def int_nvvm_read_ptx_sreg_envreg17 def int_nvvm_read_ptx_sreg_envreg17
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg17">, "llvm.nvvm.read.ptx.sreg.envreg17">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg17">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg17">;
def int_nvvm_read_ptx_sreg_envreg18 def int_nvvm_read_ptx_sreg_envreg18
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg18">, "llvm.nvvm.read.ptx.sreg.envreg18">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg18">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg18">;
def int_nvvm_read_ptx_sreg_envreg19 def int_nvvm_read_ptx_sreg_envreg19
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg19">, "llvm.nvvm.read.ptx.sreg.envreg19">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg19">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg19">;
def int_nvvm_read_ptx_sreg_envreg20 def int_nvvm_read_ptx_sreg_envreg20
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg20">, "llvm.nvvm.read.ptx.sreg.envreg20">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg20">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg20">;
def int_nvvm_read_ptx_sreg_envreg21 def int_nvvm_read_ptx_sreg_envreg21
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg21">, "llvm.nvvm.read.ptx.sreg.envreg21">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg21">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg21">;
def int_nvvm_read_ptx_sreg_envreg22 def int_nvvm_read_ptx_sreg_envreg22
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg22">, "llvm.nvvm.read.ptx.sreg.envreg22">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg22">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg22">;
def int_nvvm_read_ptx_sreg_envreg23 def int_nvvm_read_ptx_sreg_envreg23
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg23">, "llvm.nvvm.read.ptx.sreg.envreg23">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg23">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg23">;
def int_nvvm_read_ptx_sreg_envreg24 def int_nvvm_read_ptx_sreg_envreg24
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg24">, "llvm.nvvm.read.ptx.sreg.envreg24">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg24">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg24">;
def int_nvvm_read_ptx_sreg_envreg25 def int_nvvm_read_ptx_sreg_envreg25
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg25">, "llvm.nvvm.read.ptx.sreg.envreg25">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg25">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg25">;
def int_nvvm_read_ptx_sreg_envreg26 def int_nvvm_read_ptx_sreg_envreg26
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg26">, "llvm.nvvm.read.ptx.sreg.envreg26">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg26">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg26">;
def int_nvvm_read_ptx_sreg_envreg27 def int_nvvm_read_ptx_sreg_envreg27
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg27">, "llvm.nvvm.read.ptx.sreg.envreg27">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg27">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg27">;
def int_nvvm_read_ptx_sreg_envreg28 def int_nvvm_read_ptx_sreg_envreg28
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg28">, "llvm.nvvm.read.ptx.sreg.envreg28">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg28">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg28">;
def int_nvvm_read_ptx_sreg_envreg29 def int_nvvm_read_ptx_sreg_envreg29
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg29">, "llvm.nvvm.read.ptx.sreg.envreg29">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg29">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg29">;
def int_nvvm_read_ptx_sreg_envreg30 def int_nvvm_read_ptx_sreg_envreg30
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg30">, "llvm.nvvm.read.ptx.sreg.envreg30">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg30">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg30">;
def int_nvvm_read_ptx_sreg_envreg31 def int_nvvm_read_ptx_sreg_envreg31
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable], : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>],
"llvm.nvvm.read.ptx.sreg.envreg31">, "llvm.nvvm.read.ptx.sreg.envreg31">,
ClangBuiltin<"__nvvm_read_ptx_sreg_envreg31">; ClangBuiltin<"__nvvm_read_ptx_sreg_envreg31">;
// Texture Fetch // Texture Fetch
// texmode_independent // texmode_independent
def int_nvvm_tex_1d_v4f32_s32 def int_nvvm_tex_1d_v4f32_s32
: Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
▲ Show 20 Lines • Show All 2,616 Lines • ▼ Show 20 Lines
// Accessing special registers. // Accessing special registers.
multiclass PTXReadSRegIntrinsic_v4i32<string regname> { multiclass PTXReadSRegIntrinsic_v4i32<string regname> {
// FIXME: Do we need the 128-bit integer type version? // FIXME: Do we need the 128-bit integer type version?
// def _r64 : Intrinsic<[llvm_i128_ty], [], [IntrNoMem, IntrSpeculatable]>; // def _r64 : Intrinsic<[llvm_i128_ty], [], [IntrNoMem, IntrSpeculatable]>;
// FIXME: Enable this once v4i32 support is enabled in back-end. // FIXME: Enable this once v4i32 support is enabled in back-end.
// def _v4i16 : Intrinsic<[llvm_v4i32_ty], [], [IntrNoMem, IntrSpeculatable]>; // def _v4i16 : Intrinsic<[llvm_v4i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
def _x : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, def _x : DefaultAttrsIntrinsic<[llvm_i32_ty], [],
[IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_x">; ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_x">;
def _y : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, def _y : DefaultAttrsIntrinsic<[llvm_i32_ty], [],
[IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_y">; ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_y">;
def _z : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, def _z : DefaultAttrsIntrinsic<[llvm_i32_ty], [],
[IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_z">; ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_z">;
def _w : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, def _w : DefaultAttrsIntrinsic<[llvm_i32_ty], [],
[IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_w">; ClangBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_w">;
} }
class PTXReadSRegIntrinsic_r32<string name> class PTXReadSRegIntrinsic_r32<string name>
: DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>, : DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>; ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>;
class PTXReadSRegIntrinsic_r64<string name> class PTXReadSRegIntrinsic_r64<string name>
: DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>, : DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>; ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>;
// Intrinsics to read registers with non-constant values. E.g. the values that // Intrinsics to read registers with non-constant values. E.g. the values that
// do change over the kernel lifetime. Such reads should not be CSE'd. // do change over the kernel lifetime. Such reads should not be CSE'd.
class PTXReadNCSRegIntrinsic_r32<string name> class PTXReadNCSRegIntrinsic_r32<string name>
: Intrinsic<[llvm_i32_ty], [], [IntrInaccessibleMemOnly, IntrNoCallback]>, : Intrinsic<[llvm_i32_ty], [], [IntrInaccessibleMemOnly, IntrNoCallback, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>; ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>;
class PTXReadNCSRegIntrinsic_r64<string name> class PTXReadNCSRegIntrinsic_r64<string name>
: Intrinsic<[llvm_i64_ty], [], [IntrInaccessibleMemOnly, IntrNoCallback]>, : Intrinsic<[llvm_i64_ty], [], [IntrInaccessibleMemOnly, IntrNoCallback, NoUndef<RetIndex>]>,
ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>; ClangBuiltin<"__nvvm_read_ptx_sreg_" # name>;
defm int_nvvm_read_ptx_sreg_tid : PTXReadSRegIntrinsic_v4i32<"tid">; defm int_nvvm_read_ptx_sreg_tid : PTXReadSRegIntrinsic_v4i32<"tid">;
defm int_nvvm_read_ptx_sreg_ntid : PTXReadSRegIntrinsic_v4i32<"ntid">; defm int_nvvm_read_ptx_sreg_ntid : PTXReadSRegIntrinsic_v4i32<"ntid">;
def int_nvvm_read_ptx_sreg_laneid : PTXReadSRegIntrinsic_r32<"laneid">; def int_nvvm_read_ptx_sreg_laneid : PTXReadSRegIntrinsic_r32<"laneid">;
def int_nvvm_read_ptx_sreg_warpid : PTXReadSRegIntrinsic_r32<"warpid">; def int_nvvm_read_ptx_sreg_warpid : PTXReadSRegIntrinsic_r32<"warpid">;
def int_nvvm_read_ptx_sreg_nwarpid : PTXReadSRegIntrinsic_r32<"nwarpid">; def int_nvvm_read_ptx_sreg_nwarpid : PTXReadSRegIntrinsic_r32<"nwarpid">;
▲ Show 20 Lines • Show All 278 Lines • Show Last 20 Lines