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