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; Index: llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll =================================================================== --- llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll +++ llvm/test/CodeGen/AMDGPU/llvm.amdgcn.buffer.wbinvl1.vol.ll @@ -5,13 +5,14 @@ ; GCN-LABEL: {{^}}test_buffer_wbinvl1_vol: ; GCN-NEXT: ; %bb.0: -; CI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00] -; VI-NEXT: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00] -; GCN: s_endpgm -define amdgpu_kernel void @test_buffer_wbinvl1_vol() #0 { +; CI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xc0,0xe1,0x00,0x00,0x00,0x00] +; VI: buffer_wbinvl1_vol ; encoding: [0x00,0x00,0xfc,0xe0,0x00,0x00,0x00,0x00] +; GCN: _store_byte +; GCN-NEXT: s_endpgm +define amdgpu_kernel void @test_buffer_wbinvl1_vol(i8 addrspace(1)* %ptr) #0 { call void @llvm.amdgcn.buffer.wbinvl1.vol() ; This used to crash in hazard recognizer - store i8 0, i8 addrspace(1)* undef, align 1 + store i8 0, i8 addrspace(1)* %ptr, align 1 ret void } Index: llvm/test/CodeGen/AMDGPU/scheduler-handle-move-bundle.mir =================================================================== --- llvm/test/CodeGen/AMDGPU/scheduler-handle-move-bundle.mir +++ llvm/test/CodeGen/AMDGPU/scheduler-handle-move-bundle.mir @@ -23,8 +23,8 @@ ; GCN: [[V_MOV_B32_e32_:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 1, implicit $exec ; GCN: [[V_MOV_B32_e32_1:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 0, implicit $exec ; GCN: $vcc_hi = IMPLICIT_DEF - ; GCN: DS_WRITE_B32_gfx9 [[V_MOV_B32_e32_1]], [[V_MOV_B32_e32_]], 0, 0, implicit $exec :: (store 4, addrspace 3) ; GCN: [[V_MOV_B32_e32_2:%[0-9]+]]:vgpr_32 = V_MOV_B32_e32 2, implicit $exec + ; GCN: DS_WRITE_B32_gfx9 [[V_MOV_B32_e32_1]], [[V_MOV_B32_e32_]], 0, 0, implicit $exec :: (store 4, addrspace 3) ; GCN: $m0 = S_MOV_B32 0 ; GCN: $vgpr0 = COPY [[S_LOAD_DWORD_IMM]] ; GCN: BUNDLE implicit $vgpr0, implicit $m0, implicit $exec {