This is an archive of the discontinued LLVM Phabricator instance.

Differential D41028

AMDGPU: Set IntrReadMem on memtime intrinsics
ClosedPublic

Authored by arsenm on Dec 8 2017, 11:12 AM.

Details

Reviewers
rampitec

Diff Detail

Event Timeline

arsenm created this revision.Dec 8 2017, 11:12 AM
Herald added subscribers: t-tye, tpr, dstuttard and 4 others. · View Herald TranscriptDec 8 2017, 11:12 AM
rampitec accepted this revision.Dec 8 2017, 11:15 AM

LGTM

This revision is now accepted and ready to land.Dec 8 2017, 11:15 AM
arsenm closed this revision.Dec 8 2017, 12:02 PM

r320188

Revision Contents

PathSize
include/
llvm/
IR/
4 lines
lib/
Target/
AMDGPU/
7 lines

Diff 126185

include/llvm/IR/IntrinsicsAMDGPU.td

Show First 20 Lines • Show All 564 Lines • ▼ Show 20 Linesdef int_amdgcn_buffer_wbinvl1 :
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
def int_amdgcn_s_dcache_inv : def int_amdgcn_s_dcache_inv :
GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">, GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
def int_amdgcn_s_memtime : def int_amdgcn_s_memtime :
GCCBuiltin<"__builtin_amdgcn_s_memtime">, GCCBuiltin<"__builtin_amdgcn_s_memtime">,
Intrinsic<[llvm_i64_ty], [], []>; Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
def int_amdgcn_s_sleep : def int_amdgcn_s_sleep :
GCCBuiltin<"__builtin_amdgcn_s_sleep">, GCCBuiltin<"__builtin_amdgcn_s_sleep">,
Intrinsic<[], [llvm_i32_ty], []> { Intrinsic<[], [llvm_i32_ty], []> {
} }
def int_amdgcn_s_incperflevel : def int_amdgcn_s_incperflevel :
GCCBuiltin<"__builtin_amdgcn_s_incperflevel">, GCCBuiltin<"__builtin_amdgcn_s_incperflevel">,
▲ Show 20 Lines • Show All 229 Lines • ▼ Show 20 Linesdef int_amdgcn_s_dcache_wb :
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
def int_amdgcn_s_dcache_wb_vol : def int_amdgcn_s_dcache_wb_vol :
GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">, GCCBuiltin<"__builtin_amdgcn_s_dcache_wb_vol">,
Intrinsic<[], [], []>; Intrinsic<[], [], []>;
def int_amdgcn_s_memrealtime : def int_amdgcn_s_memrealtime :
GCCBuiltin<"__builtin_amdgcn_s_memrealtime">, GCCBuiltin<"__builtin_amdgcn_s_memrealtime">,
Intrinsic<[llvm_i64_ty], [], []>; Intrinsic<[llvm_i64_ty], [], [IntrReadMem]>;
// llvm.amdgcn.ds.permute <index> <src> // llvm.amdgcn.ds.permute <index> <src>
def int_amdgcn_ds_permute : def int_amdgcn_ds_permute :
GCCBuiltin<"__builtin_amdgcn_ds_permute">, GCCBuiltin<"__builtin_amdgcn_ds_permute">,
Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>; Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty], [IntrNoMem, IntrConvergent]>;
// llvm.amdgcn.ds.bpermute <index> <src> // llvm.amdgcn.ds.bpermute <index> <src>
def int_amdgcn_ds_bpermute : def int_amdgcn_ds_bpermute :
▲ Show 20 Lines • Show All 44 Lines • Show Last 20 Lines

lib/Target/AMDGPU/SMInstructions.td

Show First 20 Lines • Show All 123 Lines • ▼ Show 20 Lines def _SGPR : SM_Store_Pseudo <opName,
let PseudoInstr = opName # "_SGPR"; let PseudoInstr = opName # "_SGPR";
} }
} }
class SM_Time_Pseudo<string opName, SDPatternOperator node> : SM_Pseudo< class SM_Time_Pseudo<string opName, SDPatternOperator node> : SM_Pseudo<
opName, (outs SReg_64_XEXEC:$sdst), (ins), opName, (outs SReg_64_XEXEC:$sdst), (ins),
" $sdst", [(set i64:$sdst, (node))]> { " $sdst", [(set i64:$sdst, (node))]> {
let hasSideEffects = 1; let hasSideEffects = 1;
// FIXME: mayStore = ? is a workaround for tablegen bug for different let mayStore = 0;
// inferred mayStore flags for the instruction pattern vs. standalone let mayLoad = 1;
// Pat. Each considers the other contradictory.
let mayStore = ?;
let mayLoad = ?;
let has_sbase = 0; let has_sbase = 0;
let has_offset = 0; let has_offset = 0;
} }
class SM_Inval_Pseudo <string opName, SDPatternOperator node> : SM_Pseudo< class SM_Inval_Pseudo <string opName, SDPatternOperator node> : SM_Pseudo<
opName, (outs), (ins), "", [(node)]> { opName, (outs), (ins), "", [(node)]> {
let hasSideEffects = 1; let hasSideEffects = 1;
let mayStore = 1; let mayStore = 1;
▲ Show 20 Lines • Show All 381 Lines • Show Last 20 Lines