[PATCH] D120494: [NVPTX][AsmPrinter] Respect metadata 'align' for aggregate input parameters
Kristina Bessonova via Phabricator via llvm-commits
llvm-commits at lists.llvm.org
Thu Feb 24 08:43:44 PST 2022
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>
More information about the llvm-commits
mailing list