diff options
| -rw-r--r-- | llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 4 | ||||
| -rw-r--r-- | llvm/lib/Target/AMDGPU/SMInstructions.td | 7 |
2 files changed, 4 insertions, 7 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index d72a8c6798a..d7999cd3323 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -570,7 +570,7 @@ def int_amdgcn_s_dcache_inv : def int_amdgcn_s_memtime : GCCBuiltin<"__builtin_amdgcn_s_memtime">, - Intrinsic<[llvm_i64_ty], [], []>; + Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>; def int_amdgcn_s_sleep : GCCBuiltin<"__builtin_amdgcn_s_sleep">, @@ -816,7 +816,7 @@ def int_amdgcn_s_dcache_wb_vol : def int_amdgcn_s_memrealtime : GCCBuiltin<"__builtin_amdgcn_s_memrealtime">, - Intrinsic<[llvm_i64_ty], [], []>; + Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>; // llvm.amdgcn.ds.permute <index> <src> def int_amdgcn_ds_permute : diff --git a/llvm/lib/Target/AMDGPU/SMInstructions.td b/llvm/lib/Target/AMDGPU/SMInstructions.td index 5e72a2e8828..8f347986eb8 100644 --- a/llvm/lib/Target/AMDGPU/SMInstructions.td +++ b/llvm/lib/Target/AMDGPU/SMInstructions.td @@ -129,11 +129,8 @@ class SM_Time_Pseudo<string opName, SDPatternOperator node> : SM_Pseudo< opName, (outs SReg_64_XEXEC:$sdst), (ins), " $sdst", [(set i64:$sdst, (node))]> { let hasSideEffects = 1; - // FIXME: mayStore = ? is a workaround for tablegen bug for different - // inferred mayStore flags for the instruction pattern vs. standalone - // Pat. Each considers the other contradictory. - let mayStore = ?; - let mayLoad = ?; + let mayStore = 0; + let mayLoad = 1; let has_sbase = 0; let has_offset = 0; } |
