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