[llvm] [AMDGPU] Split struct kernel arguments (PR #133786)

[Shilei Tian via llvm-commits]

================
@@ -0,0 +1,120 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --check-attributes --check-globals all --version 5
+; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -passes=amdgpu-split-kernel-arguments -amdgpu-enable-split-kernel-args -verify-machineinstrs < %s | FileCheck %s
+; RUN: opt -S -mtriple=amdgcn-amd-amdhsa -passes=amdgpu-split-kernel-arguments -amdgpu-enable-split-kernel-args -verify-machineinstrs < %s | llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx1100 -verify-machineinstrs | FileCheck --check-prefix=GCN %s
+;
+; The LLVM IR is from the following HIP program:
+;
+; struct A {
+; int i;
+; char c;
+; long l;
+; int *p;
+; };
+
+; struct B {
+; char c;
+; A a1;
+; int i;
+; A a2;
+; };
+;
+; __global__ void test(int *out, int i, A a, char c, B b) {
+;  *out = i + a.l + c + a.l + b.a1.c;
+;  b.a2.p[2] = a.l + b.a2.c;
+;}
+;
+%struct.A = type { i32, i8, i64, ptr }
----------------
shiltian wrote:

Do we (compiler and/or runtime) have special handling of 64-bit byval? In OpenMP we pass all arguments as `void *`, which is 64-bit so if anything smaller than that, it can be passed as byval and then we do a type punning in the code. Do we have something similar in HIP?

https://github.com/llvm/llvm-project/pull/133786