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

[Daniil Kovalev via Phabricator via llvm-commits]

kovdan01 marked 6 inline comments as done.
kovdan01 added inline comments.


================
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()
+
----------------
tra wrote:
> 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.
Used the approach with separate RUN lines, thanks!


================
Comment at: llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp:1362
+    Align AlignCandidate = getFunctionParamOptimizedAlign(F, ETy, DL);
+    ParamByValAlign = std::max(ParamByValAlign, AlignCandidate);
+
----------------
tra wrote:
> ```
> 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.
> 
Looks too unwieldy for inlining IMHO, so kept in current state.


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

https://reviews.llvm.org/D120129