[PATCH] D120129: [NVPTX] Enhance vectorization of ld.param & st.param

[Artem Belevich via Phabricator via llvm-commits]

tra accepted this revision.
tra added a comment.
This revision is now accepted and ready to land.

Few cosmetic&style nits. Otherwise LGTM.



================
Comment at: clang/test/CodeGenCUDA/device-fun-linkage.cu:23-24
+static __device__ void static_func() {}
+// NORDC-NOT: define{{.*}} void @_ZL13static_funcv()
+// RDC-NOT:   define{{.*}} void @_ZL13static_funcv()
+
----------------
Negative testing is tricky. The check here only ensures that the function does not appear between the positive matches around them. If we emit this function as the first or last function in the file, these checks will succeed, even though they should not.

One way to deal with this is to run negative tests separately, with their own RUN lines and check tags.


================
Comment at: llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp:1423
 
         unsigned sz = DL.getTypeAllocSize(Ty);
+        O << "\t.param .align " << OptimalAlign.value() << " .b8 ";
----------------
I'd inline it into printout, too.


================
Comment at: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp:1307
   for (unsigned i = 0, e = Args.size(); i != e; ++i, ++OIdx) {
+    const Function *F = CB.getFunction();
     Type *Ty = Args[i].Ty;
----------------
This could be moved out of the loop.


================
Comment at: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp:1321
+          ParamAlign = getFunctionParamOptimizedAlign(F, Ty, DL).value();
         unsigned sz = DL.getTypeAllocSize(Ty);
+        O << ".param .align " << ParamAlign << " .b8 ";
----------------
No need for the variable, just inline the call into printout statement. Makes it easier to see what we're printing.


================
Comment at: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp:1362
+    Align AlignCandidate = getFunctionParamOptimizedAlign(F, ETy, DL);
+    ParamByValAlign = std::max(ParamByValAlign, AlignCandidate);
+
----------------
```
Align ParamByValAlign = std::max(Outs[OIdx].Flags.getNonZeroByValAlign(), 
                                                           getFunctionParamOptimizedAlign(F, ETy, DL));
```

We could try inlining it into the printout, but it may be too unwieldy for that. Up to you.



================
Comment at: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp:1364
+
     unsigned sz = Outs[OIdx].Flags.getByValSize();
+    O << ".param .align " << ParamByValAlign.value() << " .b8 ";
----------------
ditto.


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D120129/new/

https://reviews.llvm.org/D120129