squiid/llvm
1
0
Fork 0
Code Issues Pull requests Activity

[AMDGPU] Mark s_get_waveid_in_workgroup as not reading memory

Browse source 
It is already marked as having side effects, at least in MIR. It does
not interact with anything else that is modelled as a memory access
either in IR or MachineIR.

Differential Revision: https://reviews.llvm.org/D125985
This commit is contained in:
Jay Foad 2022-05-19 16:52:41 +01:00
parent 86b55edab6
commit 9ece051847
2 changed files with 2 additions and 2 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View file

@ -1860,7 +1860,7 @@ def int_amdgcn_mov_dpp8 :
def int_amdgcn_s_get_waveid_in_workgroup :
GCCBuiltin<"__builtin_amdgcn_s_get_waveid_in_workgroup">,
Intrinsic<[llvm_i32_ty], [],
[IntrReadMem, IntrInaccessibleMemOnly, IntrWillReturn]>;
[IntrNoMem, IntrHasSideEffects, IntrWillReturn]>;
class AMDGPUGlobalAtomicRtn<LLVMType vt> : Intrinsic <
[vt],

2
llvm/lib/Target/AMDGPU/SMInstructions.td
View file

@ -235,7 +235,7 @@ class SM_WaveId_Pseudo<string opName, SDPatternOperator node> : SM_Pseudo<
" $sdst", [(set i32:$sdst, (node))]> {
let hasSideEffects = 1;
let mayStore = 0;
let mayLoad = 1;
let mayLoad = 0;
let has_sbase = 0;
}