1; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s 2; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s 3; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s 4; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s 5; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s 6; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -amdgpu-dump-hsa-metadata -amdgpu-verify-hsa-metadata -filetype=obj -o - < %s 2>&1 | FileCheck --check-prefix=PARSER %s 7 8%struct.A = type { i8, float } 9%opencl.image1d_t = type opaque 10%opencl.image2d_t = type opaque 11%opencl.image3d_t = type opaque 12%opencl.queue_t = type opaque 13%opencl.pipe_t = type opaque 14%struct.B = type { ptr addrspace(1) } 15%opencl.clk_event_t = type opaque 16 17@__test_block_invoke_kernel_runtime_handle = external addrspace(1) externally_initialized constant ptr addrspace(1) 18 19; CHECK: --- 20; CHECK-NEXT: amdhsa.kernels: 21; CHECK-NEXT: - .args: 22; CHECK-NEXT: - .name: a 23; CHECK-NEXT: .offset: 0 24; CHECK-NEXT: .size: 1 25; CHECK-NEXT: .type_name: char 26; CHECK-NEXT: .value_kind: by_value 27; CHECK-NEXT: - .offset: 8 28; CHECK-NEXT: .size: 8 29; CHECK-NEXT: .value_kind: hidden_global_offset_x 30; CHECK-NEXT: - .offset: 16 31; CHECK-NEXT: .size: 8 32; CHECK-NEXT: .value_kind: hidden_global_offset_y 33; CHECK-NEXT: - .offset: 24 34; CHECK-NEXT: .size: 8 35; CHECK-NEXT: .value_kind: hidden_global_offset_z 36; CHECK-NEXT: - .offset: 32 37; CHECK-NEXT: .size: 8 38; CHECK-NOT: .value_kind: hidden_default_queue 39; CHECK-NOT: .value_kind: hidden_completion_action 40; CHECK-NOT: .value_kind: hidden_hostcall_buffer 41; CHECK-NEXT: .value_kind: hidden_printf_buffer 42; CHECK: .value_kind: hidden_multigrid_sync_arg 43; CHECK: .language: OpenCL C 44; CHECK-NEXT: .language_version: 45; CHECK-NEXT: - 2 46; CHECK-NEXT: - 0 47; CHECK: .name: test_char 48; CHECK: .symbol: test_char.kd 49define amdgpu_kernel void @test_char(i8 %a) #0 50 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9 51 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 { 52 ret void 53} 54 55; CHECK: - .args: 56; CHECK-NEXT: - .name: a 57; CHECK-NEXT: .offset: 0 58; CHECK-NEXT: .size: 1 59; CHECK-NEXT: .type_name: char 60; CHECK-NEXT: .value_kind: by_value 61; CHECK-NEXT: - .offset: 8 62; CHECK-NEXT: .size: 8 63; CHECK-NEXT: .value_kind: hidden_global_offset_x 64; CHECK-NEXT: - .offset: 16 65; CHECK-NEXT: .size: 8 66; CHECK-NEXT: .value_kind: hidden_global_offset_y 67; CHECK-NEXT: - .offset: 24 68; CHECK-NEXT: .size: 8 69; CHECK-NEXT: .value_kind: hidden_global_offset_z 70; CHECK-NEXT: - .offset: 32 71; CHECK-NEXT: .size: 8 72; CHECK-NOT: .value_kind: hidden_default_queue 73; CHECK-NOT: .value_kind: hidden_completion_action 74; CHECK-NOT: .value_kind: hidden_hostcall_buffer 75; CHECK-NEXT: .value_kind: hidden_printf_buffer 76; CHECK: .value_kind: hidden_multigrid_sync_arg 77; CHECK: .language: OpenCL C 78; CHECK-NEXT: .language_version: 79; CHECK-NEXT: - 2 80; CHECK-NEXT: - 0 81; CHECK: .name: test_char_byref_constant 82; CHECK: .symbol: test_char_byref_constant.kd 83define amdgpu_kernel void @test_char_byref_constant(ptr addrspace(4) byref(i8) %a) #0 84 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9 85 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 { 86 ret void 87} 88 89; CHECK: - .args: 90; CHECK-NEXT: - .offset: 0 91; CHECK-NEXT: .size: 1 92; CHECK-NEXT: .type_name: char 93; CHECK-NEXT: .value_kind: by_value 94; CHECK-NEXT: - .name: a 95; CHECK-NEXT: .offset: 512 96; CHECK-NEXT: .size: 1 97; CHECK-NEXT: .type_name: char 98; CHECK-NEXT: .value_kind: by_value 99; CHECK-NEXT: - .offset: 520 100; CHECK-NEXT: .size: 8 101; CHECK-NEXT: .value_kind: hidden_global_offset_x 102; CHECK-NEXT: - .offset: 528 103; CHECK-NEXT: .size: 8 104; CHECK-NEXT: .value_kind: hidden_global_offset_y 105; CHECK-NEXT: - .offset: 536 106; CHECK-NEXT: .size: 8 107; CHECK-NEXT: .value_kind: hidden_global_offset_z 108; CHECK-NEXT: - .offset: 544 109; CHECK-NEXT: .size: 8 110; CHECK-NOT: .value_kind: hidden_default_queue 111; CHECK-NOT: .value_kind: hidden_completion_action 112; CHECK-NOT: .value_kind: hidden_hostcall_buffer 113; CHECK-NEXT: .value_kind: hidden_printf_buffer 114; CHECK: .value_kind: hidden_multigrid_sync_arg 115; CHECK: .language: OpenCL C 116; CHECK-NEXT: .language_version: 117; CHECK-NEXT: - 2 118; CHECK-NEXT: - 0 119; CHECK: .name: test_char_byref_constant_align512 120; CHECK: .symbol: test_char_byref_constant_align512.kd 121define amdgpu_kernel void @test_char_byref_constant_align512(i8, ptr addrspace(4) byref(i8) align(512) %a) #0 122 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !111 123 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 { 124 ret void 125} 126 127; CHECK: - .args: 128; CHECK-NEXT: - .name: a 129; CHECK-NEXT: .offset: 0 130; CHECK-NEXT: .size: 4 131; CHECK-NEXT: .type_name: ushort2 132; CHECK-NEXT: .value_kind: by_value 133; CHECK-NEXT: - .offset: 8 134; CHECK-NEXT: .size: 8 135; CHECK-NEXT: .value_kind: hidden_global_offset_x 136; CHECK-NEXT: - .offset: 16 137; CHECK-NEXT: .size: 8 138; CHECK-NEXT: .value_kind: hidden_global_offset_y 139; CHECK-NEXT: - .offset: 24 140; CHECK-NEXT: .size: 8 141; CHECK-NEXT: .value_kind: hidden_global_offset_z 142; CHECK-NEXT: - .offset: 32 143; CHECK-NEXT: .size: 8 144; CHECK-NEXT: .value_kind: hidden_printf_buffer 145; CHECK-NEXT: - .offset: 40 146; CHECK-NEXT: .size: 8 147; CHECK-NEXT: .value_kind: hidden_none 148; CHECK-NEXT: - .offset: 48 149; CHECK-NEXT: .size: 8 150; CHECK-NEXT: .value_kind: hidden_none 151; CHECK-NEXT: - .offset: 56 152; CHECK-NEXT: .size: 8 153; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 154; CHECK: .language: OpenCL C 155; CHECK-NEXT: .language_version: 156; CHECK-NEXT: - 2 157; CHECK-NEXT: - 0 158; CHECK: .name: test_ushort2 159; CHECK: .symbol: test_ushort2.kd 160define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0 161 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10 162 !kernel_arg_base_type !10 !kernel_arg_type_qual !4 { 163 ret void 164} 165 166; CHECK: - .args: 167; CHECK-NEXT: - .name: a 168; CHECK-NEXT: .offset: 0 169; CHECK-NEXT: .size: 16 170; CHECK-NEXT: .type_name: int3 171; CHECK-NEXT: .value_kind: by_value 172; CHECK-NEXT: - .offset: 16 173; CHECK-NEXT: .size: 8 174; CHECK-NEXT: .value_kind: hidden_global_offset_x 175; CHECK-NEXT: - .offset: 24 176; CHECK-NEXT: .size: 8 177; CHECK-NEXT: .value_kind: hidden_global_offset_y 178; CHECK-NEXT: - .offset: 32 179; CHECK-NEXT: .size: 8 180; CHECK-NEXT: .value_kind: hidden_global_offset_z 181; CHECK-NEXT: - .offset: 40 182; CHECK-NEXT: .size: 8 183; CHECK-NEXT: .value_kind: hidden_printf_buffer 184; CHECK-NEXT: - .offset: 48 185; CHECK-NEXT: .size: 8 186; CHECK-NEXT: .value_kind: hidden_none 187; CHECK-NEXT: - .offset: 56 188; CHECK-NEXT: .size: 8 189; CHECK-NEXT: .value_kind: hidden_none 190; CHECK-NEXT: - .offset: 64 191; CHECK-NEXT: .size: 8 192; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 193; CHECK: .language: OpenCL C 194; CHECK-NEXT: .language_version: 195; CHECK-NEXT: - 2 196; CHECK-NEXT: - 0 197; CHECK: .name: test_int3 198; CHECK: .symbol: test_int3.kd 199define amdgpu_kernel void @test_int3(<3 x i32> %a) #0 200 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11 201 !kernel_arg_base_type !11 !kernel_arg_type_qual !4 { 202 ret void 203} 204 205; CHECK: - .args: 206; CHECK-NEXT: - .name: a 207; CHECK-NEXT: .offset: 0 208; CHECK-NEXT: .size: 32 209; CHECK-NEXT: .type_name: ulong4 210; CHECK-NEXT: .value_kind: by_value 211; CHECK-NEXT: - .offset: 32 212; CHECK-NEXT: .size: 8 213; CHECK-NEXT: .value_kind: hidden_global_offset_x 214; CHECK-NEXT: - .offset: 40 215; CHECK-NEXT: .size: 8 216; CHECK-NEXT: .value_kind: hidden_global_offset_y 217; CHECK-NEXT: - .offset: 48 218; CHECK-NEXT: .size: 8 219; CHECK-NEXT: .value_kind: hidden_global_offset_z 220; CHECK-NEXT: - .offset: 56 221; CHECK-NEXT: .size: 8 222; CHECK-NEXT: .value_kind: hidden_printf_buffer 223; CHECK-NEXT: - .offset: 64 224; CHECK-NEXT: .size: 8 225; CHECK-NEXT: .value_kind: hidden_none 226; CHECK-NEXT: - .offset: 72 227; CHECK-NEXT: .size: 8 228; CHECK-NEXT: .value_kind: hidden_none 229; CHECK-NEXT: - .offset: 80 230; CHECK-NEXT: .size: 8 231; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 232; CHECK: .language: OpenCL C 233; CHECK-NEXT: .language_version: 234; CHECK-NEXT: - 2 235; CHECK-NEXT: - 0 236; CHECK: .name: test_ulong4 237; CHECK: .symbol: test_ulong4.kd 238define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0 239 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12 240 !kernel_arg_base_type !12 !kernel_arg_type_qual !4 { 241 ret void 242} 243 244; CHECK: - .args: 245; CHECK-NEXT: - .name: a 246; CHECK-NEXT: .offset: 0 247; CHECK-NEXT: .size: 16 248; CHECK-NEXT: .type_name: half8 249; CHECK-NEXT: .value_kind: by_value 250; CHECK-NEXT: - .offset: 16 251; CHECK-NEXT: .size: 8 252; CHECK-NEXT: .value_kind: hidden_global_offset_x 253; CHECK-NEXT: - .offset: 24 254; CHECK-NEXT: .size: 8 255; CHECK-NEXT: .value_kind: hidden_global_offset_y 256; CHECK-NEXT: - .offset: 32 257; CHECK-NEXT: .size: 8 258; CHECK-NEXT: .value_kind: hidden_global_offset_z 259; CHECK-NEXT: - .offset: 40 260; CHECK-NEXT: .size: 8 261; CHECK-NEXT: .value_kind: hidden_printf_buffer 262; CHECK-NEXT: - .offset: 48 263; CHECK-NEXT: .size: 8 264; CHECK-NEXT: .value_kind: hidden_none 265; CHECK-NEXT: - .offset: 56 266; CHECK-NEXT: .size: 8 267; CHECK-NEXT: .value_kind: hidden_none 268; CHECK-NEXT: - .offset: 64 269; CHECK-NEXT: .size: 8 270; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 271; CHECK: .language: OpenCL C 272; CHECK-NEXT: .language_version: 273; CHECK-NEXT: - 2 274; CHECK-NEXT: - 0 275; CHECK: .name: test_half8 276; CHECK: .symbol: test_half8.kd 277define amdgpu_kernel void @test_half8(<8 x half> %a) #0 278 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13 279 !kernel_arg_base_type !13 !kernel_arg_type_qual !4 { 280 ret void 281} 282 283; CHECK: - .args: 284; CHECK-NEXT: - .name: a 285; CHECK-NEXT: .offset: 0 286; CHECK-NEXT: .size: 64 287; CHECK-NEXT: .type_name: float16 288; CHECK-NEXT: .value_kind: by_value 289; CHECK-NEXT: - .offset: 64 290; CHECK-NEXT: .size: 8 291; CHECK-NEXT: .value_kind: hidden_global_offset_x 292; CHECK-NEXT: - .offset: 72 293; CHECK-NEXT: .size: 8 294; CHECK-NEXT: .value_kind: hidden_global_offset_y 295; CHECK-NEXT: - .offset: 80 296; CHECK-NEXT: .size: 8 297; CHECK-NEXT: .value_kind: hidden_global_offset_z 298; CHECK-NEXT: - .offset: 88 299; CHECK-NEXT: .size: 8 300; CHECK-NEXT: .value_kind: hidden_printf_buffer 301; CHECK-NEXT: - .offset: 96 302; CHECK-NEXT: .size: 8 303; CHECK-NEXT: .value_kind: hidden_none 304; CHECK-NEXT: - .offset: 104 305; CHECK-NEXT: .size: 8 306; CHECK-NEXT: .value_kind: hidden_none 307; CHECK-NEXT: - .offset: 112 308; CHECK-NEXT: .size: 8 309; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 310; CHECK: .language: OpenCL C 311; CHECK-NEXT: .language_version: 312; CHECK-NEXT: - 2 313; CHECK-NEXT: - 0 314; CHECK: .name: test_float16 315; CHECK: .symbol: test_float16.kd 316define amdgpu_kernel void @test_float16(<16 x float> %a) #0 317 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14 318 !kernel_arg_base_type !14 !kernel_arg_type_qual !4 { 319 ret void 320} 321 322; CHECK: - .args: 323; CHECK-NEXT: - .name: a 324; CHECK-NEXT: .offset: 0 325; CHECK-NEXT: .size: 128 326; CHECK-NEXT: .type_name: double16 327; CHECK-NEXT: .value_kind: by_value 328; CHECK-NEXT: - .offset: 128 329; CHECK-NEXT: .size: 8 330; CHECK-NEXT: .value_kind: hidden_global_offset_x 331; CHECK-NEXT: - .offset: 136 332; CHECK-NEXT: .size: 8 333; CHECK-NEXT: .value_kind: hidden_global_offset_y 334; CHECK-NEXT: - .offset: 144 335; CHECK-NEXT: .size: 8 336; CHECK-NEXT: .value_kind: hidden_global_offset_z 337; CHECK-NEXT: - .offset: 152 338; CHECK-NEXT: .size: 8 339; CHECK-NEXT: .value_kind: hidden_printf_buffer 340; CHECK-NEXT: - .offset: 160 341; CHECK-NEXT: .size: 8 342; CHECK-NEXT: .value_kind: hidden_none 343; CHECK-NEXT: - .offset: 168 344; CHECK-NEXT: .size: 8 345; CHECK-NEXT: .value_kind: hidden_none 346; CHECK-NEXT: - .offset: 176 347; CHECK-NEXT: .size: 8 348; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 349; CHECK: .language: OpenCL C 350; CHECK-NEXT: .language_version: 351; CHECK-NEXT: - 2 352; CHECK-NEXT: - 0 353; CHECK: .name: test_double16 354; CHECK: .symbol: test_double16.kd 355define amdgpu_kernel void @test_double16(<16 x double> %a) #0 356 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15 357 !kernel_arg_base_type !15 !kernel_arg_type_qual !4 { 358 ret void 359} 360 361; CHECK: - .args: 362; CHECK-NEXT: - .address_space: global 363; CHECK-NEXT: .name: a 364; CHECK-NEXT: .offset: 0 365; CHECK-NEXT: .size: 8 366; CHECK-NEXT: .type_name: 'int addrspace(5)*' 367; CHECK-NEXT: .value_kind: global_buffer 368; CHECK-NEXT: - .offset: 8 369; CHECK-NEXT: .size: 8 370; CHECK-NEXT: .value_kind: hidden_global_offset_x 371; CHECK-NEXT: - .offset: 16 372; CHECK-NEXT: .size: 8 373; CHECK-NEXT: .value_kind: hidden_global_offset_y 374; CHECK-NEXT: - .offset: 24 375; CHECK-NEXT: .size: 8 376; CHECK-NEXT: .value_kind: hidden_global_offset_z 377; CHECK-NEXT: - .offset: 32 378; CHECK-NEXT: .size: 8 379; CHECK-NEXT: .value_kind: hidden_printf_buffer 380; CHECK-NEXT: - .offset: 40 381; CHECK-NEXT: .size: 8 382; CHECK-NEXT: .value_kind: hidden_none 383; CHECK-NEXT: - .offset: 48 384; CHECK-NEXT: .size: 8 385; CHECK-NEXT: .value_kind: hidden_none 386; CHECK-NEXT: - .offset: 56 387; CHECK-NEXT: .size: 8 388; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 389; CHECK: .language: OpenCL C 390; CHECK-NEXT: .language_version: 391; CHECK-NEXT: - 2 392; CHECK-NEXT: - 0 393; CHECK: .name: test_pointer 394; CHECK: .symbol: test_pointer.kd 395define amdgpu_kernel void @test_pointer(ptr addrspace(1) %a) #0 396 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16 397 !kernel_arg_base_type !16 !kernel_arg_type_qual !4 { 398 ret void 399} 400 401; CHECK: - .args: 402; CHECK-NEXT: - .name: a 403; CHECK-NEXT: .offset: 0 404; CHECK-NEXT: .size: 8 405; CHECK-NEXT: .type_name: image2d_t 406; CHECK-NEXT: .value_kind: image 407; CHECK-NEXT: - .offset: 8 408; CHECK-NEXT: .size: 8 409; CHECK-NEXT: .value_kind: hidden_global_offset_x 410; CHECK-NEXT: - .offset: 16 411; CHECK-NEXT: .size: 8 412; CHECK-NEXT: .value_kind: hidden_global_offset_y 413; CHECK-NEXT: - .offset: 24 414; CHECK-NEXT: .size: 8 415; CHECK-NEXT: .value_kind: hidden_global_offset_z 416; CHECK-NEXT: - .offset: 32 417; CHECK-NEXT: .size: 8 418; CHECK-NEXT: .value_kind: hidden_printf_buffer 419; CHECK-NEXT: - .offset: 40 420; CHECK-NEXT: .size: 8 421; CHECK-NEXT: .value_kind: hidden_none 422; CHECK-NEXT: - .offset: 48 423; CHECK-NEXT: .size: 8 424; CHECK-NEXT: .value_kind: hidden_none 425; CHECK-NEXT: - .offset: 56 426; CHECK-NEXT: .size: 8 427; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 428; CHECK: .language: OpenCL C 429; CHECK-NEXT: .language_version: 430; CHECK-NEXT: - 2 431; CHECK-NEXT: - 0 432; CHECK: .name: test_image 433; CHECK: .symbol: test_image.kd 434define amdgpu_kernel void @test_image(ptr addrspace(1) %a) #0 435 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17 436 !kernel_arg_base_type !17 !kernel_arg_type_qual !4 { 437 ret void 438} 439 440; CHECK: - .args: 441; CHECK-NEXT: - .name: a 442; CHECK-NEXT: .offset: 0 443; CHECK-NEXT: .size: 4 444; CHECK-NEXT: .type_name: sampler_t 445; CHECK-NEXT: .value_kind: sampler 446; CHECK-NEXT: - .offset: 8 447; CHECK-NEXT: .size: 8 448; CHECK-NEXT: .value_kind: hidden_global_offset_x 449; CHECK-NEXT: - .offset: 16 450; CHECK-NEXT: .size: 8 451; CHECK-NEXT: .value_kind: hidden_global_offset_y 452; CHECK-NEXT: - .offset: 24 453; CHECK-NEXT: .size: 8 454; CHECK-NEXT: .value_kind: hidden_global_offset_z 455; CHECK-NEXT: - .offset: 32 456; CHECK-NEXT: .size: 8 457; CHECK-NEXT: .value_kind: hidden_printf_buffer 458; CHECK-NEXT: - .offset: 40 459; CHECK-NEXT: .size: 8 460; CHECK-NEXT: .value_kind: hidden_none 461; CHECK-NEXT: - .offset: 48 462; CHECK-NEXT: .size: 8 463; CHECK-NEXT: .value_kind: hidden_none 464; CHECK-NEXT: - .offset: 56 465; CHECK-NEXT: .size: 8 466; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 467; CHECK: .language: OpenCL C 468; CHECK-NEXT: .language_version: 469; CHECK-NEXT: - 2 470; CHECK-NEXT: - 0 471; CHECK: .name: test_sampler 472; CHECK: .symbol: test_sampler.kd 473define amdgpu_kernel void @test_sampler(i32 %a) #0 474 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18 475 !kernel_arg_base_type !18 !kernel_arg_type_qual !4 { 476 ret void 477} 478 479; CHECK: - .args: 480; CHECK-NEXT: - .name: a 481; CHECK-NEXT: .offset: 0 482; CHECK-NEXT: .size: 8 483; CHECK-NEXT: .type_name: queue_t 484; CHECK-NEXT: .value_kind: queue 485; CHECK-NEXT: - .offset: 8 486; CHECK-NEXT: .size: 8 487; CHECK-NEXT: .value_kind: hidden_global_offset_x 488; CHECK-NEXT: - .offset: 16 489; CHECK-NEXT: .size: 8 490; CHECK-NEXT: .value_kind: hidden_global_offset_y 491; CHECK-NEXT: - .offset: 24 492; CHECK-NEXT: .size: 8 493; CHECK-NEXT: .value_kind: hidden_global_offset_z 494; CHECK-NEXT: - .offset: 32 495; CHECK-NEXT: .size: 8 496; CHECK-NEXT: .value_kind: hidden_printf_buffer 497; CHECK-NEXT: - .offset: 40 498; CHECK-NEXT: .size: 8 499; CHECK-NEXT: .value_kind: hidden_none 500; CHECK-NEXT: - .offset: 48 501; CHECK-NEXT: .size: 8 502; CHECK-NEXT: .value_kind: hidden_none 503; CHECK-NEXT: - .offset: 56 504; CHECK-NEXT: .size: 8 505; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 506; CHECK: .language: OpenCL C 507; CHECK-NEXT: .language_version: 508; CHECK-NEXT: - 2 509; CHECK-NEXT: - 0 510; CHECK: .name: test_queue 511; CHECK: .symbol: test_queue.kd 512define amdgpu_kernel void @test_queue(ptr addrspace(1) %a) #0 513 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19 514 !kernel_arg_base_type !19 !kernel_arg_type_qual !4 { 515 ret void 516} 517 518; CHECK: - .args: 519; CHECK-NEXT: .name: a 520; CHECK-NEXT: .offset: 0 521; CHECK-NEXT: .size: 8 522; CHECK-NEXT: .type_name: struct A 523; CHECK-NEXT: .value_kind: by_value 524; CHECK-NEXT: - .offset: 8 525; CHECK-NEXT: .size: 8 526; CHECK-NEXT: .value_kind: hidden_global_offset_x 527; CHECK-NEXT: - .offset: 16 528; CHECK-NEXT: .size: 8 529; CHECK-NEXT: .value_kind: hidden_global_offset_y 530; CHECK-NEXT: - .offset: 24 531; CHECK-NEXT: .size: 8 532; CHECK-NEXT: .value_kind: hidden_global_offset_z 533; CHECK-NEXT: - .offset: 32 534; CHECK-NEXT: .size: 8 535; CHECK-NEXT: .value_kind: hidden_printf_buffer 536; CHECK-NEXT: - .offset: 40 537; CHECK-NEXT: .size: 8 538; CHECK-NEXT: .value_kind: hidden_none 539; CHECK-NEXT: - .offset: 48 540; CHECK-NEXT: .size: 8 541; CHECK-NEXT: .value_kind: hidden_none 542; CHECK-NEXT: - .offset: 56 543; CHECK-NEXT: .size: 8 544; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 545; CHECK: .language: OpenCL C 546; CHECK-NEXT: .language_version: 547; CHECK-NEXT: - 2 548; CHECK-NEXT: - 0 549; CHECK: .name: test_struct 550; CHECK: .symbol: test_struct.kd 551define amdgpu_kernel void @test_struct(%struct.A %a) #0 552 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20 553 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 { 554 ret void 555} 556 557; CHECK: - .args: 558; CHECK-NEXT: .name: a 559; CHECK-NEXT: .offset: 0 560; CHECK-NEXT: .size: 8 561; CHECK-NEXT: .type_name: struct A 562; CHECK-NEXT: .value_kind: by_value 563; CHECK-NEXT: - .offset: 8 564; CHECK-NEXT: .size: 8 565; CHECK-NEXT: .value_kind: hidden_global_offset_x 566; CHECK-NEXT: - .offset: 16 567; CHECK-NEXT: .size: 8 568; CHECK-NEXT: .value_kind: hidden_global_offset_y 569; CHECK-NEXT: - .offset: 24 570; CHECK-NEXT: .size: 8 571; CHECK-NEXT: .value_kind: hidden_global_offset_z 572; CHECK-NEXT: - .offset: 32 573; CHECK-NEXT: .size: 8 574; CHECK-NEXT: .value_kind: hidden_printf_buffer 575; CHECK-NEXT: - .offset: 40 576; CHECK-NEXT: .size: 8 577; CHECK-NEXT: .value_kind: hidden_none 578; CHECK-NEXT: - .offset: 48 579; CHECK-NEXT: .size: 8 580; CHECK-NEXT: .value_kind: hidden_none 581; CHECK-NEXT: - .offset: 56 582; CHECK-NEXT: .size: 8 583; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 584; CHECK: .language: OpenCL C 585; CHECK-NEXT: .language_version: 586; CHECK-NEXT: - 2 587; CHECK-NEXT: - 0 588; CHECK: .name: test_struct_byref_constant 589; CHECK: .symbol: test_struct_byref_constant.kd 590define amdgpu_kernel void @test_struct_byref_constant(ptr addrspace(4) byref(%struct.A) %a) #0 591 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20 592 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 { 593 ret void 594} 595 596; CHECK: - .args: 597; CHECK-NEXT: .name: a 598; CHECK-NEXT: .offset: 0 599; CHECK-NEXT: .size: 32 600; CHECK-NEXT: .type_name: struct A 601; CHECK-NEXT: .value_kind: by_value 602; CHECK-NEXT: - .offset: 32 603; CHECK-NEXT: .size: 8 604; CHECK-NEXT: .value_kind: hidden_global_offset_x 605; CHECK-NEXT: - .offset: 40 606; CHECK-NEXT: .size: 8 607; CHECK-NEXT: .value_kind: hidden_global_offset_y 608; CHECK-NEXT: - .offset: 48 609; CHECK-NEXT: .size: 8 610; CHECK-NEXT: .value_kind: hidden_global_offset_z 611; CHECK-NEXT: - .offset: 56 612; CHECK-NEXT: .size: 8 613; CHECK-NEXT: .value_kind: hidden_printf_buffer 614; CHECK-NEXT: - .offset: 64 615; CHECK-NEXT: .size: 8 616; CHECK-NEXT: .value_kind: hidden_none 617; CHECK-NEXT: - .offset: 72 618; CHECK-NEXT: .size: 8 619; CHECK-NEXT: .value_kind: hidden_none 620; CHECK-NEXT: - .offset: 80 621; CHECK-NEXT: .size: 8 622; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 623; CHECK: .language: OpenCL C 624; CHECK-NEXT: .language_version: 625; CHECK-NEXT: - 2 626; CHECK-NEXT: - 0 627; CHECK: .name: test_array 628; CHECK: .symbol: test_array.kd 629define amdgpu_kernel void @test_array([32 x i8] %a) #0 630 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20 631 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 { 632 ret void 633} 634 635; CHECK: - .args: 636; CHECK-NEXT: .name: a 637; CHECK-NEXT: .offset: 0 638; CHECK-NEXT: .size: 32 639; CHECK-NEXT: .type_name: struct A 640; CHECK-NEXT: .value_kind: by_value 641; CHECK-NEXT: - .offset: 32 642; CHECK-NEXT: .size: 8 643; CHECK-NEXT: .value_kind: hidden_global_offset_x 644; CHECK-NEXT: - .offset: 40 645; CHECK-NEXT: .size: 8 646; CHECK-NEXT: .value_kind: hidden_global_offset_y 647; CHECK-NEXT: - .offset: 48 648; CHECK-NEXT: .size: 8 649; CHECK-NEXT: .value_kind: hidden_global_offset_z 650; CHECK-NEXT: - .offset: 56 651; CHECK-NEXT: .size: 8 652; CHECK-NEXT: .value_kind: hidden_printf_buffer 653; CHECK-NEXT: - .offset: 64 654; CHECK-NEXT: .size: 8 655; CHECK-NEXT: .value_kind: hidden_none 656; CHECK-NEXT: - .offset: 72 657; CHECK-NEXT: .size: 8 658; CHECK-NEXT: .value_kind: hidden_none 659; CHECK-NEXT: - .offset: 80 660; CHECK-NEXT: .size: 8 661; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 662; CHECK: .language: OpenCL C 663; CHECK-NEXT: .language_version: 664; CHECK-NEXT: - 2 665; CHECK-NEXT: - 0 666; CHECK: .name: test_array_byref_constant 667; CHECK: .symbol: test_array_byref_constant.kd 668define amdgpu_kernel void @test_array_byref_constant(ptr addrspace(4) byref([32 x i8]) %a) #0 669 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20 670 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 { 671 ret void 672} 673 674; CHECK: - .args: 675; CHECK-NEXT: - .name: a 676; CHECK-NEXT: .offset: 0 677; CHECK-NEXT: .size: 16 678; CHECK-NEXT: .type_name: i128 679; CHECK-NEXT: .value_kind: by_value 680; CHECK-NEXT: - .offset: 16 681; CHECK-NEXT: .size: 8 682; CHECK-NEXT: .value_kind: hidden_global_offset_x 683; CHECK-NEXT: - .offset: 24 684; CHECK-NEXT: .size: 8 685; CHECK-NEXT: .value_kind: hidden_global_offset_y 686; CHECK-NEXT: - .offset: 32 687; CHECK-NEXT: .size: 8 688; CHECK-NEXT: .value_kind: hidden_global_offset_z 689; CHECK-NEXT: - .offset: 40 690; CHECK-NEXT: .size: 8 691; CHECK-NEXT: .value_kind: hidden_printf_buffer 692; CHECK-NEXT: - .offset: 48 693; CHECK-NEXT: .size: 8 694; CHECK-NEXT: .value_kind: hidden_none 695; CHECK-NEXT: - .offset: 56 696; CHECK-NEXT: .size: 8 697; CHECK-NEXT: .value_kind: hidden_none 698; CHECK-NEXT: - .offset: 64 699; CHECK-NEXT: .size: 8 700; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 701; CHECK: .language: OpenCL C 702; CHECK-NEXT: .language_version: 703; CHECK-NEXT: - 2 704; CHECK-NEXT: - 0 705; CHECK: .name: test_i128 706; CHECK: .symbol: test_i128.kd 707define amdgpu_kernel void @test_i128(i128 %a) #0 708 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21 709 !kernel_arg_base_type !21 !kernel_arg_type_qual !4 { 710 ret void 711} 712 713; CHECK: - .args: 714; CHECK-NEXT: - .name: a 715; CHECK-NEXT: .offset: 0 716; CHECK-NEXT: .size: 4 717; CHECK-NEXT: .type_name: int 718; CHECK-NEXT: .value_kind: by_value 719; CHECK-NEXT: - .name: b 720; CHECK-NEXT: .offset: 4 721; CHECK-NEXT: .size: 4 722; CHECK-NEXT: .type_name: short2 723; CHECK-NEXT: .value_kind: by_value 724; CHECK-NEXT: - .name: c 725; CHECK-NEXT: .offset: 8 726; CHECK-NEXT: .size: 4 727; CHECK-NEXT: .type_name: char3 728; CHECK-NEXT: .value_kind: by_value 729; CHECK-NEXT: - .offset: 16 730; CHECK-NEXT: .size: 8 731; CHECK-NEXT: .value_kind: hidden_global_offset_x 732; CHECK-NEXT: - .offset: 24 733; CHECK-NEXT: .size: 8 734; CHECK-NEXT: .value_kind: hidden_global_offset_y 735; CHECK-NEXT: - .offset: 32 736; CHECK-NEXT: .size: 8 737; CHECK-NEXT: .value_kind: hidden_global_offset_z 738; CHECK-NEXT: - .offset: 40 739; CHECK-NEXT: .size: 8 740; CHECK-NEXT: .value_kind: hidden_printf_buffer 741; CHECK-NEXT: - .offset: 48 742; CHECK-NEXT: .size: 8 743; CHECK-NEXT: .value_kind: hidden_none 744; CHECK-NEXT: - .offset: 56 745; CHECK-NEXT: .size: 8 746; CHECK-NEXT: .value_kind: hidden_none 747; CHECK-NEXT: - .offset: 64 748; CHECK-NEXT: .size: 8 749; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 750; CHECK: .language: OpenCL C 751; CHECK-NEXT: .language_version: 752; CHECK-NEXT: - 2 753; CHECK-NEXT: - 0 754; CHECK: .name: test_multi_arg 755; CHECK: .symbol: test_multi_arg.kd 756define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0 757 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24 758 !kernel_arg_base_type !24 !kernel_arg_type_qual !25 { 759 ret void 760} 761 762; CHECK: - .args: 763; CHECK-NEXT: - .address_space: global 764; CHECK-NEXT: .name: g 765; CHECK-NEXT: .offset: 0 766; CHECK-NEXT: .size: 8 767; CHECK-NEXT: .type_name: 'int addrspace(5)*' 768; CHECK-NEXT: .value_kind: global_buffer 769; CHECK-NEXT: - .address_space: constant 770; CHECK-NEXT: .name: c 771; CHECK-NEXT: .offset: 8 772; CHECK-NEXT: .size: 8 773; CHECK-NEXT: .type_name: 'int addrspace(5)*' 774; CHECK-NEXT: .value_kind: global_buffer 775; CHECK-NEXT: - .address_space: local 776; CHECK-NEXT: .name: l 777; CHECK-NEXT: .offset: 16 778; CHECK-NEXT: .pointee_align: 4 779; CHECK-NEXT: .size: 4 780; CHECK-NEXT: .type_name: 'int addrspace(5)*' 781; CHECK-NEXT: .value_kind: dynamic_shared_pointer 782; CHECK-NEXT: - .offset: 24 783; CHECK-NEXT: .size: 8 784; CHECK-NEXT: .value_kind: hidden_global_offset_x 785; CHECK-NEXT: - .offset: 32 786; CHECK-NEXT: .size: 8 787; CHECK-NEXT: .value_kind: hidden_global_offset_y 788; CHECK-NEXT: - .offset: 40 789; CHECK-NEXT: .size: 8 790; CHECK-NEXT: .value_kind: hidden_global_offset_z 791; CHECK-NEXT: - .offset: 48 792; CHECK-NEXT: .size: 8 793; CHECK-NEXT: .value_kind: hidden_printf_buffer 794; CHECK-NEXT: - .offset: 56 795; CHECK-NEXT: .size: 8 796; CHECK-NEXT: .value_kind: hidden_none 797; CHECK-NEXT: - .offset: 64 798; CHECK-NEXT: .size: 8 799; CHECK-NEXT: .value_kind: hidden_none 800; CHECK-NEXT: - .offset: 72 801; CHECK-NEXT: .size: 8 802; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 803; CHECK: .language: OpenCL C 804; CHECK-NEXT: .language_version: 805; CHECK-NEXT: - 2 806; CHECK-NEXT: - 0 807; CHECK: .name: test_addr_space 808; CHECK: .symbol: test_addr_space.kd 809define amdgpu_kernel void @test_addr_space(ptr addrspace(1) %g, 810 ptr addrspace(4) %c, 811 ptr addrspace(3) align 4 %l) #0 812 !kernel_arg_addr_space !50 !kernel_arg_access_qual !23 !kernel_arg_type !51 813 !kernel_arg_base_type !51 !kernel_arg_type_qual !25 { 814 ret void 815} 816 817; CHECK: - .args: 818; CHECK-NEXT: - .address_space: global 819; CHECK-NEXT: .is_volatile: true 820; CHECK-NEXT: .name: a 821; CHECK-NEXT: .offset: 0 822; CHECK-NEXT: .size: 8 823; CHECK-NEXT: .type_name: 'int addrspace(5)*' 824; CHECK-NEXT: .value_kind: global_buffer 825; CHECK-NEXT: - .address_space: global 826; CHECK-NEXT: .is_const: true 827; CHECK-NEXT: .is_restrict: true 828; CHECK-NEXT: .name: b 829; CHECK-NEXT: .offset: 8 830; CHECK-NEXT: .size: 8 831; CHECK-NEXT: .type_name: 'int addrspace(5)*' 832; CHECK-NEXT: .value_kind: global_buffer 833; CHECK-NEXT: - .is_pipe: true 834; CHECK-NEXT: .name: c 835; CHECK-NEXT: .offset: 16 836; CHECK-NEXT: .size: 8 837; CHECK-NEXT: .type_name: 'int addrspace(5)*' 838; CHECK-NEXT: .value_kind: pipe 839; CHECK-NEXT: - .offset: 24 840; CHECK-NEXT: .size: 8 841; CHECK-NEXT: .value_kind: hidden_global_offset_x 842; CHECK-NEXT: - .offset: 32 843; CHECK-NEXT: .size: 8 844; CHECK-NEXT: .value_kind: hidden_global_offset_y 845; CHECK-NEXT: - .offset: 40 846; CHECK-NEXT: .size: 8 847; CHECK-NEXT: .value_kind: hidden_global_offset_z 848; CHECK-NEXT: - .offset: 48 849; CHECK-NEXT: .size: 8 850; CHECK-NEXT: .value_kind: hidden_printf_buffer 851; CHECK-NEXT: - .offset: 56 852; CHECK-NEXT: .size: 8 853; CHECK-NEXT: .value_kind: hidden_none 854; CHECK-NEXT: - .offset: 64 855; CHECK-NEXT: .size: 8 856; CHECK-NEXT: .value_kind: hidden_none 857; CHECK-NEXT: - .offset: 72 858; CHECK-NEXT: .size: 8 859; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 860; CHECK: .language: OpenCL C 861; CHECK-NEXT: .language_version: 862; CHECK-NEXT: - 2 863; CHECK-NEXT: - 0 864; CHECK: .name: test_type_qual 865; CHECK: .symbol: test_type_qual.kd 866define amdgpu_kernel void @test_type_qual(ptr addrspace(1) %a, 867 ptr addrspace(1) %b, 868 ptr addrspace(1) %c) #0 869 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !51 870 !kernel_arg_base_type !51 !kernel_arg_type_qual !70 { 871 ret void 872} 873 874; CHECK: - .args: 875; CHECK-NEXT: - .access: read_only 876; CHECK-NEXT: .name: ro 877; CHECK-NEXT: .offset: 0 878; CHECK-NEXT: .size: 8 879; CHECK-NEXT: .type_name: image1d_t 880; CHECK-NEXT: .value_kind: image 881; CHECK-NEXT: - .access: write_only 882; CHECK-NEXT: .name: wo 883; CHECK-NEXT: .offset: 8 884; CHECK-NEXT: .size: 8 885; CHECK-NEXT: .type_name: image2d_t 886; CHECK-NEXT: .value_kind: image 887; CHECK-NEXT: - .access: read_write 888; CHECK-NEXT: .name: rw 889; CHECK-NEXT: .offset: 16 890; CHECK-NEXT: .size: 8 891; CHECK-NEXT: .type_name: image3d_t 892; CHECK-NEXT: .value_kind: image 893; CHECK-NEXT: - .offset: 24 894; CHECK-NEXT: .size: 8 895; CHECK-NEXT: .value_kind: hidden_global_offset_x 896; CHECK-NEXT: - .offset: 32 897; CHECK-NEXT: .size: 8 898; CHECK-NEXT: .value_kind: hidden_global_offset_y 899; CHECK-NEXT: - .offset: 40 900; CHECK-NEXT: .size: 8 901; CHECK-NEXT: .value_kind: hidden_global_offset_z 902; CHECK-NEXT: - .offset: 48 903; CHECK-NEXT: .size: 8 904; CHECK-NEXT: .value_kind: hidden_printf_buffer 905; CHECK-NEXT: - .offset: 56 906; CHECK-NEXT: .size: 8 907; CHECK-NEXT: .value_kind: hidden_none 908; CHECK-NEXT: - .offset: 64 909; CHECK-NEXT: .size: 8 910; CHECK-NEXT: .value_kind: hidden_none 911; CHECK-NEXT: - .offset: 72 912; CHECK-NEXT: .size: 8 913; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 914; CHECK: .language: OpenCL C 915; CHECK-NEXT: .language_version: 916; CHECK-NEXT: - 2 917; CHECK-NEXT: - 0 918; CHECK: .name: test_access_qual 919; CHECK: .symbol: test_access_qual.kd 920define amdgpu_kernel void @test_access_qual(ptr addrspace(1) %ro, 921 ptr addrspace(1) %wo, 922 ptr addrspace(1) %rw) #0 923 !kernel_arg_addr_space !60 !kernel_arg_access_qual !61 !kernel_arg_type !62 924 !kernel_arg_base_type !62 !kernel_arg_type_qual !25 { 925 ret void 926} 927 928; CHECK: - .args: 929; CHECK-NEXT: - .name: a 930; CHECK-NEXT: .offset: 0 931; CHECK-NEXT: .size: 4 932; CHECK-NEXT: .type_name: int 933; CHECK-NEXT: .value_kind: by_value 934; CHECK-NEXT: - .offset: 8 935; CHECK-NEXT: .size: 8 936; CHECK-NEXT: .value_kind: hidden_global_offset_x 937; CHECK-NEXT: - .offset: 16 938; CHECK-NEXT: .size: 8 939; CHECK-NEXT: .value_kind: hidden_global_offset_y 940; CHECK-NEXT: - .offset: 24 941; CHECK-NEXT: .size: 8 942; CHECK-NEXT: .value_kind: hidden_global_offset_z 943; CHECK-NEXT: - .offset: 32 944; CHECK-NEXT: .size: 8 945; CHECK-NEXT: .value_kind: hidden_printf_buffer 946; CHECK-NEXT: - .offset: 40 947; CHECK-NEXT: .size: 8 948; CHECK-NEXT: .value_kind: hidden_none 949; CHECK-NEXT: - .offset: 48 950; CHECK-NEXT: .size: 8 951; CHECK-NEXT: .value_kind: hidden_none 952; CHECK-NEXT: - .offset: 56 953; CHECK-NEXT: .size: 8 954; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 955; CHECK: .language: OpenCL C 956; CHECK-NEXT: .language_version: 957; CHECK-NEXT: - 2 958; CHECK-NEXT: - 0 959; CHECK: .name: test_vec_type_hint_half 960; CHECK: .symbol: test_vec_type_hint_half.kd 961; CHECK: .vec_type_hint: half 962define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0 963 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 964 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !26 { 965 ret void 966} 967 968; CHECK: - .args: 969; CHECK-NEXT: - .name: a 970; CHECK-NEXT: .offset: 0 971; CHECK-NEXT: .size: 4 972; CHECK-NEXT: .type_name: int 973; CHECK-NEXT: .value_kind: by_value 974; CHECK-NEXT: - .offset: 8 975; CHECK-NEXT: .size: 8 976; CHECK-NEXT: .value_kind: hidden_global_offset_x 977; CHECK-NEXT: - .offset: 16 978; CHECK-NEXT: .size: 8 979; CHECK-NEXT: .value_kind: hidden_global_offset_y 980; CHECK-NEXT: - .offset: 24 981; CHECK-NEXT: .size: 8 982; CHECK-NEXT: .value_kind: hidden_global_offset_z 983; CHECK-NEXT: - .offset: 32 984; CHECK-NEXT: .size: 8 985; CHECK-NEXT: .value_kind: hidden_printf_buffer 986; CHECK-NEXT: - .offset: 40 987; CHECK-NEXT: .size: 8 988; CHECK-NEXT: .value_kind: hidden_none 989; CHECK-NEXT: - .offset: 48 990; CHECK-NEXT: .size: 8 991; CHECK-NEXT: .value_kind: hidden_none 992; CHECK-NEXT: - .offset: 56 993; CHECK-NEXT: .size: 8 994; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 995; CHECK: .language: OpenCL C 996; CHECK-NEXT: .language_version: 997; CHECK-NEXT: - 2 998; CHECK-NEXT: - 0 999; CHECK: .name: test_vec_type_hint_float 1000; CHECK: .symbol: test_vec_type_hint_float.kd 1001; CHECK: .vec_type_hint: float 1002define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #0 1003 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1004 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !27 { 1005 ret void 1006} 1007 1008; CHECK: - .args: 1009; CHECK-NEXT: - .name: a 1010; CHECK-NEXT: .offset: 0 1011; CHECK-NEXT: .size: 4 1012; CHECK-NEXT: .type_name: int 1013; CHECK-NEXT: .value_kind: by_value 1014; CHECK-NEXT: - .offset: 8 1015; CHECK-NEXT: .size: 8 1016; CHECK-NEXT: .value_kind: hidden_global_offset_x 1017; CHECK-NEXT: - .offset: 16 1018; CHECK-NEXT: .size: 8 1019; CHECK-NEXT: .value_kind: hidden_global_offset_y 1020; CHECK-NEXT: - .offset: 24 1021; CHECK-NEXT: .size: 8 1022; CHECK-NEXT: .value_kind: hidden_global_offset_z 1023; CHECK-NEXT: - .offset: 32 1024; CHECK-NEXT: .size: 8 1025; CHECK-NEXT: .value_kind: hidden_printf_buffer 1026; CHECK-NEXT: - .offset: 40 1027; CHECK-NEXT: .size: 8 1028; CHECK-NEXT: .value_kind: hidden_none 1029; CHECK-NEXT: - .offset: 48 1030; CHECK-NEXT: .size: 8 1031; CHECK-NEXT: .value_kind: hidden_none 1032; CHECK-NEXT: - .offset: 56 1033; CHECK-NEXT: .size: 8 1034; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1035; CHECK: .language: OpenCL C 1036; CHECK-NEXT: .language_version: 1037; CHECK-NEXT: - 2 1038; CHECK-NEXT: - 0 1039; CHECK: .name: test_vec_type_hint_double 1040; CHECK: .symbol: test_vec_type_hint_double.kd 1041; CHECK: .vec_type_hint: double 1042define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #0 1043 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1044 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !28 { 1045 ret void 1046} 1047 1048; CHECK: - .args: 1049; CHECK-NEXT: - .name: a 1050; CHECK-NEXT: .offset: 0 1051; CHECK-NEXT: .size: 4 1052; CHECK-NEXT: .type_name: int 1053; CHECK-NEXT: .value_kind: by_value 1054; CHECK-NEXT: - .offset: 8 1055; CHECK-NEXT: .size: 8 1056; CHECK-NEXT: .value_kind: hidden_global_offset_x 1057; CHECK-NEXT: - .offset: 16 1058; CHECK-NEXT: .size: 8 1059; CHECK-NEXT: .value_kind: hidden_global_offset_y 1060; CHECK-NEXT: - .offset: 24 1061; CHECK-NEXT: .size: 8 1062; CHECK-NEXT: .value_kind: hidden_global_offset_z 1063; CHECK-NEXT: - .offset: 32 1064; CHECK-NEXT: .size: 8 1065; CHECK-NEXT: .value_kind: hidden_printf_buffer 1066; CHECK-NEXT: - .offset: 40 1067; CHECK-NEXT: .size: 8 1068; CHECK-NEXT: .value_kind: hidden_none 1069; CHECK-NEXT: - .offset: 48 1070; CHECK-NEXT: .size: 8 1071; CHECK-NEXT: .value_kind: hidden_none 1072; CHECK-NEXT: - .offset: 56 1073; CHECK-NEXT: .size: 8 1074; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1075; CHECK: .language: OpenCL C 1076; CHECK-NEXT: .language_version: 1077; CHECK-NEXT: - 2 1078; CHECK-NEXT: - 0 1079; CHECK: .name: test_vec_type_hint_char 1080; CHECK: .symbol: test_vec_type_hint_char.kd 1081; CHECK: .vec_type_hint: char 1082define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #0 1083 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1084 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !29 { 1085 ret void 1086} 1087 1088; CHECK: - .args: 1089; CHECK-NEXT: - .name: a 1090; CHECK-NEXT: .offset: 0 1091; CHECK-NEXT: .size: 4 1092; CHECK-NEXT: .type_name: int 1093; CHECK-NEXT: .value_kind: by_value 1094; CHECK-NEXT: - .offset: 8 1095; CHECK-NEXT: .size: 8 1096; CHECK-NEXT: .value_kind: hidden_global_offset_x 1097; CHECK-NEXT: - .offset: 16 1098; CHECK-NEXT: .size: 8 1099; CHECK-NEXT: .value_kind: hidden_global_offset_y 1100; CHECK-NEXT: - .offset: 24 1101; CHECK-NEXT: .size: 8 1102; CHECK-NEXT: .value_kind: hidden_global_offset_z 1103; CHECK-NEXT: - .offset: 32 1104; CHECK-NEXT: .size: 8 1105; CHECK-NEXT: .value_kind: hidden_printf_buffer 1106; CHECK-NEXT: - .offset: 40 1107; CHECK-NEXT: .size: 8 1108; CHECK-NEXT: .value_kind: hidden_none 1109; CHECK-NEXT: - .offset: 48 1110; CHECK-NEXT: .size: 8 1111; CHECK-NEXT: .value_kind: hidden_none 1112; CHECK-NEXT: - .offset: 56 1113; CHECK-NEXT: .size: 8 1114; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1115; CHECK: .language: OpenCL C 1116; CHECK-NEXT: .language_version: 1117; CHECK-NEXT: - 2 1118; CHECK-NEXT: - 0 1119; CHECK: .name: test_vec_type_hint_short 1120; CHECK: .symbol: test_vec_type_hint_short.kd 1121; CHECK: .vec_type_hint: short 1122define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #0 1123 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1124 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !30 { 1125 ret void 1126} 1127 1128; CHECK: - .args: 1129; CHECK-NEXT: - .name: a 1130; CHECK-NEXT: .offset: 0 1131; CHECK-NEXT: .size: 4 1132; CHECK-NEXT: .type_name: int 1133; CHECK-NEXT: .value_kind: by_value 1134; CHECK-NEXT: - .offset: 8 1135; CHECK-NEXT: .size: 8 1136; CHECK-NEXT: .value_kind: hidden_global_offset_x 1137; CHECK-NEXT: - .offset: 16 1138; CHECK-NEXT: .size: 8 1139; CHECK-NEXT: .value_kind: hidden_global_offset_y 1140; CHECK-NEXT: - .offset: 24 1141; CHECK-NEXT: .size: 8 1142; CHECK-NEXT: .value_kind: hidden_global_offset_z 1143; CHECK-NEXT: - .offset: 32 1144; CHECK-NEXT: .size: 8 1145; CHECK-NEXT: .value_kind: hidden_printf_buffer 1146; CHECK-NEXT: - .offset: 40 1147; CHECK-NEXT: .size: 8 1148; CHECK-NEXT: .value_kind: hidden_none 1149; CHECK-NEXT: - .offset: 48 1150; CHECK-NEXT: .size: 8 1151; CHECK-NEXT: .value_kind: hidden_none 1152; CHECK-NEXT: - .offset: 56 1153; CHECK-NEXT: .size: 8 1154; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1155; CHECK: .language: OpenCL C 1156; CHECK-NEXT: .language_version: 1157; CHECK-NEXT: - 2 1158; CHECK-NEXT: - 0 1159; CHECK: .name: test_vec_type_hint_long 1160; CHECK: .symbol: test_vec_type_hint_long.kd 1161; CHECK: .vec_type_hint: long 1162define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #0 1163 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1164 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !31 { 1165 ret void 1166} 1167 1168; CHECK: - .args: 1169; CHECK-NEXT: - .name: a 1170; CHECK-NEXT: .offset: 0 1171; CHECK-NEXT: .size: 4 1172; CHECK-NEXT: .type_name: int 1173; CHECK-NEXT: .value_kind: by_value 1174; CHECK-NEXT: - .offset: 8 1175; CHECK-NEXT: .size: 8 1176; CHECK-NEXT: .value_kind: hidden_global_offset_x 1177; CHECK-NEXT: - .offset: 16 1178; CHECK-NEXT: .size: 8 1179; CHECK-NEXT: .value_kind: hidden_global_offset_y 1180; CHECK-NEXT: - .offset: 24 1181; CHECK-NEXT: .size: 8 1182; CHECK-NEXT: .value_kind: hidden_global_offset_z 1183; CHECK-NEXT: - .offset: 32 1184; CHECK-NEXT: .size: 8 1185; CHECK-NEXT: .value_kind: hidden_printf_buffer 1186; CHECK-NEXT: - .offset: 40 1187; CHECK-NEXT: .size: 8 1188; CHECK-NEXT: .value_kind: hidden_none 1189; CHECK-NEXT: - .offset: 48 1190; CHECK-NEXT: .size: 8 1191; CHECK-NEXT: .value_kind: hidden_none 1192; CHECK-NEXT: - .offset: 56 1193; CHECK-NEXT: .size: 8 1194; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1195; CHECK: .language: OpenCL C 1196; CHECK-NEXT: .language_version: 1197; CHECK-NEXT: - 2 1198; CHECK-NEXT: - 0 1199; CHECK: .name: test_vec_type_hint_unknown 1200; CHECK: .symbol: test_vec_type_hint_unknown.kd 1201; CHECK: .vec_type_hint: unknown 1202define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #0 1203 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1204 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !32 { 1205 ret void 1206} 1207 1208; CHECK: - .args: 1209; CHECK-NEXT: - .name: a 1210; CHECK-NEXT: .offset: 0 1211; CHECK-NEXT: .size: 4 1212; CHECK-NEXT: .type_name: int 1213; CHECK-NEXT: .value_kind: by_value 1214; CHECK-NEXT: - .offset: 8 1215; CHECK-NEXT: .size: 8 1216; CHECK-NEXT: .value_kind: hidden_global_offset_x 1217; CHECK-NEXT: - .offset: 16 1218; CHECK-NEXT: .size: 8 1219; CHECK-NEXT: .value_kind: hidden_global_offset_y 1220; CHECK-NEXT: - .offset: 24 1221; CHECK-NEXT: .size: 8 1222; CHECK-NEXT: .value_kind: hidden_global_offset_z 1223; CHECK-NEXT: - .offset: 32 1224; CHECK-NEXT: .size: 8 1225; CHECK-NEXT: .value_kind: hidden_printf_buffer 1226; CHECK-NEXT: - .offset: 40 1227; CHECK-NEXT: .size: 8 1228; CHECK-NEXT: .value_kind: hidden_none 1229; CHECK-NEXT: - .offset: 48 1230; CHECK-NEXT: .size: 8 1231; CHECK-NEXT: .value_kind: hidden_none 1232; CHECK-NEXT: - .offset: 56 1233; CHECK-NEXT: .size: 8 1234; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1235; CHECK: .language: OpenCL C 1236; CHECK-NEXT: .language_version: 1237; CHECK-NEXT: - 2 1238; CHECK-NEXT: - 0 1239; CHECK: .name: test_reqd_wgs_vec_type_hint 1240; CHECK: .reqd_workgroup_size: 1241; CHECK-NEXT: - 1 1242; CHECK-NEXT: - 2 1243; CHECK-NEXT: - 4 1244; CHECK: .symbol: test_reqd_wgs_vec_type_hint.kd 1245; CHECK: .vec_type_hint: int 1246define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a) #0 1247 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1248 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !5 1249 !reqd_work_group_size !6 { 1250 ret void 1251} 1252 1253; CHECK: - .args: 1254; CHECK-NEXT: - .name: a 1255; CHECK-NEXT: .offset: 0 1256; CHECK-NEXT: .size: 4 1257; CHECK-NEXT: .type_name: int 1258; CHECK-NEXT: .value_kind: by_value 1259; CHECK-NEXT: - .offset: 8 1260; CHECK-NEXT: .size: 8 1261; CHECK-NEXT: .value_kind: hidden_global_offset_x 1262; CHECK-NEXT: - .offset: 16 1263; CHECK-NEXT: .size: 8 1264; CHECK-NEXT: .value_kind: hidden_global_offset_y 1265; CHECK-NEXT: - .offset: 24 1266; CHECK-NEXT: .size: 8 1267; CHECK-NEXT: .value_kind: hidden_global_offset_z 1268; CHECK-NEXT: - .offset: 32 1269; CHECK-NEXT: .size: 8 1270; CHECK-NEXT: .value_kind: hidden_printf_buffer 1271; CHECK-NEXT: - .offset: 40 1272; CHECK-NEXT: .size: 8 1273; CHECK-NEXT: .value_kind: hidden_none 1274; CHECK-NEXT: - .offset: 48 1275; CHECK-NEXT: .size: 8 1276; CHECK-NEXT: .value_kind: hidden_none 1277; CHECK-NEXT: - .offset: 56 1278; CHECK-NEXT: .size: 8 1279; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1280; CHECK: .language: OpenCL C 1281; CHECK-NEXT: .language_version: 1282; CHECK-NEXT: - 2 1283; CHECK-NEXT: - 0 1284; CHECK: .name: test_wgs_hint_vec_type_hint 1285; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd 1286; CHECK: .vec_type_hint: uint4 1287; CHECK: .workgroup_size_hint: 1288; CHECK-NEXT: - 8 1289; CHECK-NEXT: - 16 1290; CHECK-NEXT: - 32 1291define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a) #0 1292 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 1293 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !7 1294 !work_group_size_hint !8 { 1295 ret void 1296} 1297 1298; CHECK: - .args: 1299; CHECK-NEXT: - .address_space: global 1300; CHECK-NEXT: .name: a 1301; CHECK-NEXT: .offset: 0 1302; CHECK-NEXT: .size: 8 1303; CHECK-NEXT: .type_name: 'int addrspace(5)* addrspace(5)*' 1304; CHECK-NEXT: .value_kind: global_buffer 1305; CHECK-NEXT: - .offset: 8 1306; CHECK-NEXT: .size: 8 1307; CHECK-NEXT: .value_kind: hidden_global_offset_x 1308; CHECK-NEXT: - .offset: 16 1309; CHECK-NEXT: .size: 8 1310; CHECK-NEXT: .value_kind: hidden_global_offset_y 1311; CHECK-NEXT: - .offset: 24 1312; CHECK-NEXT: .size: 8 1313; CHECK-NEXT: .value_kind: hidden_global_offset_z 1314; CHECK-NEXT: - .offset: 32 1315; CHECK-NEXT: .size: 8 1316; CHECK-NEXT: .value_kind: hidden_printf_buffer 1317; CHECK-NEXT: - .offset: 40 1318; CHECK-NEXT: .size: 8 1319; CHECK-NEXT: .value_kind: hidden_none 1320; CHECK-NEXT: - .offset: 48 1321; CHECK-NEXT: .size: 8 1322; CHECK-NEXT: .value_kind: hidden_none 1323; CHECK-NEXT: - .offset: 56 1324; CHECK-NEXT: .size: 8 1325; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1326; CHECK: .language: OpenCL C 1327; CHECK-NEXT: .language_version: 1328; CHECK-NEXT: - 2 1329; CHECK-NEXT: - 0 1330; CHECK: .name: test_arg_ptr_to_ptr 1331; CHECK: .symbol: test_arg_ptr_to_ptr.kd 1332define amdgpu_kernel void @test_arg_ptr_to_ptr(ptr addrspace(1) %a) #0 1333 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80 1334 !kernel_arg_base_type !80 !kernel_arg_type_qual !4 { 1335 ret void 1336} 1337 1338; CHECK: - .args: 1339; CHECK-NEXT: .name: a 1340; CHECK-NEXT: .offset: 0 1341; CHECK-NEXT: .size: 8 1342; CHECK-NEXT: .type_name: struct B 1343; CHECK-NEXT: .value_kind: by_value 1344; CHECK-NEXT: - .offset: 8 1345; CHECK-NEXT: .size: 8 1346; CHECK-NEXT: .value_kind: hidden_global_offset_x 1347; CHECK-NEXT: - .offset: 16 1348; CHECK-NEXT: .size: 8 1349; CHECK-NEXT: .value_kind: hidden_global_offset_y 1350; CHECK-NEXT: - .offset: 24 1351; CHECK-NEXT: .size: 8 1352; CHECK-NEXT: .value_kind: hidden_global_offset_z 1353; CHECK-NEXT: - .offset: 32 1354; CHECK-NEXT: .size: 8 1355; CHECK-NEXT: .value_kind: hidden_printf_buffer 1356; CHECK-NEXT: - .offset: 40 1357; CHECK-NEXT: .size: 8 1358; CHECK-NEXT: .value_kind: hidden_none 1359; CHECK-NEXT: - .offset: 48 1360; CHECK-NEXT: .size: 8 1361; CHECK-NEXT: .value_kind: hidden_none 1362; CHECK-NEXT: - .offset: 56 1363; CHECK-NEXT: .size: 8 1364; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1365; CHECK: .language: OpenCL C 1366; CHECK-NEXT: .language_version: 1367; CHECK-NEXT: - 2 1368; CHECK-NEXT: - 0 1369; CHECK: .name: test_arg_struct_contains_ptr 1370; CHECK: .symbol: test_arg_struct_contains_ptr.kd 1371define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B %a) #0 1372 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !82 1373 !kernel_arg_base_type !82 !kernel_arg_type_qual !4 { 1374 ret void 1375} 1376 1377; CHECK: - .args: 1378; CHECK-NEXT: - .name: a 1379; CHECK-NEXT: .offset: 0 1380; CHECK-NEXT: .size: 16 1381; CHECK-NEXT: .type_name: 'global int addrspace(5)* __attribute__((ext_vector_type(2)))' 1382; CHECK-NEXT: .value_kind: by_value 1383; CHECK-NEXT: - .offset: 16 1384; CHECK-NEXT: .size: 8 1385; CHECK-NEXT: .value_kind: hidden_global_offset_x 1386; CHECK-NEXT: - .offset: 24 1387; CHECK-NEXT: .size: 8 1388; CHECK-NEXT: .value_kind: hidden_global_offset_y 1389; CHECK-NEXT: - .offset: 32 1390; CHECK-NEXT: .size: 8 1391; CHECK-NEXT: .value_kind: hidden_global_offset_z 1392; CHECK-NEXT: - .offset: 40 1393; CHECK-NEXT: .size: 8 1394; CHECK-NEXT: .value_kind: hidden_printf_buffer 1395; CHECK-NEXT: - .offset: 48 1396; CHECK-NEXT: .size: 8 1397; CHECK-NEXT: .value_kind: hidden_none 1398; CHECK-NEXT: - .offset: 56 1399; CHECK-NEXT: .size: 8 1400; CHECK-NEXT: .value_kind: hidden_none 1401; CHECK-NEXT: - .offset: 64 1402; CHECK-NEXT: .size: 8 1403; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1404; CHECK: .language: OpenCL C 1405; CHECK-NEXT: .language_version: 1406; CHECK-NEXT: - 2 1407; CHECK-NEXT: - 0 1408; CHECK: .name: test_arg_vector_of_ptr 1409; CHECK: .symbol: test_arg_vector_of_ptr.kd 1410define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x ptr addrspace(1)> %a) #0 1411 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83 1412 !kernel_arg_base_type !83 !kernel_arg_type_qual !4 { 1413 ret void 1414} 1415 1416; CHECK: - .args: 1417; CHECK-NEXT: - .address_space: global 1418; CHECK-NEXT: .name: a 1419; CHECK-NEXT: .offset: 0 1420; CHECK-NEXT: .size: 8 1421; CHECK-NEXT: .type_name: clk_event_t 1422; CHECK-NEXT: .value_kind: global_buffer 1423; CHECK-NEXT: - .offset: 8 1424; CHECK-NEXT: .size: 8 1425; CHECK-NEXT: .value_kind: hidden_global_offset_x 1426; CHECK-NEXT: - .offset: 16 1427; CHECK-NEXT: .size: 8 1428; CHECK-NEXT: .value_kind: hidden_global_offset_y 1429; CHECK-NEXT: - .offset: 24 1430; CHECK-NEXT: .size: 8 1431; CHECK-NEXT: .value_kind: hidden_global_offset_z 1432; CHECK-NEXT: - .offset: 32 1433; CHECK-NEXT: .size: 8 1434; CHECK-NEXT: .value_kind: hidden_printf_buffer 1435; CHECK-NEXT: - .offset: 40 1436; CHECK-NEXT: .size: 8 1437; CHECK-NEXT: .value_kind: hidden_none 1438; CHECK-NEXT: - .offset: 48 1439; CHECK-NEXT: .size: 8 1440; CHECK-NEXT: .value_kind: hidden_none 1441; CHECK-NEXT: - .offset: 56 1442; CHECK-NEXT: .size: 8 1443; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1444; CHECK: .language: OpenCL C 1445; CHECK-NEXT: .language_version: 1446; CHECK-NEXT: - 2 1447; CHECK-NEXT: - 0 1448; CHECK: .name: test_arg_unknown_builtin_type 1449; CHECK: .symbol: test_arg_unknown_builtin_type.kd 1450define amdgpu_kernel void @test_arg_unknown_builtin_type( 1451 ptr addrspace(1) %a) #0 1452 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !84 1453 !kernel_arg_base_type !84 !kernel_arg_type_qual !4 { 1454 ret void 1455} 1456 1457; CHECK: - .args: 1458; CHECK-NEXT: - .address_space: global 1459; CHECK-NEXT: .name: a 1460; CHECK-NEXT: .offset: 0 1461; CHECK-NEXT: .size: 8 1462; CHECK-NEXT: .type_name: 'long addrspace(5)*' 1463; CHECK-NEXT: .value_kind: global_buffer 1464; CHECK-NEXT: - .address_space: local 1465; CHECK-NEXT: .name: b 1466; CHECK-NEXT: .offset: 8 1467; CHECK-NEXT: .pointee_align: 1 1468; CHECK-NEXT: .size: 4 1469; CHECK-NEXT: .type_name: 'char addrspace(5)*' 1470; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1471; CHECK-NEXT: - .address_space: local 1472; CHECK-NEXT: .name: c 1473; CHECK-NEXT: .offset: 12 1474; CHECK-NEXT: .pointee_align: 2 1475; CHECK-NEXT: .size: 4 1476; CHECK-NEXT: .type_name: 'char2 addrspace(5)*' 1477; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1478; CHECK-NEXT: - .address_space: local 1479; CHECK-NEXT: .name: d 1480; CHECK-NEXT: .offset: 16 1481; CHECK-NEXT: .pointee_align: 4 1482; CHECK-NEXT: .size: 4 1483; CHECK-NEXT: .type_name: 'char3 addrspace(5)*' 1484; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1485; CHECK-NEXT: - .address_space: local 1486; CHECK-NEXT: .name: e 1487; CHECK-NEXT: .offset: 20 1488; CHECK-NEXT: .pointee_align: 4 1489; CHECK-NEXT: .size: 4 1490; CHECK-NEXT: .type_name: 'char4 addrspace(5)*' 1491; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1492; CHECK-NEXT: - .address_space: local 1493; CHECK-NEXT: .name: f 1494; CHECK-NEXT: .offset: 24 1495; CHECK-NEXT: .pointee_align: 8 1496; CHECK-NEXT: .size: 4 1497; CHECK-NEXT: .type_name: 'char8 addrspace(5)*' 1498; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1499; CHECK-NEXT: - .address_space: local 1500; CHECK-NEXT: .name: g 1501; CHECK-NEXT: .offset: 28 1502; CHECK-NEXT: .pointee_align: 16 1503; CHECK-NEXT: .size: 4 1504; CHECK-NEXT: .type_name: 'char16 addrspace(5)*' 1505; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1506; CHECK-NEXT: - .address_space: local 1507; CHECK-NEXT: .name: h 1508; CHECK-NEXT: .offset: 32 1509; CHECK-NEXT: .pointee_align: 1 1510; CHECK-NEXT: .size: 4 1511; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1512; CHECK-NEXT: - .offset: 40 1513; CHECK-NEXT: .size: 8 1514; CHECK-NEXT: .value_kind: hidden_global_offset_x 1515; CHECK-NEXT: - .offset: 48 1516; CHECK-NEXT: .size: 8 1517; CHECK-NEXT: .value_kind: hidden_global_offset_y 1518; CHECK-NEXT: - .offset: 56 1519; CHECK-NEXT: .size: 8 1520; CHECK-NEXT: .value_kind: hidden_global_offset_z 1521; CHECK-NEXT: - .offset: 64 1522; CHECK-NEXT: .size: 8 1523; CHECK-NEXT: .value_kind: hidden_printf_buffer 1524; CHECK-NEXT: - .offset: 72 1525; CHECK-NEXT: .size: 8 1526; CHECK-NEXT: .value_kind: hidden_none 1527; CHECK-NEXT: - .offset: 80 1528; CHECK-NEXT: .size: 8 1529; CHECK-NEXT: .value_kind: hidden_none 1530; CHECK-NEXT: - .offset: 88 1531; CHECK-NEXT: .size: 8 1532; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1533; CHECK: .language: OpenCL C 1534; CHECK-NEXT: .language_version: 1535; CHECK-NEXT: - 2 1536; CHECK-NEXT: - 0 1537; CHECK: .name: test_pointee_align 1538; CHECK: .symbol: test_pointee_align.kd 1539define amdgpu_kernel void @test_pointee_align(ptr addrspace(1) %a, 1540 ptr addrspace(3) %b, 1541 ptr addrspace(3) align 2 %c, 1542 ptr addrspace(3) align 4 %d, 1543 ptr addrspace(3) align 4 %e, 1544 ptr addrspace(3) align 8 %f, 1545 ptr addrspace(3) align 16 %g, 1546 ptr addrspace(3) %h) #0 1547 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93 1548 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 { 1549 ret void 1550} 1551 1552; CHECK: - .args: 1553; CHECK-NEXT: - .address_space: global 1554; CHECK-NEXT: .name: a 1555; CHECK-NEXT: .offset: 0 1556; CHECK-NEXT: .size: 8 1557; CHECK-NEXT: .type_name: 'long addrspace(5)*' 1558; CHECK-NEXT: .value_kind: global_buffer 1559; CHECK-NEXT: - .address_space: local 1560; CHECK-NEXT: .name: b 1561; CHECK-NEXT: .offset: 8 1562; CHECK-NEXT: .pointee_align: 8 1563; CHECK-NEXT: .size: 4 1564; CHECK-NEXT: .type_name: 'char addrspace(5)*' 1565; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1566; CHECK-NEXT: - .address_space: local 1567; CHECK-NEXT: .name: c 1568; CHECK-NEXT: .offset: 12 1569; CHECK-NEXT: .pointee_align: 32 1570; CHECK-NEXT: .size: 4 1571; CHECK-NEXT: .type_name: 'char2 addrspace(5)*' 1572; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1573; CHECK-NEXT: - .address_space: local 1574; CHECK-NEXT: .name: d 1575; CHECK-NEXT: .offset: 16 1576; CHECK-NEXT: .pointee_align: 64 1577; CHECK-NEXT: .size: 4 1578; CHECK-NEXT: .type_name: 'char3 addrspace(5)*' 1579; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1580; CHECK-NEXT: - .address_space: local 1581; CHECK-NEXT: .name: e 1582; CHECK-NEXT: .offset: 20 1583; CHECK-NEXT: .pointee_align: 256 1584; CHECK-NEXT: .size: 4 1585; CHECK-NEXT: .type_name: 'char4 addrspace(5)*' 1586; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1587; CHECK-NEXT: - .address_space: local 1588; CHECK-NEXT: .name: f 1589; CHECK-NEXT: .offset: 24 1590; CHECK-NEXT: .pointee_align: 128 1591; CHECK-NEXT: .size: 4 1592; CHECK-NEXT: .type_name: 'char8 addrspace(5)*' 1593; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1594; CHECK-NEXT: - .address_space: local 1595; CHECK-NEXT: .name: g 1596; CHECK-NEXT: .offset: 28 1597; CHECK-NEXT: .pointee_align: 1024 1598; CHECK-NEXT: .size: 4 1599; CHECK-NEXT: .type_name: 'char16 addrspace(5)*' 1600; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1601; CHECK-NEXT: - .address_space: local 1602; CHECK-NEXT: .name: h 1603; CHECK-NEXT: .offset: 32 1604; CHECK-NEXT: .pointee_align: 16 1605; CHECK-NEXT: .size: 4 1606; CHECK-NEXT: .value_kind: dynamic_shared_pointer 1607; CHECK-NEXT: - .offset: 40 1608; CHECK-NEXT: .size: 8 1609; CHECK-NEXT: .value_kind: hidden_global_offset_x 1610; CHECK-NEXT: - .offset: 48 1611; CHECK-NEXT: .size: 8 1612; CHECK-NEXT: .value_kind: hidden_global_offset_y 1613; CHECK-NEXT: - .offset: 56 1614; CHECK-NEXT: .size: 8 1615; CHECK-NEXT: .value_kind: hidden_global_offset_z 1616; CHECK-NEXT: - .offset: 64 1617; CHECK-NEXT: .size: 8 1618; CHECK-NEXT: .value_kind: hidden_printf_buffer 1619; CHECK-NEXT: - .offset: 72 1620; CHECK-NEXT: .size: 8 1621; CHECK-NEXT: .value_kind: hidden_none 1622; CHECK-NEXT: - .offset: 80 1623; CHECK-NEXT: .size: 8 1624; CHECK-NEXT: .value_kind: hidden_none 1625; CHECK-NEXT: - .offset: 88 1626; CHECK-NEXT: .size: 8 1627; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1628; CHECK: .language: OpenCL C 1629; CHECK-NEXT: .language_version: 1630; CHECK-NEXT: - 2 1631; CHECK-NEXT: - 0 1632; CHECK: .name: test_pointee_align_attribute 1633; CHECK: .symbol: test_pointee_align_attribute.kd 1634define amdgpu_kernel void @test_pointee_align_attribute(ptr addrspace(1) align 16 %a, 1635 ptr addrspace(3) align 8 %b, 1636 ptr addrspace(3) align 32 %c, 1637 ptr addrspace(3) align 64 %d, 1638 ptr addrspace(3) align 256 %e, 1639 ptr addrspace(3) align 128 %f, 1640 ptr addrspace(3) align 1024 %g, 1641 ptr addrspace(3) align 16 %h) #0 1642 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !93 1643 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 { 1644 ret void 1645} 1646; CHECK: - .args: 1647; CHECK-NEXT: - .name: arg 1648; CHECK-NEXT: .offset: 0 1649; CHECK-NEXT: .size: 25 1650; CHECK-NEXT: .type_name: __block_literal 1651; CHECK-NEXT: .value_kind: by_value 1652; CHECK-NEXT: - .offset: 32 1653; CHECK-NEXT: .size: 8 1654; CHECK-NEXT: .value_kind: hidden_global_offset_x 1655; CHECK-NEXT: - .offset: 40 1656; CHECK-NEXT: .size: 8 1657; CHECK-NEXT: .value_kind: hidden_global_offset_y 1658; CHECK-NEXT: - .offset: 48 1659; CHECK-NEXT: .size: 8 1660; CHECK-NEXT: .value_kind: hidden_global_offset_z 1661; CHECK-NEXT: - .offset: 56 1662; CHECK-NEXT: .size: 8 1663; CHECK-NEXT: .value_kind: hidden_printf_buffer 1664; CHECK-NEXT: - .offset: 64 1665; CHECK-NEXT: .size: 8 1666; CHECK-NEXT: .value_kind: hidden_none 1667; CHECK-NEXT: - .offset: 72 1668; CHECK-NEXT: .size: 8 1669; CHECK-NEXT: .value_kind: hidden_none 1670; CHECK-NEXT: - .offset: 80 1671; CHECK-NEXT: .size: 8 1672; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1673; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle 1674; CHECK: .language: OpenCL C 1675; CHECK-NEXT: .language_version: 1676; CHECK-NEXT: - 2 1677; CHECK-NEXT: - 0 1678; CHECK: .name: __test_block_invoke_kernel 1679; CHECK: .symbol: __test_block_invoke_kernel.kd 1680define amdgpu_kernel void @__test_block_invoke_kernel( 1681 <{ i32, i32, ptr, ptr addrspace(1), i8 }> %arg) #1 1682 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110 1683 !kernel_arg_base_type !110 !kernel_arg_type_qual !4 { 1684 ret void 1685} 1686 1687; CHECK: - .args: 1688; CHECK-NEXT: - .name: a 1689; CHECK-NEXT: .offset: 0 1690; CHECK-NEXT: .size: 1 1691; CHECK-NEXT: .type_name: char 1692; CHECK-NEXT: .value_kind: by_value 1693; CHECK-NEXT: - .offset: 8 1694; CHECK-NEXT: .size: 8 1695; CHECK-NEXT: .value_kind: hidden_global_offset_x 1696; CHECK-NEXT: - .offset: 16 1697; CHECK-NEXT: .size: 8 1698; CHECK-NEXT: .value_kind: hidden_global_offset_y 1699; CHECK-NEXT: - .offset: 24 1700; CHECK-NEXT: .size: 8 1701; CHECK-NEXT: .value_kind: hidden_global_offset_z 1702; CHECK-NEXT: - .offset: 32 1703; CHECK-NEXT: .size: 8 1704; CHECK-NEXT: .value_kind: hidden_printf_buffer 1705; CHECK-NEXT: - .offset: 40 1706; CHECK-NEXT: .size: 8 1707; CHECK-NEXT: .value_kind: hidden_default_queue 1708; CHECK-NEXT: - .offset: 48 1709; CHECK-NEXT: .size: 8 1710; CHECK-NEXT: .value_kind: hidden_completion_action 1711; CHECK-NEXT: - .offset: 56 1712; CHECK-NEXT: .size: 8 1713; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg 1714; CHECK: .language: OpenCL C 1715; CHECK-NEXT: .language_version: 1716; CHECK-NEXT: - 2 1717; CHECK-NEXT: - 0 1718; CHECK: .name: test_enqueue_kernel_caller 1719; CHECK: .symbol: test_enqueue_kernel_caller.kd 1720define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2 1721 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9 1722 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 { 1723 ret void 1724} 1725 1726; CHECK: - .args: 1727; CHECK-NEXT: - .name: ptr 1728; CHECK-NEXT: .offset: 0 1729; CHECK-NEXT: .size: 8 1730; CHECK-NEXT: .value_kind: global_buffer 1731; CHECK: .name: unknown_addrspace_kernarg 1732; CHECK: .symbol: unknown_addrspace_kernarg.kd 1733define amdgpu_kernel void @unknown_addrspace_kernarg(ptr addrspace(12345) %ptr) #0 { 1734 ret void 1735} 1736 1737; CHECK: amdhsa.printf: 1738; CHECK-NEXT: - '1:1:4:%d\n' 1739; CHECK-NEXT: - '2:1:8:%g\n' 1740; CHECK: amdhsa.version: 1741; CHECK-NEXT: - 1 1742; CHECK-NEXT: - 1 1743 1744attributes #0 = { optnone noinline "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-implicitarg-num-bytes"="56" } 1745attributes #1 = { optnone noinline "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-implicitarg-num-bytes"="56" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" } 1746attributes #2 = { optnone noinline "amdgpu-implicitarg-num-bytes"="56" } 1747 1748!llvm.module.flags = !{!0} 1749!0 = !{i32 1, !"amdhsa_code_object_version", i32 400} 1750 1751!llvm.printf.fmts = !{!100, !101} 1752 1753!1 = !{i32 0} 1754!2 = !{!"none"} 1755!3 = !{!"int"} 1756!4 = !{!""} 1757!5 = !{i32 undef, i32 1} 1758!6 = !{i32 1, i32 2, i32 4} 1759!7 = !{<4 x i32> undef, i32 0} 1760!8 = !{i32 8, i32 16, i32 32} 1761!9 = !{!"char"} 1762!10 = !{!"ushort2"} 1763!11 = !{!"int3"} 1764!12 = !{!"ulong4"} 1765!13 = !{!"half8"} 1766!14 = !{!"float16"} 1767!15 = !{!"double16"} 1768!16 = !{!"int addrspace(5)*"} 1769!17 = !{!"image2d_t"} 1770!18 = !{!"sampler_t"} 1771!19 = !{!"queue_t"} 1772!20 = !{!"struct A"} 1773!21 = !{!"i128"} 1774!22 = !{i32 0, i32 0, i32 0} 1775!23 = !{!"none", !"none", !"none"} 1776!24 = !{!"int", !"short2", !"char3"} 1777!25 = !{!"", !"", !""} 1778!26 = !{half undef, i32 1} 1779!27 = !{float undef, i32 1} 1780!28 = !{double undef, i32 1} 1781!29 = !{i8 undef, i32 1} 1782!30 = !{i16 undef, i32 1} 1783!31 = !{i64 undef, i32 1} 1784!32 = !{ptr addrspace(5) undef, i32 1} 1785!50 = !{i32 1, i32 2, i32 3} 1786!51 = !{!"int addrspace(5)*", !"int addrspace(5)*", !"int addrspace(5)*"} 1787!60 = !{i32 1, i32 1, i32 1} 1788!61 = !{!"read_only", !"write_only", !"read_write"} 1789!62 = !{!"image1d_t", !"image2d_t", !"image3d_t"} 1790!70 = !{!"volatile", !"const restrict", !"pipe"} 1791!80 = !{!"int addrspace(5)* addrspace(5)*"} 1792!81 = !{i32 1} 1793!82 = !{!"struct B"} 1794!83 = !{!"global int addrspace(5)* __attribute__((ext_vector_type(2)))"} 1795!84 = !{!"clk_event_t"} 1796!opencl.ocl.version = !{!90} 1797!90 = !{i32 2, i32 0} 1798!91 = !{i32 0, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3} 1799!92 = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none"} 1800!93 = !{!"long addrspace(5)*", !"char addrspace(5)*", !"char2 addrspace(5)*", !"char3 addrspace(5)*", !"char4 addrspace(5)*", !"char8 addrspace(5)*", !"char16 addrspace(5)*"} 1801!94 = !{!"", !"", !"", !"", !"", !"", !""} 1802!100 = !{!"1:1:4:%d\5Cn"} 1803!101 = !{!"2:1:8:%g\5Cn"} 1804!110 = !{!"__block_literal"} 1805!111 = !{!"char", !"char"} 1806 1807; PARSER: AMDGPU HSA Metadata Parser Test: PASS 1808