From 47e3ade685cbc66c55bcf58923ef37ef4b5b49f7 Mon Sep 17 00:00:00 2001 From: Justin Lebar Date: Tue, 1 Mar 2016 19:44:22 +0000 Subject: [PATCH] [NVPTX] Annotate param loads/stores as mayLoad/mayStore. Summary: Tablegen was unable to determine that param loads/stores were actually reading or writing from memory. I think this isn't a problem in practice for param stores, because those occur in a block right before we make our call. But param loads don't have to at the very beginning of a function, so should be annotated as mayLoad so we don't incorrectly optimize them. Reviewers: jholewinski Subscribers: jholewinski, llvm-commits Differential Revision: http://reviews.llvm.org/D17471 git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@262381 91177308-0d34-0410-b5e6-96231b3b80d8 --- lib/Target/NVPTX/NVPTXInstrInfo.td | 116 +++++++++++++++++++----------------- lib/Target/NVPTX/NVPTXIntrinsics.td | 8 +++ 2 files changed, 68 insertions(+), 56 deletions(-) diff --git a/lib/Target/NVPTX/NVPTXInstrInfo.td b/lib/Target/NVPTX/NVPTXInstrInfo.td index 4f4e23c3449..6a2c3282868 100644 --- a/lib/Target/NVPTX/NVPTXInstrInfo.td +++ b/lib/Target/NVPTX/NVPTXInstrInfo.td @@ -1764,68 +1764,72 @@ def RETURNNode : SDNode<"NVPTXISD::RETURN", SDTCallArgMarkProfile, [SDNPHasChain, SDNPSideEffect]>; -class LoadParamMemInst : - NVPTXInst<(outs regclass:$dst), (ins i32imm:$b), - !strconcat(!strconcat("ld.param", opstr), - "\t$dst, [retval0+$b];"), - []>; +let mayLoad = 1 in { + class LoadParamMemInst : + NVPTXInst<(outs regclass:$dst), (ins i32imm:$b), + !strconcat(!strconcat("ld.param", opstr), + "\t$dst, [retval0+$b];"), + []>; + + class LoadParamV2MemInst : + NVPTXInst<(outs regclass:$dst, regclass:$dst2), (ins i32imm:$b), + !strconcat("ld.param.v2", opstr, + "\t{{$dst, $dst2}}, [retval0+$b];"), []>; + + class LoadParamV4MemInst : + NVPTXInst<(outs regclass:$dst, regclass:$dst2, regclass:$dst3, + regclass:$dst4), + (ins i32imm:$b), + !strconcat("ld.param.v4", opstr, + "\t{{$dst, $dst2, $dst3, $dst4}}, [retval0+$b];"), + []>; +} class LoadParamRegInst : NVPTXInst<(outs regclass:$dst), (ins i32imm:$b), !strconcat("mov", opstr, "\t$dst, retval$b;"), [(set regclass:$dst, (LoadParam (i32 0), (i32 imm:$b)))]>; -class LoadParamV2MemInst : - NVPTXInst<(outs regclass:$dst, regclass:$dst2), (ins i32imm:$b), - !strconcat("ld.param.v2", opstr, - "\t{{$dst, $dst2}}, [retval0+$b];"), []>; - -class LoadParamV4MemInst : - NVPTXInst<(outs regclass:$dst, regclass:$dst2, regclass:$dst3, - regclass:$dst4), - (ins i32imm:$b), - !strconcat("ld.param.v4", opstr, - "\t{{$dst, $dst2, $dst3, $dst4}}, [retval0+$b];"), - []>; - -class StoreParamInst : - NVPTXInst<(outs), (ins regclass:$val, i32imm:$a, i32imm:$b), - !strconcat("st.param", opstr, "\t[param$a+$b], $val;"), - []>; - -class StoreParamV2Inst : - NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, - i32imm:$a, i32imm:$b), - !strconcat("st.param.v2", opstr, - "\t[param$a+$b], {{$val, $val2}};"), - []>; - -class StoreParamV4Inst : - NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, regclass:$val3, - regclass:$val4, i32imm:$a, - i32imm:$b), - !strconcat("st.param.v4", opstr, - "\t[param$a+$b], {{$val, $val2, $val3, $val4}};"), - []>; - -class StoreRetvalInst : - NVPTXInst<(outs), (ins regclass:$val, i32imm:$a), - !strconcat("st.param", opstr, "\t[func_retval0+$a], $val;"), - []>; - -class StoreRetvalV2Inst : - NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, i32imm:$a), - !strconcat("st.param.v2", opstr, - "\t[func_retval0+$a], {{$val, $val2}};"), - []>; - -class StoreRetvalV4Inst : - NVPTXInst<(outs), - (ins regclass:$val, regclass:$val2, regclass:$val3, - regclass:$val4, i32imm:$a), - !strconcat("st.param.v4", opstr, - "\t[func_retval0+$a], {{$val, $val2, $val3, $val4}};"), - []>; +let mayStore = 1 in { + class StoreParamInst : + NVPTXInst<(outs), (ins regclass:$val, i32imm:$a, i32imm:$b), + !strconcat("st.param", opstr, "\t[param$a+$b], $val;"), + []>; + + class StoreParamV2Inst : + NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, + i32imm:$a, i32imm:$b), + !strconcat("st.param.v2", opstr, + "\t[param$a+$b], {{$val, $val2}};"), + []>; + + class StoreParamV4Inst : + NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, regclass:$val3, + regclass:$val4, i32imm:$a, + i32imm:$b), + !strconcat("st.param.v4", opstr, + "\t[param$a+$b], {{$val, $val2, $val3, $val4}};"), + []>; + + class StoreRetvalInst : + NVPTXInst<(outs), (ins regclass:$val, i32imm:$a), + !strconcat("st.param", opstr, "\t[func_retval0+$a], $val;"), + []>; + + class StoreRetvalV2Inst : + NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, i32imm:$a), + !strconcat("st.param.v2", opstr, + "\t[func_retval0+$a], {{$val, $val2}};"), + []>; + + class StoreRetvalV4Inst : + NVPTXInst<(outs), + (ins regclass:$val, regclass:$val2, regclass:$val3, + regclass:$val4, i32imm:$a), + !strconcat("st.param.v4", opstr, + "\t[func_retval0+$a], {{$val, $val2, $val3, $val4}};"), + []>; +} let isCall=1 in { multiclass CALL { diff --git a/lib/Target/NVPTX/NVPTXIntrinsics.td b/lib/Target/NVPTX/NVPTXIntrinsics.td index 208731f4738..1aec2619ec6 100644 --- a/lib/Target/NVPTX/NVPTXIntrinsics.td +++ b/lib/Target/NVPTX/NVPTXIntrinsics.td @@ -1381,7 +1381,11 @@ def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs, // Support for ldu on sm_20 or later //----------------------------------- +// Don't annotate ldu instructions as mayLoad, as they load from memory that is +// read-only in a kernel. + // Scalar + multiclass LDU_G { def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src), !strconcat("ldu.global.", TyStr), @@ -1477,6 +1481,10 @@ defm INT_PTX_LDU_G_v4f32_ELE // Support for ldg on sm_35 or later //----------------------------------- +// Don't annotate ld.global.nc as mayLoad, because these loads go through the +// non-coherent texture cache, and therefore the values read must be read-only +// during the lifetime of the kernel. + multiclass LDG_G { def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src), !strconcat("ld.global.nc.", TyStr), -- 2.11.0