1// REQUIRES: amdgpu-registered-target
2// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
3
4// CHECK-NOT: %struct.single_element_struct_arg = type { i32 }
5typedef struct single_element_struct_arg
6{
7    int i;
8} single_element_struct_arg_t;
9
10// CHECK: %struct.struct_arg = type { i32, float, i32 }
11typedef struct struct_arg
12{
13    int i1;
14    float f;
15    int i2;
16} struct_arg_t;
17
18// CHECK: %struct.struct_of_arrays_arg = type { [2 x i32], float, [4 x i32], [3 x float], i32 }
19typedef struct struct_of_arrays_arg
20{
21    int i1[2];
22    float f1;
23    int i2[4];
24    float f2[3];
25    int i3;
26} struct_of_arrays_arg_t;
27
28// CHECK: %struct.struct_of_structs_arg = type { i32, float, %struct.struct_arg, i32 }
29typedef struct struct_of_structs_arg
30{
31    int i1;
32    float f1;
33    struct_arg_t s1;
34    int i2;
35} struct_of_structs_arg_t;
36
37// CHECK-LABEL: @test_single_element_struct_arg
38// CHECK: i32 %arg1.coerce
39__kernel void test_single_element_struct_arg(single_element_struct_arg_t arg1)
40{
41}
42
43// CHECK-LABEL: @test_struct_arg
44// CHECK: %struct.struct_arg %arg1.coerce
45__kernel void test_struct_arg(struct_arg_t arg1)
46{
47}
48
49// CHECK-LABEL: @test_struct_of_arrays_arg
50// CHECK: %struct.struct_of_arrays_arg %arg1.coerce
51__kernel void test_struct_of_arrays_arg(struct_of_arrays_arg_t arg1)
52{
53}
54
55// CHECK-LABEL: @test_struct_of_structs_arg
56// CHECK: %struct.struct_of_structs_arg %arg1.coerce
57__kernel void test_struct_of_structs_arg(struct_of_structs_arg_t arg1)
58{
59}
60
61// CHECK-LABEL: @test_non_kernel_struct_arg
62// CHECK-NOT: %struct.struct_arg %arg1.coerce
63// CHECK: %struct.struct_arg* byval
64void test_non_kernel_struct_arg(struct_arg_t arg1)
65{
66}
67