[NVPTX][AsmPrinter] Respect metadata 'align' for aggregate input parameters

[tra]

What is the problem the patch is intended to solve? To allow better vectorization of param-space loads/stores ?

llvm/test/CodeGen/NVPTX/align-annotation.ll

27

Does clang generate this metadata? AFAICT, it does not. https://godbolt.org/z/s9EYbdfjT
Where is this metadata expected to come from?

[krisb]

In D120494#3343880, @tra wrote:

What is the problem the patch is intended to solve? To allow better vectorization of param-space loads/stores ?

I guess it may help vectorization, but I do not have a particular use case.
My main motivation is to make this consistent: we already handle this metadata while emitting an aggregate as a return value, but ignore it for parameters, which is a bit unexpected if this is the same aggregate.

llvm/test/CodeGen/NVPTX/align-annotation.ll

27

Hmm, I looked at clang, and, surprisingly, it indeed doesn't generate align property, yet it generates all the other properties of nvvm.annotations metadata. We might add align to clang for the sake of completeness.
However, clang support does not matter much in this context (it could be different frontends or users could write IR manually as align is defined in NVVM IR specification). It appears that the backend already supports the metadata; so this patch does not introduce a new one, it just makes the existing a bit more consistent.

[tra]

llvm/test/CodeGen/NVPTX/align-annotation.ll

27

We might add align to clang for the sake of completeness.

Let me play a devil's advocate here.

If we lived without it till now and there's no demonstrable benefit of having it supported, I'd argue that we do not need it and we should remove it instead.

NVVM metadata is a bit of a hack, IMO. Some of the things it's used for should be done via the standard LLVM mechanisms. E.g. kernels should be using a special calling convention, types do not have the concept of alignment in LLVM, only memory accesses do. What are we supposed to do with NVVM alignment info if the user-generated code provides conflicting alignment info about the loads/stores of the data? IMO it creates more problems than it was supposed to solve.

it could be different frontends or users could write IR manually as align is defined in NVVM IR specification

NVVM IR specification is somewhat irrelevant here. LLVM != NVVM. We have already diverged quite a bit since the time NVIDIA has contributed NVPTX back-end.

llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

if (isKernelFunction(*F)) {
       }
     }
 
     if (!PAL.hasParamAttr(paramIndex, Attribute::ByVal)) {
       if (Ty->isAggregateType() || Ty->isVectorTy() || Ty->isIntegerTy(128)) {
         // Just print .param .align <a> .b8 .param[size];
         // <a> = PAL.getparamalignment
         // size = typeallocsize of element type
-        const Align align = DL.getValueOrABITypeAlignment(
+        unsigned alignVal = 0;
+        const Align align = getAlign(*F, paramIndex + 1, alignVal)
+                                ? Align(alignVal)
+                                : DL.getValueOrABITypeAlignment(
                                       PAL.getParamAlignment(paramIndex), Ty);
 
         unsigned sz = DL.getTypeAllocSize(Ty);
         O << "\t.param .align " << align.value() << " .b8 ";
         printParamName(I, paramIndex, O);
         O << "[" << sz << "]";
 
         continue;
       }

llvm/test/CodeGen/NVPTX/align-annotation.ll

+; RUN: llc < %s -march=nvptx | FileCheck %s
+
+; Compiled from:
+;
+; struct __align__(32) S1_t {
+;   float a;
+;   float b;
+; };
+;
+; __device__ __noinline__ S1_t dummy(S1_t in)
+; {
+;   in.a = in.a + 1;
+;   return in;
+; }
+
+%struct.S1_t = type { float, float }
+
+; CHECK: .visible .func  (.param .align 32 .b8 func_retval0[8]) _Z5dummy4S1_t(
+; CHECK-NEXT: .param .align 32 .b8 _Z5dummy4S1_t_param_0[8]
+
+define %struct.S1_t @_Z5dummy4S1_t(%struct.S1_t %in) {
+  ret %struct.S1_t %in
+}
+
+!nvvm.annotations = !{!5}
+
+!5 = !{%struct.S1_t (%struct.S1_t)* @_Z5dummy4S1_t, !"align", i32 32, !"align", i32 65568}