Index: llvm/include/llvm/IR/IntrinsicsAMDGPU.td =================================================================== --- llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -209,10 +209,10 @@ Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent]>; def int_amdgcn_wave_barrier : GCCBuiltin<"__builtin_amdgcn_wave_barrier">, - Intrinsic<[], [], [IntrConvergent]>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects, IntrConvergent]>; def int_amdgcn_s_waitcnt : GCCBuiltin<"__builtin_amdgcn_s_waitcnt">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>]>; + Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects]>; def int_amdgcn_div_scale : Intrinsic< // 1st parameter: Numerator @@ -1172,15 +1172,15 @@ def int_amdgcn_buffer_wbinvl1_sc : GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_sc">, - Intrinsic<[], [], []>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; def int_amdgcn_buffer_wbinvl1 : GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1">, - Intrinsic<[], [], []>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; def int_amdgcn_s_dcache_inv : GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">, - Intrinsic<[], [], []>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; def int_amdgcn_s_memtime : GCCBuiltin<"__builtin_amdgcn_s_memtime">, @@ -1188,17 +1188,17 @@ def int_amdgcn_s_sleep : GCCBuiltin<"__builtin_amdgcn_s_sleep">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>]> { + Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects]> { } def int_amdgcn_s_incperflevel : GCCBuiltin<"__builtin_amdgcn_s_incperflevel">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>]> { + Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects]> { } def int_amdgcn_s_decperflevel : GCCBuiltin<"__builtin_amdgcn_s_decperflevel">, - Intrinsic<[], [llvm_i32_ty], [ImmArg>]> { + Intrinsic<[], [llvm_i32_ty], [ImmArg>, IntrNoMem, IntrHasSideEffects]> { } def int_amdgcn_s_getreg : @@ -1476,6 +1476,7 @@ >; // If false, set EXEC=0 for the current thread until the end of program. +// FIXME: Should this be IntrNoMem, IntrHasSideEffects? def int_amdgcn_kill : Intrinsic<[], [llvm_i1_ty], []>; // Copies the active channels of the source value to the destination value, @@ -1515,11 +1516,11 @@ def int_amdgcn_s_dcache_inv_vol : GCCBuiltin<"__builtin_amdgcn_s_dcache_inv_vol">, - Intrinsic<[], [], []>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; def int_amdgcn_buffer_wbinvl1_vol : GCCBuiltin<"__builtin_amdgcn_buffer_wbinvl1_vol">, - Intrinsic<[], [], []>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; //===----------------------------------------------------------------------===// // VI Intrinsics @@ -1545,11 +1546,11 @@ def int_amdgcn_s_dcache_wb : GCCBuiltin<"__builtin_amdgcn_s_dcache_wb">, - Intrinsic<[], [], []>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; def int_amdgcn_s_dcache_wb_vol : GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">, - Intrinsic<[], [], []>; + Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>; def int_amdgcn_s_memrealtime : GCCBuiltin<"__builtin_amdgcn_s_memrealtime">, Index: llvm/lib/Target/AMDGPU/BUFInstructions.td =================================================================== --- llvm/lib/Target/AMDGPU/BUFInstructions.td +++ llvm/lib/Target/AMDGPU/BUFInstructions.td @@ -374,7 +374,8 @@ let AsmMatchConverter = ""; let hasSideEffects = 1; - let mayStore = 1; + let mayLoad = 0; + let mayStore = 0; // Set everything to 0. let offen = 0; Index: llvm/lib/Target/AMDGPU/SIInstructions.td =================================================================== --- llvm/lib/Target/AMDGPU/SIInstructions.td +++ llvm/lib/Target/AMDGPU/SIInstructions.td @@ -274,13 +274,14 @@ def S_ANDN2_B32_term : WrapTerminatorInst; } + def WAVE_BARRIER : SPseudoInstSI<(outs), (ins), [(int_amdgcn_wave_barrier)]> { let SchedRW = []; let hasNoSchedulingInfo = 1; let hasSideEffects = 1; - let mayLoad = 1; - let mayStore = 1; + let mayLoad = 0; + let mayStore = 0; let isConvergent = 1; let FixedSize = 1; let Size = 0; Index: llvm/lib/Target/AMDGPU/SMInstructions.td =================================================================== --- llvm/lib/Target/AMDGPU/SMInstructions.td +++ llvm/lib/Target/AMDGPU/SMInstructions.td @@ -187,7 +187,7 @@ class SM_Inval_Pseudo : SM_Pseudo< opName, (outs), (ins), "", [(node)]> { let hasSideEffects = 1; - let mayStore = 1; + let mayStore = 0; let has_sdst = 0; let has_sbase = 0; let has_offset = 0; Index: llvm/lib/Target/AMDGPU/SOPInstructions.td =================================================================== --- llvm/lib/Target/AMDGPU/SOPInstructions.td +++ llvm/lib/Target/AMDGPU/SOPInstructions.td @@ -1140,7 +1140,7 @@ let mayStore = 1; } -let mayLoad = 1, mayStore = 1, hasSideEffects = 1 in +let mayLoad = 0, mayStore = 0, hasSideEffects = 1 in def S_WAITCNT : SOPP <0x0000000c, (ins WAIT_FLAG:$simm16), "s_waitcnt $simm16", [(int_amdgcn_s_waitcnt timm:$simm16)]>; def S_SETHALT : SOPP <0x0000000d, (ins i16imm:$simm16), "s_sethalt $simm16">; @@ -1153,8 +1153,8 @@ def S_SLEEP : SOPP <0x0000000e, (ins i32imm:$simm16), "s_sleep $simm16", [(int_amdgcn_s_sleep timm:$simm16)]> { let hasSideEffects = 1; - let mayLoad = 1; - let mayStore = 1; + let mayLoad = 0; + let mayStore = 0; } def S_SETPRIO : SOPP <0x0000000f, (ins i16imm:$simm16), "s_setprio $simm16">; @@ -1179,14 +1179,14 @@ def S_INCPERFLEVEL : SOPP <0x00000014, (ins i32imm:$simm16), "s_incperflevel $simm16", [(int_amdgcn_s_incperflevel timm:$simm16)]> { let hasSideEffects = 1; - let mayLoad = 1; - let mayStore = 1; + let mayLoad = 0; + let mayStore = 0; } def S_DECPERFLEVEL : SOPP <0x00000015, (ins i32imm:$simm16), "s_decperflevel $simm16", [(int_amdgcn_s_decperflevel timm:$simm16)]> { let hasSideEffects = 1; - let mayLoad = 1; - let mayStore = 1; + let mayLoad = 0; + let mayStore = 0; } def S_TTRACEDATA : SOPP <0x00000016, (ins), "s_ttracedata"> { let simm16 = 0;