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

[Yaxun Liu 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 }
----------------
yxsamliu wrote:

https://godbolt.org/z/3aGMEjvMK

If a struct size is not greater than 64 bit and has only one field, and in HIP source code it is passed by value, clang will change it to a scalar parameter in the kernel signature. In this case, it is a normal scalar argument in LLVM IR and this pass does not handle it.

Otherwise, clang use byref ptr to struct type for it in the kernel signature, and my pass will handle it.

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