diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -1379,11 +1379,11 @@ Intrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem, IntrHasSideEffects, IntrWillReturn]>; +// This is IntrHasSideEffects so it can be used to read cycle counters. def int_amdgcn_s_getreg : GCCBuiltin<"__builtin_amdgcn_s_getreg">, Intrinsic<[llvm_i32_ty], [llvm_i32_ty], - [IntrInaccessibleMemOnly, IntrReadMem, IntrSpeculatable, - IntrWillReturn, ImmArg<ArgIndex<0>>] + [IntrNoMem, IntrHasSideEffects, IntrWillReturn, ImmArg<ArgIndex<0>>] >; // Note this can be used to set FP environment properties that are diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td --- a/llvm/lib/Target/AMDGPU/SOPInstructions.td +++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td @@ -873,9 +873,7 @@ "$sdst, $simm16" >; -let mayLoad = 1 in { -// s_getreg_b32 should use hasSideEffects = 1 for tablegen to allow -// its use in the readcyclecounter selection. +// This is hasSideEffects to allow its use in readcyclecounter selection. // FIXME: Need to truncate immediate to 16-bits. def S_GETREG_B32 : SOPK_Pseudo < "s_getreg_b32", @@ -885,7 +883,6 @@ let SOPKZext = 1; let hasSideEffects = 1; } -} // End mayLoad = 1 let Defs = [MODE], Uses = [MODE] in {