diff options
| author | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-12-08 20:01:02 +0000 |
|---|---|---|
| committer | Matt Arsenault <Matthew.Arsenault@amd.com> | 2017-12-08 20:01:02 +0000 |
| commit | 73ce93b08b86121e3f485531d621ffd8667cbaf5 (patch) | |
| tree | 8a109df44a55ab2e472761ab30f33e886c96c328 /llvm | |
| parent | 856777d8c972ef9dbf61306931447b4421c11a9e (diff) | |
| download | bcm5719-llvm-73ce93b08b86121e3f485531d621ffd8667cbaf5.tar.gz bcm5719-llvm-73ce93b08b86121e3f485531d621ffd8667cbaf5.zip | |
AMDGPU: Set IntrReadMem on memtime intrinsics
llvm-svn: 320188
Diffstat (limited to 'llvm')
| -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; } |
