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

[Kristina Bessonova via Phabricator via llvm-commits]

krisb created this revision.
krisb added reviewers: tra, jholewinski.
Herald added subscribers: asavonic, hiraditya.
krisb requested review of this revision.
Herald added a project: LLVM.
Herald added a subscriber: llvm-commits.

Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D120494

Files:
  llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
  llvm/test/CodeGen/NVPTX/align-annotation.ll


Index: llvm/test/CodeGen/NVPTX/align-annotation.ll
===================================================================
--- /dev/null
+++ llvm/test/CodeGen/NVPTX/align-annotation.ll
@@ -0,0 +1,27 @@
+; 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}
Index: llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
===================================================================
--- llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -1407,8 +1407,11 @@
         // Just print .param .align <a> .b8 .param[size];
         // <a> = PAL.getparamalignment
         // size = typeallocsize of element type
-        const Align align = DL.getValueOrABITypeAlignment(
-            PAL.getParamAlignment(paramIndex), Ty);
+        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 ";


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D120494.411147.patch
Type: text/x-patch
Size: 1721 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20220224/9efc1bb9/attachment.bin>