clang/test/CodeGenOpenCL/amdgpu-abi-struct-coerce.cl

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
   3 // RUN: %clang_cc1 -no-opaque-pointers -triple r600-unknown-unknown -S -emit-llvm -o - %s | FileCheck %s
   4
   5 typedef __attribute__(( ext_vector_type(2) )) char char2;
   6 typedef __attribute__(( ext_vector_type(3) )) char char3;
   7 typedef __attribute__(( ext_vector_type(4) )) char char4;
   8
   9 typedef __attribute__(( ext_vector_type(2) )) short short2;
  10 typedef __attribute__(( ext_vector_type(3) )) short short3;
  11 typedef __attribute__(( ext_vector_type(4) )) short short4;
  12
  13 typedef __attribute__(( ext_vector_type(2) )) int int2;
  14 typedef __attribute__(( ext_vector_type(3) )) int int3;
  15 typedef __attribute__(( ext_vector_type(4) )) int int4;
  16 typedef __attribute__(( ext_vector_type(16) )) int int16;
  17 typedef __attribute__(( ext_vector_type(32) )) int int32;
  18
  19 // CHECK: %struct.empty_struct = type {}
  20 typedef struct empty_struct
  21 {
  22 } empty_struct;
  23
  24 // CHECK-NOT: %struct.single_element_struct_arg
  25 typedef struct single_element_struct_arg
  26 {
  27     int i;
  28 } single_element_struct_arg_t;
  29
  30 // CHECK-NOT: %struct.nested_single_element_struct_arg
  31 typedef struct nested_single_element_struct_arg
  32 {
  33   single_element_struct_arg_t i;
  34 } nested_single_element_struct_arg_t;
  35
  36 // CHECK: %struct.struct_arg = type { i32, float, i32 }
  37 typedef struct struct_arg
  38 {
  39     int i1;
  40     float f;
  41     int i2;
  42 } struct_arg_t;
  43
  44 // CHECK: %struct.struct_padding_arg = type { i8, i64 }
  45 typedef struct struct_padding_arg
  46 {
  47   char i1;
  48   long f;
  49 } struct_padding_arg;
  50
  51 // CHECK: %struct.struct_of_arrays_arg = type { [2 x i32], float, [4 x i32], [3 x float], i32 }
  52 typedef struct struct_of_arrays_arg
  53 {
  54     int i1[2];
  55     float f1;
  56     int i2[4];
  57     float f2[3];
  58     int i3;
  59 } struct_of_arrays_arg_t;
  60
  61 // CHECK: %struct.struct_of_structs_arg = type { i32, float, %struct.struct_arg, i32 }
  62 typedef struct struct_of_structs_arg
  63 {
  64     int i1;
  65     float f1;
  66     struct_arg_t s1;
  67     int i2;
  68 } struct_of_structs_arg_t;
  69
  70 typedef union
  71 {
  72   int b1;
  73   float b2;
  74 } transparent_u __attribute__((__transparent_union__));
  75
  76 // CHECK: %struct.single_array_element_struct_arg = type { [4 x i32] }
  77 typedef struct single_array_element_struct_arg
  78 {
  79     int i[4];
  80 } single_array_element_struct_arg_t;
  81
  82 // CHECK: %struct.single_struct_element_struct_arg = type { %struct.inner }
  83 // CHECK: %struct.inner = type { i32, i64 }
  84 typedef struct single_struct_element_struct_arg
  85 {
  86   struct inner {
  87     int a;
  88     long b;
  89   } s;
  90 } single_struct_element_struct_arg_t;
  91
  92 // CHECK: %struct.different_size_type_pair
  93 typedef struct different_size_type_pair {
  94   long l;
  95   int i;
  96 } different_size_type_pair;
  97
  98 // CHECK: %struct.flexible_array = type { i32, [0 x i32] }
  99 typedef struct flexible_array
 100 {
 101   int i;
 102   int flexible[];
 103 } flexible_array;
 104
 105 // CHECK: %struct.struct_arr16 = type { [16 x i32] }
 106 typedef struct struct_arr16
 107 {
 108     int arr[16];
 109 } struct_arr16;
 110
 111 // CHECK: %struct.struct_arr32 = type { [32 x i32] }
 112 typedef struct struct_arr32
 113 {
 114     int arr[32];
 115 } struct_arr32;
 116
 117 // CHECK: %struct.struct_arr33 = type { [33 x i32] }
 118 typedef struct struct_arr33
 119 {
 120     int arr[33];
 121 } struct_arr33;
 122
 123 // CHECK: %struct.struct_char_arr32 = type { [32 x i8] }
 124 typedef struct struct_char_arr32
 125 {
 126   char arr[32];
 127 } struct_char_arr32;
 128
 129 // CHECK-NOT: %struct.struct_char_x8
 130 typedef struct struct_char_x8 {
 131   char x, y, z, w;
 132   char a, b, c, d;
 133 } struct_char_x8;
 134
 135 // CHECK-NOT: %struct.struct_char_x4
 136 typedef struct struct_char_x4 {
 137   char x, y, z, w;
 138 } struct_char_x4;
 139
 140 // CHECK-NOT: %struct.struct_char_x3
 141 typedef struct struct_char_x3 {
 142   char x, y, z;
 143 } struct_char_x3;
 144
 145 // CHECK-NOT: %struct.struct_char_x2
 146 typedef struct struct_char_x2 {
 147   char x, y;
 148 } struct_char_x2;
 149
 150 // CHECK-NOT: %struct.struct_char_x1
 151 typedef struct struct_char_x1 {
 152   char x;
 153 } struct_char_x1;
 154
 155 // 4 registers from fields, 5 if padding included.
 156 // CHECK: %struct.nested = type { i8, i64 }
 157 // CHECK: %struct.num_regs_nested_struct = type { i32, %struct.nested }
 158 typedef struct num_regs_nested_struct {
 159   int x;
 160   struct nested {
 161     char z;
 162     long y;
 163   } inner;
 164 } num_regs_nested_struct;
 165
 166 // CHECK: %struct.double_nested = type { %struct.inner_inner }
 167 // CHECK: %struct.inner_inner = type { i8, i32, i8 }
 168 // CHECK: %struct.double_nested_struct = type { i32, %struct.double_nested, i16 }
 169 typedef struct double_nested_struct {
 170   int x;
 171   struct double_nested {
 172     struct inner_inner {
 173       char y;
 174       int q;
 175       char z;
 176     } inner_inner;
 177   } inner;
 178
 179   short w;
 180 } double_nested_struct;
 181
 182 // This is a large struct, but uses fewer registers than the limit.
 183 // CHECK: %struct.large_struct_padding = type { i8, i32, i8, i32, i8, i8, i16, i16, [3 x i8], i64, i32, i8, i32, i16, i8 }
 184 typedef struct large_struct_padding {
 185   char e0;
 186   int e1;
 187   char e2;
 188   int e3;
 189   char e4;
 190   char e5;
 191   short e6;
 192   short e7;
 193   char e8[3];
 194   long e9;
 195   int e10;
 196   char e11;
 197   int e12;
 198   short e13;
 199   char e14;
 200 } large_struct_padding;
 201
 202 // CHECK: %struct.int3_pair = type { <3 x i32>, <3 x i32> }
 203 // The number of registers computed should be 6, not 8.
 204 typedef struct int3_pair {
 205         int3 dx;
 206         int3 dy;
 207 } int3_pair;
 208
 209 // CHECK: %struct.struct_4regs = type { i32, i32, i32, i32 }
 210 typedef struct struct_4regs
 211 {
 212   int x;
 213   int y;
 214   int z;
 215   int w;
 216 } struct_4regs;
 217
 218 // CHECK: void @kernel_empty_struct_arg(%struct.empty_struct %s.coerce)
 219 __kernel void kernel_empty_struct_arg(empty_struct s) { }
 220
 221 // CHECK: void @kernel_single_element_struct_arg(i32 %arg1.coerce)
 222 __kernel void kernel_single_element_struct_arg(single_element_struct_arg_t arg1) { }
 223
 224 // CHECK: void @kernel_nested_single_element_struct_arg(i32 %arg1.coerce)
 225 __kernel void kernel_nested_single_element_struct_arg(nested_single_element_struct_arg_t arg1) { }
 226
 227 // CHECK: void @kernel_struct_arg(%struct.struct_arg %arg1.coerce)
 228 __kernel void kernel_struct_arg(struct_arg_t arg1) { }
 229
 230 // CHECK: void @kernel_struct_padding_arg(%struct.struct_padding_arg %arg1.coerce)
 231 __kernel void kernel_struct_padding_arg(struct_padding_arg arg1) { }
 232
 233 // CHECK: void @kernel_test_struct_of_arrays_arg(%struct.struct_of_arrays_arg %arg1.coerce)
 234 __kernel void kernel_test_struct_of_arrays_arg(struct_of_arrays_arg_t arg1) { }
 235
 236 // CHECK: void @kernel_struct_of_structs_arg(%struct.struct_of_structs_arg %arg1.coerce)
 237 __kernel void kernel_struct_of_structs_arg(struct_of_structs_arg_t arg1) { }
 238
 239 // CHECK: void @test_kernel_transparent_union_arg(i32 %u.coerce)
 240 __kernel void test_kernel_transparent_union_arg(transparent_u u) { }
 241
 242 // CHECK: void @kernel_single_array_element_struct_arg(%struct.single_array_element_struct_arg %arg1.coerce)
 243 __kernel void kernel_single_array_element_struct_arg(single_array_element_struct_arg_t arg1) { }
 244
 245 // CHECK: void @kernel_single_struct_element_struct_arg(%struct.single_struct_element_struct_arg %arg1.coerce)
 246 __kernel void kernel_single_struct_element_struct_arg(single_struct_element_struct_arg_t arg1) { }
 247
 248 // CHECK: void @kernel_different_size_type_pair_arg(%struct.different_size_type_pair %arg1.coerce)
 249 __kernel void kernel_different_size_type_pair_arg(different_size_type_pair arg1) { }
 250
 251 // CHECK: define{{.*}} void @func_f32_arg(float noundef %arg)
 252 void func_f32_arg(float arg) { }
 253
 254 // CHECK: define{{.*}} void @func_v2i16_arg(<2 x i16> noundef %arg)
 255 void func_v2i16_arg(short2 arg) { }
 256
 257 // CHECK: define{{.*}} void @func_v3i32_arg(<3 x i32> noundef %arg)
 258 void func_v3i32_arg(int3 arg) { }
 259
 260 // CHECK: define{{.*}} void @func_v4i32_arg(<4 x i32> noundef %arg)
 261 void func_v4i32_arg(int4 arg) { }
 262
 263 // CHECK: define{{.*}} void @func_v16i32_arg(<16 x i32> noundef %arg)
 264 void func_v16i32_arg(int16 arg) { }
 265
 266 // CHECK: define{{.*}} void @func_v32i32_arg(<32 x i32> noundef %arg)
 267 void func_v32i32_arg(int32 arg) { }
 268
 269 // CHECK: define{{.*}} void @func_empty_struct_arg()
 270 void func_empty_struct_arg(empty_struct empty) { }
 271
 272 // CHECK: void @func_single_element_struct_arg(i32 %arg1.coerce)
 273 void func_single_element_struct_arg(single_element_struct_arg_t arg1) { }
 274
 275 // CHECK: void @func_nested_single_element_struct_arg(i32 %arg1.coerce)
 276 void func_nested_single_element_struct_arg(nested_single_element_struct_arg_t arg1) { }
 277
 278 // CHECK: void @func_struct_arg(i32 %arg1.coerce0, float %arg1.coerce1, i32 %arg1.coerce2)
 279 void func_struct_arg(struct_arg_t arg1) { }
 280
 281 // CHECK: void @func_struct_padding_arg(i8 %arg1.coerce0, i64 %arg1.coerce1)
 282 void func_struct_padding_arg(struct_padding_arg arg1) { }
 283
 284 // CHECK: define{{.*}} void @func_struct_char_x8([2 x i32] %arg.coerce)
 285 void func_struct_char_x8(struct_char_x8 arg) { }
 286
 287 // CHECK: define{{.*}} void @func_struct_char_x4(i32 %arg.coerce)
 288 void func_struct_char_x4(struct_char_x4 arg) { }
 289
 290 // CHECK: define{{.*}} void @func_struct_char_x3(i32 %arg.coerce)
 291 void func_struct_char_x3(struct_char_x3 arg) { }
 292
 293 // CHECK: define{{.*}} void @func_struct_char_x2(i16 %arg.coerce)
 294 void func_struct_char_x2(struct_char_x2 arg) { }
 295
 296 // CHECK: define{{.*}} void @func_struct_char_x1(i8 %arg.coerce)
 297 void func_struct_char_x1(struct_char_x1 arg) { }
 298
 299 // CHECK: void @func_transparent_union_arg(i32 %u.coerce)
 300 void func_transparent_union_arg(transparent_u u) { }
 301
 302 // CHECK: void @func_single_array_element_struct_arg([4 x i32] %arg1.coerce)
 303 void func_single_array_element_struct_arg(single_array_element_struct_arg_t arg1) { }
 304
 305 // CHECK: void @func_single_struct_element_struct_arg(%struct.inner %arg1.coerce)
 306 void func_single_struct_element_struct_arg(single_struct_element_struct_arg_t arg1) { }
 307
 308 // CHECK: void @func_different_size_type_pair_arg(i64 %arg1.coerce0, i32 %arg1.coerce1)
 309 void func_different_size_type_pair_arg(different_size_type_pair arg1) { }
 310
 311 // CHECK: void @func_flexible_array_arg(%struct.flexible_array addrspace(5)* nocapture noundef byval(%struct.flexible_array) align 4 %arg)
 312 void func_flexible_array_arg(flexible_array arg) { }
 313
 314 // CHECK: define{{.*}} float @func_f32_ret()
 315 float func_f32_ret()
 316 {
 317   return 0.0f;
 318 }
 319
 320 // CHECK: define{{.*}} void @func_empty_struct_ret()
 321 empty_struct func_empty_struct_ret()
 322 {
 323   empty_struct s = {};
 324   return s;
 325 }
 326
 327 // CHECK: define{{.*}} i32 @single_element_struct_ret()
 328 // CHECK: ret i32 0
 329 single_element_struct_arg_t single_element_struct_ret()
 330 {
 331   single_element_struct_arg_t s = { 0 };
 332   return s;
 333 }
 334
 335 // CHECK: define{{.*}} i32 @nested_single_element_struct_ret()
 336 // CHECK: ret i32 0
 337 nested_single_element_struct_arg_t nested_single_element_struct_ret()
 338 {
 339   nested_single_element_struct_arg_t s = { 0 };
 340   return s;
 341 }
 342
 343 // CHECK: define{{.*}} %struct.struct_arg @func_struct_ret()
 344 // CHECK: ret %struct.struct_arg zeroinitializer
 345 struct_arg_t func_struct_ret()
 346 {
 347   struct_arg_t s = { 0 };
 348   return s;
 349 }
 350
 351 // CHECK: define{{.*}} %struct.struct_padding_arg @func_struct_padding_ret()
 352 // CHECK: ret %struct.struct_padding_arg zeroinitializer
 353 struct_padding_arg func_struct_padding_ret()
 354 {
 355   struct_padding_arg s = { 0 };
 356   return s;
 357 }
 358
 359 // CHECK: define{{.*}} [2 x i32] @func_struct_char_x8_ret()
 360 // CHECK: ret [2 x i32] zeroinitializer
 361 struct_char_x8 func_struct_char_x8_ret()
 362 {
 363   struct_char_x8 s = { 0 };
 364   return s;
 365 }
 366
 367 // CHECK: define{{.*}} i32 @func_struct_char_x4_ret()
 368 // CHECK: ret i32 0
 369 struct_char_x4 func_struct_char_x4_ret()
 370 {
 371   struct_char_x4 s = { 0 };
 372   return s;
 373 }
 374
 375 // CHECK: define{{.*}} i32 @func_struct_char_x3_ret()
 376 // CHECK: ret i32 0
 377 struct_char_x3 func_struct_char_x3_ret()
 378 {
 379   struct_char_x3 s = { 0 };
 380   return s;
 381 }
 382
 383 // CHECK: define{{.*}} i16 @func_struct_char_x2_ret()
 384 struct_char_x2 func_struct_char_x2_ret()
 385 {
 386   struct_char_x2 s = { 0 };
 387   return s;
 388 }
 389
 390 // CHECK: define{{.*}} i8 @func_struct_char_x1_ret()
 391 // CHECK: ret i8 0
 392 struct_char_x1 func_struct_char_x1_ret()
 393 {
 394   struct_char_x1 s = { 0 };
 395   return s;
 396 }
 397
 398 // CHECK: define{{.*}} %struct.struct_arr16 @func_ret_struct_arr16()
 399 // CHECK: ret %struct.struct_arr16 zeroinitializer
 400 struct_arr16 func_ret_struct_arr16()
 401 {
 402   struct_arr16 s = { 0 };
 403   return s;
 404 }
 405
 406 // CHECK: define{{.*}} void @func_ret_struct_arr32(%struct.struct_arr32 addrspace(5)* noalias nocapture writeonly sret(%struct.struct_arr32) align 4 %agg.result)
 407 struct_arr32 func_ret_struct_arr32()
 408 {
 409   struct_arr32 s = { 0 };
 410   return s;
 411 }
 412
 413 // CHECK: define{{.*}} void @func_ret_struct_arr33(%struct.struct_arr33 addrspace(5)* noalias nocapture writeonly sret(%struct.struct_arr33) align 4 %agg.result)
 414 struct_arr33 func_ret_struct_arr33()
 415 {
 416   struct_arr33 s = { 0 };
 417   return s;
 418 }
 419
 420 // CHECK: define{{.*}} %struct.struct_char_arr32 @func_ret_struct_char_arr32()
 421 struct_char_arr32 func_ret_struct_char_arr32()
 422 {
 423   struct_char_arr32 s = { 0 };
 424   return s;
 425 }
 426
 427 // CHECK: define{{.*}} i32 @func_transparent_union_ret() local_unnamed_addr #1 {
 428 // CHECK: ret i32 0
 429 transparent_u func_transparent_union_ret()
 430 {
 431   transparent_u u = { 0 };
 432   return u;
 433 }
 434
 435 // CHECK: define{{.*}} %struct.different_size_type_pair @func_different_size_type_pair_ret()
 436 different_size_type_pair func_different_size_type_pair_ret()
 437 {
 438   different_size_type_pair s = { 0 };
 439   return s;
 440 }
 441
 442 // CHECK: define{{.*}} void @func_flexible_array_ret(%struct.flexible_array addrspace(5)* noalias nocapture writeonly sret(%struct.flexible_array) align 4 %agg.result)
 443 flexible_array func_flexible_array_ret()
 444 {
 445   flexible_array s = { 0 };
 446   return s;
 447 }
 448
 449 // CHECK: define{{.*}} void @func_reg_state_lo(<4 x i32> noundef %arg0, <4 x i32> noundef %arg1, <4 x i32> noundef %arg2, i32 noundef %arg3, i32 %s.coerce0, float %s.coerce1, i32 %s.coerce2)
 450 void func_reg_state_lo(int4 arg0, int4 arg1, int4 arg2, int arg3, struct_arg_t s) { }
 451
 452 // CHECK: define{{.*}} void @func_reg_state_hi(<4 x i32> noundef %arg0, <4 x i32> noundef %arg1, <4 x i32> noundef %arg2, i32 noundef %arg3, i32 noundef %arg4, %struct.struct_arg addrspace(5)* nocapture noundef byval(%struct.struct_arg) align 4 %s)
 453 void func_reg_state_hi(int4 arg0, int4 arg1, int4 arg2, int arg3, int arg4, struct_arg_t s) { }
 454
 455 // XXX - Why don't the inner structs flatten?
 456 // CHECK: define{{.*}} void @func_reg_state_num_regs_nested_struct(<4 x i32> noundef %arg0, i32 noundef %arg1, i32 %arg2.coerce0, %struct.nested %arg2.coerce1, i32 %arg3.coerce0, %struct.nested %arg3.coerce1, %struct.num_regs_nested_struct addrspace(5)* nocapture noundef byval(%struct.num_regs_nested_struct) align 8 %arg4)
 457 void func_reg_state_num_regs_nested_struct(int4 arg0, int arg1, num_regs_nested_struct arg2, num_regs_nested_struct arg3, num_regs_nested_struct arg4) { }
 458
 459 // CHECK: define{{.*}} void @func_double_nested_struct_arg(<4 x i32> noundef %arg0, i32 noundef %arg1, i32 %arg2.coerce0, %struct.double_nested %arg2.coerce1, i16 %arg2.coerce2)
 460 void func_double_nested_struct_arg(int4 arg0, int arg1, double_nested_struct arg2) { }
 461
 462 // CHECK: define{{.*}} %struct.double_nested_struct @func_double_nested_struct_ret(<4 x i32> noundef %arg0, i32 noundef %arg1)
 463 double_nested_struct func_double_nested_struct_ret(int4 arg0, int arg1) {
 464   double_nested_struct s = { 0 };
 465   return s;
 466 }
 467
 468 // CHECK: define{{.*}} void @func_large_struct_padding_arg_direct(i8 %arg.coerce0, i32 %arg.coerce1, i8 %arg.coerce2, i32 %arg.coerce3, i8 %arg.coerce4, i8 %arg.coerce5, i16 %arg.coerce6, i16 %arg.coerce7, [3 x i8] %arg.coerce8, i64 %arg.coerce9, i32 %arg.coerce10, i8 %arg.coerce11, i32 %arg.coerce12, i16 %arg.coerce13, i8 %arg.coerce14)
 469 void func_large_struct_padding_arg_direct(large_struct_padding arg) { }
 470
 471 // CHECK: define{{.*}} void @func_large_struct_padding_arg_store(%struct.large_struct_padding addrspace(1)* nocapture noundef writeonly %out, %struct.large_struct_padding addrspace(5)* nocapture noundef readonly byval(%struct.large_struct_padding) align 8 %arg)
 472 void func_large_struct_padding_arg_store(global large_struct_padding* out, large_struct_padding arg) {
 473   *out = arg;
 474 }
 475
 476 // CHECK: define{{.*}} void @v3i32_reg_count(<3 x i32> noundef %arg1, <3 x i32> noundef %arg2, <3 x i32> noundef %arg3, <3 x i32> noundef %arg4, i32 %arg5.coerce0, float %arg5.coerce1, i32 %arg5.coerce2)
 477 void v3i32_reg_count(int3 arg1, int3 arg2, int3 arg3, int3 arg4, struct_arg_t arg5) { }
 478
 479 // Function signature from blender, nothing should be passed byval. The v3i32
 480 // should not count as 4 passed registers.
 481 // CHECK: define{{.*}} void @v3i32_pair_reg_count(%struct.int3_pair addrspace(5)* nocapture noundef %arg0, <3 x i32> %arg1.coerce0, <3 x i32> %arg1.coerce1, <3 x i32> noundef %arg2, <3 x i32> %arg3.coerce0, <3 x i32> %arg3.coerce1, <3 x i32> noundef %arg4, float noundef %arg5)
 482 void v3i32_pair_reg_count(int3_pair *arg0, int3_pair arg1, int3 arg2, int3_pair arg3, int3 arg4, float arg5) { }
 483
 484 // Each short4 should fit pack into 2 registers.
 485 // CHECK: define{{.*}} void @v4i16_reg_count(<4 x i16> noundef %arg0, <4 x i16> noundef %arg1, <4 x i16> noundef %arg2, <4 x i16> noundef %arg3, <4 x i16> noundef %arg4, <4 x i16> noundef %arg5, i32 %arg6.coerce0, i32 %arg6.coerce1, i32 %arg6.coerce2, i32 %arg6.coerce3)
 486 void v4i16_reg_count(short4 arg0, short4 arg1, short4 arg2, short4 arg3,
 487                      short4 arg4, short4 arg5, struct_4regs arg6) { }
 488
 489 // CHECK: define{{.*}} void @v4i16_pair_reg_count_over(<4 x i16> noundef %arg0, <4 x i16> noundef %arg1, <4 x i16> noundef %arg2, <4 x i16> noundef %arg3, <4 x i16> noundef %arg4, <4 x i16> noundef %arg5, <4 x i16> noundef %arg6, %struct.struct_4regs addrspace(5)* nocapture noundef byval(%struct.struct_4regs) align 4 %arg7)
 490 void v4i16_pair_reg_count_over(short4 arg0, short4 arg1, short4 arg2, short4 arg3,
 491                                short4 arg4, short4 arg5, short4 arg6, struct_4regs arg7) { }
 492
 493 // CHECK: define{{.*}} void @v3i16_reg_count(<3 x i16> noundef %arg0, <3 x i16> noundef %arg1, <3 x i16> noundef %arg2, <3 x i16> noundef %arg3, <3 x i16> noundef %arg4, <3 x i16> noundef %arg5, i32 %arg6.coerce0, i32 %arg6.coerce1, i32 %arg6.coerce2, i32 %arg6.coerce3)
 494 void v3i16_reg_count(short3 arg0, short3 arg1, short3 arg2, short3 arg3,
 495                      short3 arg4, short3 arg5, struct_4regs arg6) { }
 496
 497 // CHECK: define{{.*}} void @v3i16_reg_count_over(<3 x i16> noundef %arg0, <3 x i16> noundef %arg1, <3 x i16> noundef %arg2, <3 x i16> noundef %arg3, <3 x i16> noundef %arg4, <3 x i16> noundef %arg5, <3 x i16> noundef %arg6, %struct.struct_4regs addrspace(5)* nocapture noundef byval(%struct.struct_4regs) align 4 %arg7)
 498 void v3i16_reg_count_over(short3 arg0, short3 arg1, short3 arg2, short3 arg3,
 499                           short3 arg4, short3 arg5, short3 arg6, struct_4regs arg7) { }
 500
 501 // CHECK: define{{.*}} void @v2i16_reg_count(<2 x i16> noundef %arg0, <2 x i16> noundef %arg1, <2 x i16> noundef %arg2, <2 x i16> noundef %arg3, <2 x i16> noundef %arg4, <2 x i16> noundef %arg5, <2 x i16> noundef %arg6, <2 x i16> noundef %arg7, <2 x i16> noundef %arg8, <2 x i16> noundef %arg9, <2 x i16> noundef %arg10, <2 x i16> noundef %arg11, i32 %arg13.coerce0, i32 %arg13.coerce1, i32 %arg13.coerce2, i32 %arg13.coerce3)
 502 void v2i16_reg_count(short2 arg0, short2 arg1, short2 arg2, short2 arg3,
 503                      short2 arg4, short2 arg5, short2 arg6, short2 arg7,
 504                      short2 arg8, short2 arg9, short2 arg10, short2 arg11,
 505                      struct_4regs arg13) { }
 506
 507 // CHECK: define{{.*}} void @v2i16_reg_count_over(<2 x i16> noundef %arg0, <2 x i16> noundef %arg1, <2 x i16> noundef %arg2, <2 x i16> noundef %arg3, <2 x i16> noundef %arg4, <2 x i16> noundef %arg5, <2 x i16> noundef %arg6, <2 x i16> noundef %arg7, <2 x i16> noundef %arg8, <2 x i16> noundef %arg9, <2 x i16> noundef %arg10, <2 x i16> noundef %arg11, <2 x i16> noundef %arg12, %struct.struct_4regs addrspace(5)* nocapture noundef byval(%struct.struct_4regs) align 4 %arg13)
 508 void v2i16_reg_count_over(short2 arg0, short2 arg1, short2 arg2, short2 arg3,
 509                           short2 arg4, short2 arg5, short2 arg6, short2 arg7,
 510                           short2 arg8, short2 arg9, short2 arg10, short2 arg11,
 511                           short2 arg12, struct_4regs arg13) { }
 512
 513 // CHECK: define{{.*}} void @v2i8_reg_count(<2 x i8> noundef %arg0, <2 x i8> noundef %arg1, <2 x i8> noundef %arg2, <2 x i8> noundef %arg3, <2 x i8> noundef %arg4, <2 x i8> noundef %arg5, i32 %arg6.coerce0, i32 %arg6.coerce1, i32 %arg6.coerce2, i32 %arg6.coerce3)
 514 void v2i8_reg_count(char2 arg0, char2 arg1, char2 arg2, char2 arg3,
 515                     char2 arg4, char2 arg5, struct_4regs arg6) { }
 516
 517 // CHECK: define{{.*}} void @v2i8_reg_count_over(<2 x i8> noundef %arg0, <2 x i8> noundef %arg1, <2 x i8> noundef %arg2, <2 x i8> noundef %arg3, <2 x i8> noundef %arg4, <2 x i8> noundef %arg5, i32 noundef %arg6, %struct.struct_4regs addrspace(5)* nocapture noundef byval(%struct.struct_4regs) align 4 %arg7)
 518 void v2i8_reg_count_over(char2 arg0, char2 arg1, char2 arg2, char2 arg3,
 519                          char2 arg4, char2 arg5, int arg6, struct_4regs arg7) { }
 520
 521 // CHECK: define{{.*}} void @num_regs_left_64bit_aggregate(<4 x i32> noundef %arg0, <4 x i32> noundef %arg1, <4 x i32> noundef %arg2, <3 x i32> noundef %arg3, [2 x i32] %arg4.coerce, i32 noundef %arg5)
 522 void num_regs_left_64bit_aggregate(int4 arg0, int4 arg1, int4 arg2, int3 arg3, struct_char_x8 arg4, int arg5) { }