[NVPTX] Annotate param loads/stores as mayLoad/mayStore.

[jlebar]

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.

[jlebar]

Friendly ping. We're looking at a correctness issue which appears that it might be related to this; would be nice to have this resolved.

[jholewinski]

Sorry, missed the original notifications.

The change looks reasonable. Do you have a particular example that shows the issue you're addressing here? The pre-call stores and post-call loads should be properly chained to prevent invalid reordering.

[jlebar]

In D17471#364891, @jholewinski wrote:

Sorry, missed the original notifications.

Do you have a particular example that shows the issue you're addressing here? The pre-call stores and post-call loads should be properly chained to prevent invalid reordering.

We're seeing some (non-param) loads suspiciously being sunk. This may be unrelated, but it's helpful to rule out...

[jholewinski]

I'm not sure I understand the relation to non-param loads/stores. Param loads/stores will never alias another address space, so arbitrary reordering with other loads/stores should be legal. Can you show an example of what you're seeing?

[jlebar]

In D17471#365176, @jholewinski wrote:

I'm not sure I understand the relation to non-param loads/stores. Param loads/stores will never alias another address space, so arbitrary reordering with other loads/stores should be legal. Can you show an example of what you're seeing?

I looked more closely at the thing we were looking at, and you're right, it's unrelated.

However, param loads/stores *can* alias with the global address space, yes? So given something like

a = ld.global [A]
st.param [B]
if (...) {
  // use a
}

we cannot safely sink the load of 'a' into the branch:

st.param [B]
if (...) {
  a = ld.global [A]
  // use a
}

because A and B may alias.

Or, even if that's not possible, then one could conceive of a situation where we somehow have a param pointer aliasing another param:

param i32 A;
param i32* B = &A;
// now we have the same problem.

Anyway I agree this is mostly academic, given the kind of code we (currently) generate.

[jholewinski]

At the hardware level, there is no 'param' address space. It's a PTX abstraction to handle parameter passing. There may be memory associated with it for complex calls, but it should never alias another address space. There's really no such thing as a pointer to a param space pointer.

Anyway, I have no issues with this change. These are technically loads/stores. Thanks!

lib/Target/NVPTX/NVPTXInstrInfo.td

SDNode<"NVPTXISD::StoreRetvalV4", SDTStoreRetvalV4Profile,
          [SDNPHasChain, SDNPSideEffect]>;
 def PseudoUseParam :
   SDNode<"NVPTXISD::PseudoUseParam", SDTPseudoUseParamProfile,
          [SDNPHasChain, SDNPOutGlue, SDNPInGlue, SDNPSideEffect]>;
 def RETURNNode :
   SDNode<"NVPTXISD::RETURN", SDTCallArgMarkProfile,
          [SDNPHasChain, SDNPSideEffect]>;
 
+let mayLoad = 1 in {
   class LoadParamMemInst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs regclass:$dst), (ins i32imm:$b),
                   !strconcat(!strconcat("ld.param", opstr),
                              "\t$dst, [retval0+$b];"),
                   []>;
 
-class LoadParamRegInst<NVPTXRegClass regclass, string opstr> :
-      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<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs regclass:$dst, regclass:$dst2), (ins i32imm:$b),
                   !strconcat("ld.param.v2", opstr,
                              "\t{{$dst, $dst2}}, [retval0+$b];"), []>;
 
   class LoadParamV4MemInst<NVPTXRegClass regclass, string opstr> :
         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<NVPTXRegClass regclass, string opstr> :
+      NVPTXInst<(outs regclass:$dst), (ins i32imm:$b),
+                !strconcat("mov", opstr, "\t$dst, retval$b;"),
+                [(set regclass:$dst, (LoadParam (i32 0), (i32 imm:$b)))]>;
+
+let mayStore = 1 in {
   class StoreParamInst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs), (ins regclass:$val, i32imm:$a, i32imm:$b),
                   !strconcat("st.param", opstr, "\t[param$a+$b], $val;"),
                   []>;
 
   class StoreParamV2Inst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs), (ins regclass:$val, regclass:$val2,
                                i32imm:$a, i32imm:$b),
                   !strconcat("st.param.v2", opstr,
                              "\t[param$a+$b], {{$val, $val2}};"),
                   []>;
 
   class StoreParamV4Inst<NVPTXRegClass regclass, string opstr> :
         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<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs), (ins regclass:$val, i32imm:$a),
                   !strconcat("st.param", opstr, "\t[func_retval0+$a], $val;"),
                   []>;
 
   class StoreRetvalV2Inst<NVPTXRegClass regclass, string opstr> :
         NVPTXInst<(outs), (ins regclass:$val, regclass:$val2, i32imm:$a),
                   !strconcat("st.param.v2", opstr,
                              "\t[func_retval0+$a], {{$val, $val2}};"),
                   []>;
 
   class StoreRetvalV4Inst<NVPTXRegClass regclass, string opstr> :
         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<string OpcStr, SDNode OpNode> {
      def PrintCallNoRetInst : NVPTXInst<(outs), (ins),
        !strconcat(OpcStr, " "), [(OpNode (i32 0))]>;
      def PrintCallRetInst1 : NVPTXInst<(outs), (ins),
        !strconcat(OpcStr, " (retval0), "), [(OpNode (i32 1))]>;
      def PrintCallRetInst2 : NVPTXInst<(outs), (ins),

lib/Target/NVPTX/NVPTXIntrinsics.td

 def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
   int_nvvm_read_ptx_sreg_warpsize>;
 
 
 //-----------------------------------
 // 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<string TyStr, NVPTXRegClass regclass> {
   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
                !strconcat("ldu.global.", TyStr),
                       []>, Requires<[hasLDU]>;
   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
                !strconcat("ldu.global.", TyStr),
                         []>, Requires<[hasLDU]>;
  def avar:  NVPTXInst<(outs regclass:$result), (ins imemAny:$src),
defm INT_PTX_LDU_G_v4f32_ELE
   : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
     Float32Regs>;
 
 
 //-----------------------------------
 // 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<string TyStr, NVPTXRegClass regclass> {
   def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
                !strconcat("ld.global.nc.", TyStr),
                       []>, Requires<[hasLDG]>;
   def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
                !strconcat("ld.global.nc.", TyStr),
                         []>, Requires<[hasLDG]>;
  def avar:  NVPTXInst<(outs regclass:$result), (ins imemAny:$src),