1835 lines · plain
1; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx700 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s2; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx802 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s3; RUN: llc -mtriple=amdgcn-amd-amdhsa -mcpu=gfx900 -filetype=obj -o - < %s | llvm-readelf --notes - | FileCheck --check-prefix=CHECK %s4; 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 %s5; 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 %s6; 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 %s7 8%struct.A = type { i8, float }9%opencl.image1d_t = type opaque10%opencl.image2d_t = type opaque11%opencl.image3d_t = type opaque12%opencl.queue_t = type opaque13%opencl.pipe_t = type opaque14%struct.B = type { ptr addrspace(1) }15%opencl.clk_event_t = type opaque16 17@__test_block_invoke_kernel_runtime_handle = external addrspace(1) externally_initialized constant ptr addrspace(1), section ".amdgpu.kernel.runtime.handle"18@not.a.handle = external addrspace(1) externally_initialized constant ptr addrspace(1)19 20; CHECK: ---21; CHECK-NEXT: amdhsa.kernels:22; CHECK-NEXT: - .args:23; CHECK-NEXT: - .name: a24; CHECK-NEXT: .offset: 025; CHECK-NEXT: .size: 126; CHECK-NEXT: .type_name: char27; CHECK-NEXT: .value_kind: by_value28; CHECK-NEXT: - .offset: 829; CHECK-NEXT: .size: 830; CHECK-NEXT: .value_kind: hidden_global_offset_x31; CHECK-NEXT: - .offset: 1632; CHECK-NEXT: .size: 833; CHECK-NEXT: .value_kind: hidden_global_offset_y34; CHECK-NEXT: - .offset: 2435; CHECK-NEXT: .size: 836; CHECK-NEXT: .value_kind: hidden_global_offset_z37; CHECK-NEXT: - .offset: 3238; CHECK-NEXT: .size: 839; CHECK-NOT: .value_kind: hidden_default_queue40; CHECK-NOT: .value_kind: hidden_completion_action41; CHECK-NOT: .value_kind: hidden_hostcall_buffer42; CHECK-NEXT: .value_kind: hidden_printf_buffer43; CHECK: .value_kind: hidden_multigrid_sync_arg44; CHECK: .language: OpenCL C45; CHECK-NEXT: .language_version:46; CHECK-NEXT: - 247; CHECK-NEXT: - 048; CHECK: .name: test_char49; CHECK: .symbol: test_char.kd50define amdgpu_kernel void @test_char(i8 %a) #051 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !952 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {53 ret void54}55 56; CHECK: - .args:57; CHECK-NEXT: - .name: a58; CHECK-NEXT: .offset: 059; CHECK-NEXT: .size: 160; CHECK-NEXT: .type_name: char61; CHECK-NEXT: .value_kind: by_value62; CHECK-NEXT: - .offset: 863; CHECK-NEXT: .size: 864; CHECK-NEXT: .value_kind: hidden_global_offset_x65; CHECK-NEXT: - .offset: 1666; CHECK-NEXT: .size: 867; CHECK-NEXT: .value_kind: hidden_global_offset_y68; CHECK-NEXT: - .offset: 2469; CHECK-NEXT: .size: 870; CHECK-NEXT: .value_kind: hidden_global_offset_z71; CHECK-NEXT: - .offset: 3272; CHECK-NEXT: .size: 873; CHECK-NOT: .value_kind: hidden_default_queue74; CHECK-NOT: .value_kind: hidden_completion_action75; CHECK-NOT: .value_kind: hidden_hostcall_buffer76; CHECK-NEXT: .value_kind: hidden_printf_buffer77; CHECK: .value_kind: hidden_multigrid_sync_arg78; CHECK: .language: OpenCL C79; CHECK-NEXT: .language_version:80; CHECK-NEXT: - 281; CHECK-NEXT: - 082; CHECK: .name: test_char_byref_constant83; CHECK: .symbol: test_char_byref_constant.kd84define amdgpu_kernel void @test_char_byref_constant(ptr addrspace(4) byref(i8) %a) #085 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !986 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {87 ret void88}89 90; CHECK: - .args:91; CHECK-NEXT: - .offset: 092; CHECK-NEXT: .size: 193; CHECK-NEXT: .type_name: char94; CHECK-NEXT: .value_kind: by_value95; CHECK-NEXT: - .name: a96; CHECK-NEXT: .offset: 51297; CHECK-NEXT: .size: 198; CHECK-NEXT: .type_name: char99; CHECK-NEXT: .value_kind: by_value100; CHECK-NEXT: - .offset: 520101; CHECK-NEXT: .size: 8102; CHECK-NEXT: .value_kind: hidden_global_offset_x103; CHECK-NEXT: - .offset: 528104; CHECK-NEXT: .size: 8105; CHECK-NEXT: .value_kind: hidden_global_offset_y106; CHECK-NEXT: - .offset: 536107; CHECK-NEXT: .size: 8108; CHECK-NEXT: .value_kind: hidden_global_offset_z109; CHECK-NEXT: - .offset: 544110; CHECK-NEXT: .size: 8111; CHECK-NOT: .value_kind: hidden_default_queue112; CHECK-NOT: .value_kind: hidden_completion_action113; CHECK-NOT: .value_kind: hidden_hostcall_buffer114; CHECK-NEXT: .value_kind: hidden_printf_buffer115; CHECK: .value_kind: hidden_multigrid_sync_arg116; CHECK: .language: OpenCL C117; CHECK-NEXT: .language_version:118; CHECK-NEXT: - 2119; CHECK-NEXT: - 0120; CHECK: .name: test_char_byref_constant_align512121; CHECK: .symbol: test_char_byref_constant_align512.kd122define amdgpu_kernel void @test_char_byref_constant_align512(i8, ptr addrspace(4) byref(i8) align(512) %a) #0123 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !111124 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {125 ret void126}127 128; CHECK: - .args:129; CHECK-NEXT: - .name: a130; CHECK-NEXT: .offset: 0131; CHECK-NEXT: .size: 4132; CHECK-NEXT: .type_name: ushort2133; CHECK-NEXT: .value_kind: by_value134; CHECK-NEXT: - .offset: 8135; CHECK-NEXT: .size: 8136; CHECK-NEXT: .value_kind: hidden_global_offset_x137; CHECK-NEXT: - .offset: 16138; CHECK-NEXT: .size: 8139; CHECK-NEXT: .value_kind: hidden_global_offset_y140; CHECK-NEXT: - .offset: 24141; CHECK-NEXT: .size: 8142; CHECK-NEXT: .value_kind: hidden_global_offset_z143; CHECK-NEXT: - .offset: 32144; CHECK-NEXT: .size: 8145; CHECK-NEXT: .value_kind: hidden_printf_buffer146; CHECK-NEXT: - .offset: 40147; CHECK-NEXT: .size: 8148; CHECK-NEXT: .value_kind: hidden_none149; CHECK-NEXT: - .offset: 48150; CHECK-NEXT: .size: 8151; CHECK-NEXT: .value_kind: hidden_none152; CHECK-NEXT: - .offset: 56153; CHECK-NEXT: .size: 8154; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg155; CHECK: .language: OpenCL C156; CHECK-NEXT: .language_version:157; CHECK-NEXT: - 2158; CHECK-NEXT: - 0159; CHECK: .name: test_ushort2160; CHECK: .symbol: test_ushort2.kd161define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0162 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10163 !kernel_arg_base_type !10 !kernel_arg_type_qual !4 {164 ret void165}166 167; CHECK: - .args:168; CHECK-NEXT: - .name: a169; CHECK-NEXT: .offset: 0170; CHECK-NEXT: .size: 16171; CHECK-NEXT: .type_name: int3172; CHECK-NEXT: .value_kind: by_value173; CHECK-NEXT: - .offset: 16174; CHECK-NEXT: .size: 8175; CHECK-NEXT: .value_kind: hidden_global_offset_x176; CHECK-NEXT: - .offset: 24177; CHECK-NEXT: .size: 8178; CHECK-NEXT: .value_kind: hidden_global_offset_y179; CHECK-NEXT: - .offset: 32180; CHECK-NEXT: .size: 8181; CHECK-NEXT: .value_kind: hidden_global_offset_z182; CHECK-NEXT: - .offset: 40183; CHECK-NEXT: .size: 8184; CHECK-NEXT: .value_kind: hidden_printf_buffer185; CHECK-NEXT: - .offset: 48186; CHECK-NEXT: .size: 8187; CHECK-NEXT: .value_kind: hidden_none188; CHECK-NEXT: - .offset: 56189; CHECK-NEXT: .size: 8190; CHECK-NEXT: .value_kind: hidden_none191; CHECK-NEXT: - .offset: 64192; CHECK-NEXT: .size: 8193; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg194; CHECK: .language: OpenCL C195; CHECK-NEXT: .language_version:196; CHECK-NEXT: - 2197; CHECK-NEXT: - 0198; CHECK: .name: test_int3199; CHECK: .symbol: test_int3.kd200define amdgpu_kernel void @test_int3(<3 x i32> %a) #0201 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11202 !kernel_arg_base_type !11 !kernel_arg_type_qual !4 {203 ret void204}205 206; CHECK: - .args:207; CHECK-NEXT: - .name: a208; CHECK-NEXT: .offset: 0209; CHECK-NEXT: .size: 32210; CHECK-NEXT: .type_name: ulong4211; CHECK-NEXT: .value_kind: by_value212; CHECK-NEXT: - .offset: 32213; CHECK-NEXT: .size: 8214; CHECK-NEXT: .value_kind: hidden_global_offset_x215; CHECK-NEXT: - .offset: 40216; CHECK-NEXT: .size: 8217; CHECK-NEXT: .value_kind: hidden_global_offset_y218; CHECK-NEXT: - .offset: 48219; CHECK-NEXT: .size: 8220; CHECK-NEXT: .value_kind: hidden_global_offset_z221; CHECK-NEXT: - .offset: 56222; CHECK-NEXT: .size: 8223; CHECK-NEXT: .value_kind: hidden_printf_buffer224; CHECK-NEXT: - .offset: 64225; CHECK-NEXT: .size: 8226; CHECK-NEXT: .value_kind: hidden_none227; CHECK-NEXT: - .offset: 72228; CHECK-NEXT: .size: 8229; CHECK-NEXT: .value_kind: hidden_none230; CHECK-NEXT: - .offset: 80231; CHECK-NEXT: .size: 8232; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg233; CHECK: .language: OpenCL C234; CHECK-NEXT: .language_version:235; CHECK-NEXT: - 2236; CHECK-NEXT: - 0237; CHECK: .name: test_ulong4238; CHECK: .symbol: test_ulong4.kd239define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0240 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12241 !kernel_arg_base_type !12 !kernel_arg_type_qual !4 {242 ret void243}244 245; CHECK: - .args:246; CHECK-NEXT: - .name: a247; CHECK-NEXT: .offset: 0248; CHECK-NEXT: .size: 16249; CHECK-NEXT: .type_name: half8250; CHECK-NEXT: .value_kind: by_value251; CHECK-NEXT: - .offset: 16252; CHECK-NEXT: .size: 8253; CHECK-NEXT: .value_kind: hidden_global_offset_x254; CHECK-NEXT: - .offset: 24255; CHECK-NEXT: .size: 8256; CHECK-NEXT: .value_kind: hidden_global_offset_y257; CHECK-NEXT: - .offset: 32258; CHECK-NEXT: .size: 8259; CHECK-NEXT: .value_kind: hidden_global_offset_z260; CHECK-NEXT: - .offset: 40261; CHECK-NEXT: .size: 8262; CHECK-NEXT: .value_kind: hidden_printf_buffer263; CHECK-NEXT: - .offset: 48264; CHECK-NEXT: .size: 8265; CHECK-NEXT: .value_kind: hidden_none266; CHECK-NEXT: - .offset: 56267; CHECK-NEXT: .size: 8268; CHECK-NEXT: .value_kind: hidden_none269; CHECK-NEXT: - .offset: 64270; CHECK-NEXT: .size: 8271; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg272; CHECK: .language: OpenCL C273; CHECK-NEXT: .language_version:274; CHECK-NEXT: - 2275; CHECK-NEXT: - 0276; CHECK: .name: test_half8277; CHECK: .symbol: test_half8.kd278define amdgpu_kernel void @test_half8(<8 x half> %a) #0279 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13280 !kernel_arg_base_type !13 !kernel_arg_type_qual !4 {281 ret void282}283 284; CHECK: - .args:285; CHECK-NEXT: - .name: a286; CHECK-NEXT: .offset: 0287; CHECK-NEXT: .size: 64288; CHECK-NEXT: .type_name: float16289; CHECK-NEXT: .value_kind: by_value290; CHECK-NEXT: - .offset: 64291; CHECK-NEXT: .size: 8292; CHECK-NEXT: .value_kind: hidden_global_offset_x293; CHECK-NEXT: - .offset: 72294; CHECK-NEXT: .size: 8295; CHECK-NEXT: .value_kind: hidden_global_offset_y296; CHECK-NEXT: - .offset: 80297; CHECK-NEXT: .size: 8298; CHECK-NEXT: .value_kind: hidden_global_offset_z299; CHECK-NEXT: - .offset: 88300; CHECK-NEXT: .size: 8301; CHECK-NEXT: .value_kind: hidden_printf_buffer302; CHECK-NEXT: - .offset: 96303; CHECK-NEXT: .size: 8304; CHECK-NEXT: .value_kind: hidden_none305; CHECK-NEXT: - .offset: 104306; CHECK-NEXT: .size: 8307; CHECK-NEXT: .value_kind: hidden_none308; CHECK-NEXT: - .offset: 112309; CHECK-NEXT: .size: 8310; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg311; CHECK: .language: OpenCL C312; CHECK-NEXT: .language_version:313; CHECK-NEXT: - 2314; CHECK-NEXT: - 0315; CHECK: .name: test_float16316; CHECK: .symbol: test_float16.kd317define amdgpu_kernel void @test_float16(<16 x float> %a) #0318 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14319 !kernel_arg_base_type !14 !kernel_arg_type_qual !4 {320 ret void321}322 323; CHECK: - .args:324; CHECK-NEXT: - .name: a325; CHECK-NEXT: .offset: 0326; CHECK-NEXT: .size: 128327; CHECK-NEXT: .type_name: double16328; CHECK-NEXT: .value_kind: by_value329; CHECK-NEXT: - .offset: 128330; CHECK-NEXT: .size: 8331; CHECK-NEXT: .value_kind: hidden_global_offset_x332; CHECK-NEXT: - .offset: 136333; CHECK-NEXT: .size: 8334; CHECK-NEXT: .value_kind: hidden_global_offset_y335; CHECK-NEXT: - .offset: 144336; CHECK-NEXT: .size: 8337; CHECK-NEXT: .value_kind: hidden_global_offset_z338; CHECK-NEXT: - .offset: 152339; CHECK-NEXT: .size: 8340; CHECK-NEXT: .value_kind: hidden_printf_buffer341; CHECK-NEXT: - .offset: 160342; CHECK-NEXT: .size: 8343; CHECK-NEXT: .value_kind: hidden_none344; CHECK-NEXT: - .offset: 168345; CHECK-NEXT: .size: 8346; CHECK-NEXT: .value_kind: hidden_none347; CHECK-NEXT: - .offset: 176348; CHECK-NEXT: .size: 8349; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg350; CHECK: .language: OpenCL C351; CHECK-NEXT: .language_version:352; CHECK-NEXT: - 2353; CHECK-NEXT: - 0354; CHECK: .name: test_double16355; CHECK: .symbol: test_double16.kd356define amdgpu_kernel void @test_double16(<16 x double> %a) #0357 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15358 !kernel_arg_base_type !15 !kernel_arg_type_qual !4 {359 ret void360}361 362; CHECK: - .args:363; CHECK-NEXT: - .address_space: global364; CHECK-NEXT: .name: a365; CHECK-NEXT: .offset: 0366; CHECK-NEXT: .size: 8367; CHECK-NEXT: .type_name: 'int addrspace(5)*'368; CHECK-NEXT: .value_kind: global_buffer369; CHECK-NEXT: - .offset: 8370; CHECK-NEXT: .size: 8371; CHECK-NEXT: .value_kind: hidden_global_offset_x372; CHECK-NEXT: - .offset: 16373; CHECK-NEXT: .size: 8374; CHECK-NEXT: .value_kind: hidden_global_offset_y375; CHECK-NEXT: - .offset: 24376; CHECK-NEXT: .size: 8377; CHECK-NEXT: .value_kind: hidden_global_offset_z378; CHECK-NEXT: - .offset: 32379; CHECK-NEXT: .size: 8380; CHECK-NEXT: .value_kind: hidden_printf_buffer381; CHECK-NEXT: - .offset: 40382; CHECK-NEXT: .size: 8383; CHECK-NEXT: .value_kind: hidden_none384; CHECK-NEXT: - .offset: 48385; CHECK-NEXT: .size: 8386; CHECK-NEXT: .value_kind: hidden_none387; CHECK-NEXT: - .offset: 56388; CHECK-NEXT: .size: 8389; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg390; CHECK: .language: OpenCL C391; CHECK-NEXT: .language_version:392; CHECK-NEXT: - 2393; CHECK-NEXT: - 0394; CHECK: .name: test_pointer395; CHECK: .symbol: test_pointer.kd396define amdgpu_kernel void @test_pointer(ptr addrspace(1) %a) #0397 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16398 !kernel_arg_base_type !16 !kernel_arg_type_qual !4 {399 ret void400}401 402; CHECK: - .args:403; CHECK-NEXT: - .name: a404; CHECK-NEXT: .offset: 0405; CHECK-NEXT: .size: 8406; CHECK-NEXT: .type_name: image2d_t407; CHECK-NEXT: .value_kind: image408; CHECK-NEXT: - .offset: 8409; CHECK-NEXT: .size: 8410; CHECK-NEXT: .value_kind: hidden_global_offset_x411; CHECK-NEXT: - .offset: 16412; CHECK-NEXT: .size: 8413; CHECK-NEXT: .value_kind: hidden_global_offset_y414; CHECK-NEXT: - .offset: 24415; CHECK-NEXT: .size: 8416; CHECK-NEXT: .value_kind: hidden_global_offset_z417; CHECK-NEXT: - .offset: 32418; CHECK-NEXT: .size: 8419; CHECK-NEXT: .value_kind: hidden_printf_buffer420; CHECK-NEXT: - .offset: 40421; CHECK-NEXT: .size: 8422; CHECK-NEXT: .value_kind: hidden_none423; CHECK-NEXT: - .offset: 48424; CHECK-NEXT: .size: 8425; CHECK-NEXT: .value_kind: hidden_none426; CHECK-NEXT: - .offset: 56427; CHECK-NEXT: .size: 8428; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg429; CHECK: .language: OpenCL C430; CHECK-NEXT: .language_version:431; CHECK-NEXT: - 2432; CHECK-NEXT: - 0433; CHECK: .name: test_image434; CHECK: .symbol: test_image.kd435define amdgpu_kernel void @test_image(ptr addrspace(1) %a) #0436 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17437 !kernel_arg_base_type !17 !kernel_arg_type_qual !4 {438 ret void439}440 441; CHECK: - .args:442; CHECK-NEXT: - .name: a443; CHECK-NEXT: .offset: 0444; CHECK-NEXT: .size: 4445; CHECK-NEXT: .type_name: sampler_t446; CHECK-NEXT: .value_kind: sampler447; CHECK-NEXT: - .offset: 8448; CHECK-NEXT: .size: 8449; CHECK-NEXT: .value_kind: hidden_global_offset_x450; CHECK-NEXT: - .offset: 16451; CHECK-NEXT: .size: 8452; CHECK-NEXT: .value_kind: hidden_global_offset_y453; CHECK-NEXT: - .offset: 24454; CHECK-NEXT: .size: 8455; CHECK-NEXT: .value_kind: hidden_global_offset_z456; CHECK-NEXT: - .offset: 32457; CHECK-NEXT: .size: 8458; CHECK-NEXT: .value_kind: hidden_printf_buffer459; CHECK-NEXT: - .offset: 40460; CHECK-NEXT: .size: 8461; CHECK-NEXT: .value_kind: hidden_none462; CHECK-NEXT: - .offset: 48463; CHECK-NEXT: .size: 8464; CHECK-NEXT: .value_kind: hidden_none465; CHECK-NEXT: - .offset: 56466; CHECK-NEXT: .size: 8467; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg468; CHECK: .language: OpenCL C469; CHECK-NEXT: .language_version:470; CHECK-NEXT: - 2471; CHECK-NEXT: - 0472; CHECK: .name: test_sampler473; CHECK: .symbol: test_sampler.kd474define amdgpu_kernel void @test_sampler(i32 %a) #0475 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18476 !kernel_arg_base_type !18 !kernel_arg_type_qual !4 {477 ret void478}479 480; CHECK: - .args:481; CHECK-NEXT: - .name: a482; CHECK-NEXT: .offset: 0483; CHECK-NEXT: .size: 8484; CHECK-NEXT: .type_name: queue_t485; CHECK-NEXT: .value_kind: queue486; CHECK-NEXT: - .offset: 8487; CHECK-NEXT: .size: 8488; CHECK-NEXT: .value_kind: hidden_global_offset_x489; CHECK-NEXT: - .offset: 16490; CHECK-NEXT: .size: 8491; CHECK-NEXT: .value_kind: hidden_global_offset_y492; CHECK-NEXT: - .offset: 24493; CHECK-NEXT: .size: 8494; CHECK-NEXT: .value_kind: hidden_global_offset_z495; CHECK-NEXT: - .offset: 32496; CHECK-NEXT: .size: 8497; CHECK-NEXT: .value_kind: hidden_printf_buffer498; CHECK-NEXT: - .offset: 40499; CHECK-NEXT: .size: 8500; CHECK-NEXT: .value_kind: hidden_none501; CHECK-NEXT: - .offset: 48502; CHECK-NEXT: .size: 8503; CHECK-NEXT: .value_kind: hidden_none504; CHECK-NEXT: - .offset: 56505; CHECK-NEXT: .size: 8506; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg507; CHECK: .language: OpenCL C508; CHECK-NEXT: .language_version:509; CHECK-NEXT: - 2510; CHECK-NEXT: - 0511; CHECK: .name: test_queue512; CHECK: .symbol: test_queue.kd513define amdgpu_kernel void @test_queue(ptr addrspace(1) %a) #0514 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19515 !kernel_arg_base_type !19 !kernel_arg_type_qual !4 {516 ret void517}518 519; CHECK: - .args:520; CHECK-NEXT: .name: a521; CHECK-NEXT: .offset: 0522; CHECK-NEXT: .size: 8523; CHECK-NEXT: .type_name: struct A524; CHECK-NEXT: .value_kind: by_value525; CHECK-NEXT: - .offset: 8526; CHECK-NEXT: .size: 8527; CHECK-NEXT: .value_kind: hidden_global_offset_x528; CHECK-NEXT: - .offset: 16529; CHECK-NEXT: .size: 8530; CHECK-NEXT: .value_kind: hidden_global_offset_y531; CHECK-NEXT: - .offset: 24532; CHECK-NEXT: .size: 8533; CHECK-NEXT: .value_kind: hidden_global_offset_z534; CHECK-NEXT: - .offset: 32535; CHECK-NEXT: .size: 8536; CHECK-NEXT: .value_kind: hidden_printf_buffer537; CHECK-NEXT: - .offset: 40538; CHECK-NEXT: .size: 8539; CHECK-NEXT: .value_kind: hidden_none540; CHECK-NEXT: - .offset: 48541; CHECK-NEXT: .size: 8542; CHECK-NEXT: .value_kind: hidden_none543; CHECK-NEXT: - .offset: 56544; CHECK-NEXT: .size: 8545; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg546; CHECK: .language: OpenCL C547; CHECK-NEXT: .language_version:548; CHECK-NEXT: - 2549; CHECK-NEXT: - 0550; CHECK: .name: test_struct551; CHECK: .symbol: test_struct.kd552define amdgpu_kernel void @test_struct(%struct.A %a) #0553 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20554 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {555 ret void556}557 558; CHECK: - .args:559; CHECK-NEXT: .name: a560; CHECK-NEXT: .offset: 0561; CHECK-NEXT: .size: 8562; CHECK-NEXT: .type_name: struct A563; CHECK-NEXT: .value_kind: by_value564; CHECK-NEXT: - .offset: 8565; CHECK-NEXT: .size: 8566; CHECK-NEXT: .value_kind: hidden_global_offset_x567; CHECK-NEXT: - .offset: 16568; CHECK-NEXT: .size: 8569; CHECK-NEXT: .value_kind: hidden_global_offset_y570; CHECK-NEXT: - .offset: 24571; CHECK-NEXT: .size: 8572; CHECK-NEXT: .value_kind: hidden_global_offset_z573; CHECK-NEXT: - .offset: 32574; CHECK-NEXT: .size: 8575; CHECK-NEXT: .value_kind: hidden_printf_buffer576; CHECK-NEXT: - .offset: 40577; CHECK-NEXT: .size: 8578; CHECK-NEXT: .value_kind: hidden_none579; CHECK-NEXT: - .offset: 48580; CHECK-NEXT: .size: 8581; CHECK-NEXT: .value_kind: hidden_none582; CHECK-NEXT: - .offset: 56583; CHECK-NEXT: .size: 8584; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg585; CHECK: .language: OpenCL C586; CHECK-NEXT: .language_version:587; CHECK-NEXT: - 2588; CHECK-NEXT: - 0589; CHECK: .name: test_struct_byref_constant590; CHECK: .symbol: test_struct_byref_constant.kd591define amdgpu_kernel void @test_struct_byref_constant(ptr addrspace(4) byref(%struct.A) %a) #0592 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20593 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {594 ret void595}596 597; CHECK: - .args:598; CHECK-NEXT: .name: a599; CHECK-NEXT: .offset: 0600; CHECK-NEXT: .size: 32601; CHECK-NEXT: .type_name: struct A602; CHECK-NEXT: .value_kind: by_value603; CHECK-NEXT: - .offset: 32604; CHECK-NEXT: .size: 8605; CHECK-NEXT: .value_kind: hidden_global_offset_x606; CHECK-NEXT: - .offset: 40607; CHECK-NEXT: .size: 8608; CHECK-NEXT: .value_kind: hidden_global_offset_y609; CHECK-NEXT: - .offset: 48610; CHECK-NEXT: .size: 8611; CHECK-NEXT: .value_kind: hidden_global_offset_z612; CHECK-NEXT: - .offset: 56613; CHECK-NEXT: .size: 8614; CHECK-NEXT: .value_kind: hidden_printf_buffer615; CHECK-NEXT: - .offset: 64616; CHECK-NEXT: .size: 8617; CHECK-NEXT: .value_kind: hidden_none618; CHECK-NEXT: - .offset: 72619; CHECK-NEXT: .size: 8620; CHECK-NEXT: .value_kind: hidden_none621; CHECK-NEXT: - .offset: 80622; CHECK-NEXT: .size: 8623; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg624; CHECK: .language: OpenCL C625; CHECK-NEXT: .language_version:626; CHECK-NEXT: - 2627; CHECK-NEXT: - 0628; CHECK: .name: test_array629; CHECK: .symbol: test_array.kd630define amdgpu_kernel void @test_array([32 x i8] %a) #0631 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20632 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {633 ret void634}635 636; CHECK: - .args:637; CHECK-NEXT: .name: a638; CHECK-NEXT: .offset: 0639; CHECK-NEXT: .size: 32640; CHECK-NEXT: .type_name: struct A641; CHECK-NEXT: .value_kind: by_value642; CHECK-NEXT: - .offset: 32643; CHECK-NEXT: .size: 8644; CHECK-NEXT: .value_kind: hidden_global_offset_x645; CHECK-NEXT: - .offset: 40646; CHECK-NEXT: .size: 8647; CHECK-NEXT: .value_kind: hidden_global_offset_y648; CHECK-NEXT: - .offset: 48649; CHECK-NEXT: .size: 8650; CHECK-NEXT: .value_kind: hidden_global_offset_z651; CHECK-NEXT: - .offset: 56652; CHECK-NEXT: .size: 8653; CHECK-NEXT: .value_kind: hidden_printf_buffer654; CHECK-NEXT: - .offset: 64655; CHECK-NEXT: .size: 8656; CHECK-NEXT: .value_kind: hidden_none657; CHECK-NEXT: - .offset: 72658; CHECK-NEXT: .size: 8659; CHECK-NEXT: .value_kind: hidden_none660; CHECK-NEXT: - .offset: 80661; CHECK-NEXT: .size: 8662; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg663; CHECK: .language: OpenCL C664; CHECK-NEXT: .language_version:665; CHECK-NEXT: - 2666; CHECK-NEXT: - 0667; CHECK: .name: test_array_byref_constant668; CHECK: .symbol: test_array_byref_constant.kd669define amdgpu_kernel void @test_array_byref_constant(ptr addrspace(4) byref([32 x i8]) %a) #0670 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20671 !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {672 ret void673}674 675; CHECK: - .args:676; CHECK-NEXT: - .name: a677; CHECK-NEXT: .offset: 0678; CHECK-NEXT: .size: 16679; CHECK-NEXT: .type_name: i128680; CHECK-NEXT: .value_kind: by_value681; CHECK-NEXT: - .offset: 16682; CHECK-NEXT: .size: 8683; CHECK-NEXT: .value_kind: hidden_global_offset_x684; CHECK-NEXT: - .offset: 24685; CHECK-NEXT: .size: 8686; CHECK-NEXT: .value_kind: hidden_global_offset_y687; CHECK-NEXT: - .offset: 32688; CHECK-NEXT: .size: 8689; CHECK-NEXT: .value_kind: hidden_global_offset_z690; CHECK-NEXT: - .offset: 40691; CHECK-NEXT: .size: 8692; CHECK-NEXT: .value_kind: hidden_printf_buffer693; CHECK-NEXT: - .offset: 48694; CHECK-NEXT: .size: 8695; CHECK-NEXT: .value_kind: hidden_none696; CHECK-NEXT: - .offset: 56697; CHECK-NEXT: .size: 8698; CHECK-NEXT: .value_kind: hidden_none699; CHECK-NEXT: - .offset: 64700; CHECK-NEXT: .size: 8701; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg702; CHECK: .language: OpenCL C703; CHECK-NEXT: .language_version:704; CHECK-NEXT: - 2705; CHECK-NEXT: - 0706; CHECK: .name: test_i128707; CHECK: .symbol: test_i128.kd708define amdgpu_kernel void @test_i128(i128 %a) #0709 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21710 !kernel_arg_base_type !21 !kernel_arg_type_qual !4 {711 ret void712}713 714; CHECK: - .args:715; CHECK-NEXT: - .name: a716; CHECK-NEXT: .offset: 0717; CHECK-NEXT: .size: 4718; CHECK-NEXT: .type_name: int719; CHECK-NEXT: .value_kind: by_value720; CHECK-NEXT: - .name: b721; CHECK-NEXT: .offset: 4722; CHECK-NEXT: .size: 4723; CHECK-NEXT: .type_name: short2724; CHECK-NEXT: .value_kind: by_value725; CHECK-NEXT: - .name: c726; CHECK-NEXT: .offset: 8727; CHECK-NEXT: .size: 4728; CHECK-NEXT: .type_name: char3729; CHECK-NEXT: .value_kind: by_value730; CHECK-NEXT: - .offset: 16731; CHECK-NEXT: .size: 8732; CHECK-NEXT: .value_kind: hidden_global_offset_x733; CHECK-NEXT: - .offset: 24734; CHECK-NEXT: .size: 8735; CHECK-NEXT: .value_kind: hidden_global_offset_y736; CHECK-NEXT: - .offset: 32737; CHECK-NEXT: .size: 8738; CHECK-NEXT: .value_kind: hidden_global_offset_z739; CHECK-NEXT: - .offset: 40740; CHECK-NEXT: .size: 8741; CHECK-NEXT: .value_kind: hidden_printf_buffer742; CHECK-NEXT: - .offset: 48743; CHECK-NEXT: .size: 8744; CHECK-NEXT: .value_kind: hidden_none745; CHECK-NEXT: - .offset: 56746; CHECK-NEXT: .size: 8747; CHECK-NEXT: .value_kind: hidden_none748; CHECK-NEXT: - .offset: 64749; CHECK-NEXT: .size: 8750; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg751; CHECK: .language: OpenCL C752; CHECK-NEXT: .language_version:753; CHECK-NEXT: - 2754; CHECK-NEXT: - 0755; CHECK: .name: test_multi_arg756; CHECK: .symbol: test_multi_arg.kd757define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0758 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24759 !kernel_arg_base_type !24 !kernel_arg_type_qual !25 {760 ret void761}762 763; CHECK: - .args:764; CHECK-NEXT: - .address_space: global765; CHECK-NEXT: .name: g766; CHECK-NEXT: .offset: 0767; CHECK-NEXT: .size: 8768; CHECK-NEXT: .type_name: 'int addrspace(5)*'769; CHECK-NEXT: .value_kind: global_buffer770; CHECK-NEXT: - .address_space: constant771; CHECK-NEXT: .name: c772; CHECK-NEXT: .offset: 8773; CHECK-NEXT: .size: 8774; CHECK-NEXT: .type_name: 'int addrspace(5)*'775; CHECK-NEXT: .value_kind: global_buffer776; CHECK-NEXT: - .address_space: local777; CHECK-NEXT: .name: l778; CHECK-NEXT: .offset: 16779; CHECK-NEXT: .pointee_align: 4780; CHECK-NEXT: .size: 4781; CHECK-NEXT: .type_name: 'int addrspace(5)*'782; CHECK-NEXT: .value_kind: dynamic_shared_pointer783; CHECK-NEXT: - .offset: 24784; CHECK-NEXT: .size: 8785; CHECK-NEXT: .value_kind: hidden_global_offset_x786; CHECK-NEXT: - .offset: 32787; CHECK-NEXT: .size: 8788; CHECK-NEXT: .value_kind: hidden_global_offset_y789; CHECK-NEXT: - .offset: 40790; CHECK-NEXT: .size: 8791; CHECK-NEXT: .value_kind: hidden_global_offset_z792; CHECK-NEXT: - .offset: 48793; CHECK-NEXT: .size: 8794; CHECK-NEXT: .value_kind: hidden_printf_buffer795; CHECK-NEXT: - .offset: 56796; CHECK-NEXT: .size: 8797; CHECK-NEXT: .value_kind: hidden_none798; CHECK-NEXT: - .offset: 64799; CHECK-NEXT: .size: 8800; CHECK-NEXT: .value_kind: hidden_none801; CHECK-NEXT: - .offset: 72802; CHECK-NEXT: .size: 8803; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg804; CHECK: .language: OpenCL C805; CHECK-NEXT: .language_version:806; CHECK-NEXT: - 2807; CHECK-NEXT: - 0808; CHECK: .name: test_addr_space809; CHECK: .symbol: test_addr_space.kd810define amdgpu_kernel void @test_addr_space(ptr addrspace(1) %g,811 ptr addrspace(4) %c,812 ptr addrspace(3) align 4 %l) #0813 !kernel_arg_addr_space !50 !kernel_arg_access_qual !23 !kernel_arg_type !51814 !kernel_arg_base_type !51 !kernel_arg_type_qual !25 {815 ret void816}817 818; CHECK: - .args:819; CHECK-NEXT: - .address_space: global820; CHECK-NEXT: .is_volatile: true821; CHECK-NEXT: .name: a822; CHECK-NEXT: .offset: 0823; CHECK-NEXT: .size: 8824; CHECK-NEXT: .type_name: 'int addrspace(5)*'825; CHECK-NEXT: .value_kind: global_buffer826; CHECK-NEXT: - .address_space: global827; CHECK-NEXT: .is_const: true828; CHECK-NEXT: .is_restrict: true829; CHECK-NEXT: .name: b830; CHECK-NEXT: .offset: 8831; CHECK-NEXT: .size: 8832; CHECK-NEXT: .type_name: 'int addrspace(5)*'833; CHECK-NEXT: .value_kind: global_buffer834; CHECK-NEXT: - .is_pipe: true835; CHECK-NEXT: .name: c836; CHECK-NEXT: .offset: 16837; CHECK-NEXT: .size: 8838; CHECK-NEXT: .type_name: 'int addrspace(5)*'839; CHECK-NEXT: .value_kind: pipe840; CHECK-NEXT: - .offset: 24841; CHECK-NEXT: .size: 8842; CHECK-NEXT: .value_kind: hidden_global_offset_x843; CHECK-NEXT: - .offset: 32844; CHECK-NEXT: .size: 8845; CHECK-NEXT: .value_kind: hidden_global_offset_y846; CHECK-NEXT: - .offset: 40847; CHECK-NEXT: .size: 8848; CHECK-NEXT: .value_kind: hidden_global_offset_z849; CHECK-NEXT: - .offset: 48850; CHECK-NEXT: .size: 8851; CHECK-NEXT: .value_kind: hidden_printf_buffer852; CHECK-NEXT: - .offset: 56853; CHECK-NEXT: .size: 8854; CHECK-NEXT: .value_kind: hidden_none855; CHECK-NEXT: - .offset: 64856; CHECK-NEXT: .size: 8857; CHECK-NEXT: .value_kind: hidden_none858; CHECK-NEXT: - .offset: 72859; CHECK-NEXT: .size: 8860; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg861; CHECK: .language: OpenCL C862; CHECK-NEXT: .language_version:863; CHECK-NEXT: - 2864; CHECK-NEXT: - 0865; CHECK: .name: test_type_qual866; CHECK: .symbol: test_type_qual.kd867define amdgpu_kernel void @test_type_qual(ptr addrspace(1) %a,868 ptr addrspace(1) %b,869 ptr addrspace(1) %c) #0870 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !51871 !kernel_arg_base_type !51 !kernel_arg_type_qual !70 {872 ret void873}874 875; CHECK: - .args:876; CHECK-NEXT: - .access: read_only877; CHECK-NEXT: .name: ro878; CHECK-NEXT: .offset: 0879; CHECK-NEXT: .size: 8880; CHECK-NEXT: .type_name: image1d_t881; CHECK-NEXT: .value_kind: image882; CHECK-NEXT: - .access: write_only883; CHECK-NEXT: .name: wo884; CHECK-NEXT: .offset: 8885; CHECK-NEXT: .size: 8886; CHECK-NEXT: .type_name: image2d_t887; CHECK-NEXT: .value_kind: image888; CHECK-NEXT: - .access: read_write889; CHECK-NEXT: .name: rw890; CHECK-NEXT: .offset: 16891; CHECK-NEXT: .size: 8892; CHECK-NEXT: .type_name: image3d_t893; CHECK-NEXT: .value_kind: image894; CHECK-NEXT: - .offset: 24895; CHECK-NEXT: .size: 8896; CHECK-NEXT: .value_kind: hidden_global_offset_x897; CHECK-NEXT: - .offset: 32898; CHECK-NEXT: .size: 8899; CHECK-NEXT: .value_kind: hidden_global_offset_y900; CHECK-NEXT: - .offset: 40901; CHECK-NEXT: .size: 8902; CHECK-NEXT: .value_kind: hidden_global_offset_z903; CHECK-NEXT: - .offset: 48904; CHECK-NEXT: .size: 8905; CHECK-NEXT: .value_kind: hidden_printf_buffer906; CHECK-NEXT: - .offset: 56907; CHECK-NEXT: .size: 8908; CHECK-NEXT: .value_kind: hidden_none909; CHECK-NEXT: - .offset: 64910; CHECK-NEXT: .size: 8911; CHECK-NEXT: .value_kind: hidden_none912; CHECK-NEXT: - .offset: 72913; CHECK-NEXT: .size: 8914; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg915; CHECK: .language: OpenCL C916; CHECK-NEXT: .language_version:917; CHECK-NEXT: - 2918; CHECK-NEXT: - 0919; CHECK: .name: test_access_qual920; CHECK: .symbol: test_access_qual.kd921define amdgpu_kernel void @test_access_qual(ptr addrspace(1) %ro,922 ptr addrspace(1) %wo,923 ptr addrspace(1) %rw) #0924 !kernel_arg_addr_space !60 !kernel_arg_access_qual !61 !kernel_arg_type !62925 !kernel_arg_base_type !62 !kernel_arg_type_qual !25 {926 ret void927}928 929; CHECK: - .args:930; CHECK-NEXT: - .name: a931; CHECK-NEXT: .offset: 0932; CHECK-NEXT: .size: 4933; CHECK-NEXT: .type_name: int934; CHECK-NEXT: .value_kind: by_value935; CHECK-NEXT: - .offset: 8936; CHECK-NEXT: .size: 8937; CHECK-NEXT: .value_kind: hidden_global_offset_x938; CHECK-NEXT: - .offset: 16939; CHECK-NEXT: .size: 8940; CHECK-NEXT: .value_kind: hidden_global_offset_y941; CHECK-NEXT: - .offset: 24942; CHECK-NEXT: .size: 8943; CHECK-NEXT: .value_kind: hidden_global_offset_z944; CHECK-NEXT: - .offset: 32945; CHECK-NEXT: .size: 8946; CHECK-NEXT: .value_kind: hidden_printf_buffer947; CHECK-NEXT: - .offset: 40948; CHECK-NEXT: .size: 8949; CHECK-NEXT: .value_kind: hidden_none950; CHECK-NEXT: - .offset: 48951; CHECK-NEXT: .size: 8952; CHECK-NEXT: .value_kind: hidden_none953; CHECK-NEXT: - .offset: 56954; CHECK-NEXT: .size: 8955; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg956; CHECK: .language: OpenCL C957; CHECK-NEXT: .language_version:958; CHECK-NEXT: - 2959; CHECK-NEXT: - 0960; CHECK: .name: test_vec_type_hint_half961; CHECK: .symbol: test_vec_type_hint_half.kd962; CHECK: .vec_type_hint: half963define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0964 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3965 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !26 {966 ret void967}968 969; CHECK: - .args:970; CHECK-NEXT: - .name: a971; CHECK-NEXT: .offset: 0972; CHECK-NEXT: .size: 4973; CHECK-NEXT: .type_name: int974; CHECK-NEXT: .value_kind: by_value975; CHECK-NEXT: - .offset: 8976; CHECK-NEXT: .size: 8977; CHECK-NEXT: .value_kind: hidden_global_offset_x978; CHECK-NEXT: - .offset: 16979; CHECK-NEXT: .size: 8980; CHECK-NEXT: .value_kind: hidden_global_offset_y981; CHECK-NEXT: - .offset: 24982; CHECK-NEXT: .size: 8983; CHECK-NEXT: .value_kind: hidden_global_offset_z984; CHECK-NEXT: - .offset: 32985; CHECK-NEXT: .size: 8986; CHECK-NEXT: .value_kind: hidden_printf_buffer987; CHECK-NEXT: - .offset: 40988; CHECK-NEXT: .size: 8989; CHECK-NEXT: .value_kind: hidden_none990; CHECK-NEXT: - .offset: 48991; CHECK-NEXT: .size: 8992; CHECK-NEXT: .value_kind: hidden_none993; CHECK-NEXT: - .offset: 56994; CHECK-NEXT: .size: 8995; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg996; CHECK: .language: OpenCL C997; CHECK-NEXT: .language_version:998; CHECK-NEXT: - 2999; CHECK-NEXT: - 01000; CHECK: .name: test_vec_type_hint_float1001; CHECK: .symbol: test_vec_type_hint_float.kd1002; CHECK: .vec_type_hint: float1003define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #01004 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31005 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !27 {1006 ret void1007}1008 1009; CHECK: - .args:1010; CHECK-NEXT: - .name: a1011; CHECK-NEXT: .offset: 01012; CHECK-NEXT: .size: 41013; CHECK-NEXT: .type_name: int1014; CHECK-NEXT: .value_kind: by_value1015; CHECK-NEXT: - .offset: 81016; CHECK-NEXT: .size: 81017; CHECK-NEXT: .value_kind: hidden_global_offset_x1018; CHECK-NEXT: - .offset: 161019; CHECK-NEXT: .size: 81020; CHECK-NEXT: .value_kind: hidden_global_offset_y1021; CHECK-NEXT: - .offset: 241022; CHECK-NEXT: .size: 81023; CHECK-NEXT: .value_kind: hidden_global_offset_z1024; CHECK-NEXT: - .offset: 321025; CHECK-NEXT: .size: 81026; CHECK-NEXT: .value_kind: hidden_printf_buffer1027; CHECK-NEXT: - .offset: 401028; CHECK-NEXT: .size: 81029; CHECK-NEXT: .value_kind: hidden_none1030; CHECK-NEXT: - .offset: 481031; CHECK-NEXT: .size: 81032; CHECK-NEXT: .value_kind: hidden_none1033; CHECK-NEXT: - .offset: 561034; CHECK-NEXT: .size: 81035; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1036; CHECK: .language: OpenCL C1037; CHECK-NEXT: .language_version:1038; CHECK-NEXT: - 21039; CHECK-NEXT: - 01040; CHECK: .name: test_vec_type_hint_double1041; CHECK: .symbol: test_vec_type_hint_double.kd1042; CHECK: .vec_type_hint: double1043define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #01044 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31045 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !28 {1046 ret void1047}1048 1049; CHECK: - .args:1050; CHECK-NEXT: - .name: a1051; CHECK-NEXT: .offset: 01052; CHECK-NEXT: .size: 41053; CHECK-NEXT: .type_name: int1054; CHECK-NEXT: .value_kind: by_value1055; CHECK-NEXT: - .offset: 81056; CHECK-NEXT: .size: 81057; CHECK-NEXT: .value_kind: hidden_global_offset_x1058; CHECK-NEXT: - .offset: 161059; CHECK-NEXT: .size: 81060; CHECK-NEXT: .value_kind: hidden_global_offset_y1061; CHECK-NEXT: - .offset: 241062; CHECK-NEXT: .size: 81063; CHECK-NEXT: .value_kind: hidden_global_offset_z1064; CHECK-NEXT: - .offset: 321065; CHECK-NEXT: .size: 81066; CHECK-NEXT: .value_kind: hidden_printf_buffer1067; CHECK-NEXT: - .offset: 401068; CHECK-NEXT: .size: 81069; CHECK-NEXT: .value_kind: hidden_none1070; CHECK-NEXT: - .offset: 481071; CHECK-NEXT: .size: 81072; CHECK-NEXT: .value_kind: hidden_none1073; CHECK-NEXT: - .offset: 561074; CHECK-NEXT: .size: 81075; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1076; CHECK: .language: OpenCL C1077; CHECK-NEXT: .language_version:1078; CHECK-NEXT: - 21079; CHECK-NEXT: - 01080; CHECK: .name: test_vec_type_hint_char1081; CHECK: .symbol: test_vec_type_hint_char.kd1082; CHECK: .vec_type_hint: char1083define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #01084 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31085 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !29 {1086 ret void1087}1088 1089; CHECK: - .args:1090; CHECK-NEXT: - .name: a1091; CHECK-NEXT: .offset: 01092; CHECK-NEXT: .size: 41093; CHECK-NEXT: .type_name: int1094; CHECK-NEXT: .value_kind: by_value1095; CHECK-NEXT: - .offset: 81096; CHECK-NEXT: .size: 81097; CHECK-NEXT: .value_kind: hidden_global_offset_x1098; CHECK-NEXT: - .offset: 161099; CHECK-NEXT: .size: 81100; CHECK-NEXT: .value_kind: hidden_global_offset_y1101; CHECK-NEXT: - .offset: 241102; CHECK-NEXT: .size: 81103; CHECK-NEXT: .value_kind: hidden_global_offset_z1104; CHECK-NEXT: - .offset: 321105; CHECK-NEXT: .size: 81106; CHECK-NEXT: .value_kind: hidden_printf_buffer1107; CHECK-NEXT: - .offset: 401108; CHECK-NEXT: .size: 81109; CHECK-NEXT: .value_kind: hidden_none1110; CHECK-NEXT: - .offset: 481111; CHECK-NEXT: .size: 81112; CHECK-NEXT: .value_kind: hidden_none1113; CHECK-NEXT: - .offset: 561114; CHECK-NEXT: .size: 81115; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1116; CHECK: .language: OpenCL C1117; CHECK-NEXT: .language_version:1118; CHECK-NEXT: - 21119; CHECK-NEXT: - 01120; CHECK: .name: test_vec_type_hint_short1121; CHECK: .symbol: test_vec_type_hint_short.kd1122; CHECK: .vec_type_hint: short1123define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #01124 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31125 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !30 {1126 ret void1127}1128 1129; CHECK: - .args:1130; CHECK-NEXT: - .name: a1131; CHECK-NEXT: .offset: 01132; CHECK-NEXT: .size: 41133; CHECK-NEXT: .type_name: int1134; CHECK-NEXT: .value_kind: by_value1135; CHECK-NEXT: - .offset: 81136; CHECK-NEXT: .size: 81137; CHECK-NEXT: .value_kind: hidden_global_offset_x1138; CHECK-NEXT: - .offset: 161139; CHECK-NEXT: .size: 81140; CHECK-NEXT: .value_kind: hidden_global_offset_y1141; CHECK-NEXT: - .offset: 241142; CHECK-NEXT: .size: 81143; CHECK-NEXT: .value_kind: hidden_global_offset_z1144; CHECK-NEXT: - .offset: 321145; CHECK-NEXT: .size: 81146; CHECK-NEXT: .value_kind: hidden_printf_buffer1147; CHECK-NEXT: - .offset: 401148; CHECK-NEXT: .size: 81149; CHECK-NEXT: .value_kind: hidden_none1150; CHECK-NEXT: - .offset: 481151; CHECK-NEXT: .size: 81152; CHECK-NEXT: .value_kind: hidden_none1153; CHECK-NEXT: - .offset: 561154; CHECK-NEXT: .size: 81155; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1156; CHECK: .language: OpenCL C1157; CHECK-NEXT: .language_version:1158; CHECK-NEXT: - 21159; CHECK-NEXT: - 01160; CHECK: .name: test_vec_type_hint_long1161; CHECK: .symbol: test_vec_type_hint_long.kd1162; CHECK: .vec_type_hint: long1163define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #01164 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31165 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !31 {1166 ret void1167}1168 1169; CHECK: - .args:1170; CHECK-NEXT: - .name: a1171; CHECK-NEXT: .offset: 01172; CHECK-NEXT: .size: 41173; CHECK-NEXT: .type_name: int1174; CHECK-NEXT: .value_kind: by_value1175; CHECK-NEXT: - .offset: 81176; CHECK-NEXT: .size: 81177; CHECK-NEXT: .value_kind: hidden_global_offset_x1178; CHECK-NEXT: - .offset: 161179; CHECK-NEXT: .size: 81180; CHECK-NEXT: .value_kind: hidden_global_offset_y1181; CHECK-NEXT: - .offset: 241182; CHECK-NEXT: .size: 81183; CHECK-NEXT: .value_kind: hidden_global_offset_z1184; CHECK-NEXT: - .offset: 321185; CHECK-NEXT: .size: 81186; CHECK-NEXT: .value_kind: hidden_printf_buffer1187; CHECK-NEXT: - .offset: 401188; CHECK-NEXT: .size: 81189; CHECK-NEXT: .value_kind: hidden_none1190; CHECK-NEXT: - .offset: 481191; CHECK-NEXT: .size: 81192; CHECK-NEXT: .value_kind: hidden_none1193; CHECK-NEXT: - .offset: 561194; CHECK-NEXT: .size: 81195; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1196; CHECK: .language: OpenCL C1197; CHECK-NEXT: .language_version:1198; CHECK-NEXT: - 21199; CHECK-NEXT: - 01200; CHECK: .name: test_vec_type_hint_unknown1201; CHECK: .symbol: test_vec_type_hint_unknown.kd1202; CHECK: .vec_type_hint: unknown1203define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #01204 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31205 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !32 {1206 ret void1207}1208 1209; CHECK: - .args:1210; CHECK-NEXT: - .name: a1211; CHECK-NEXT: .offset: 01212; CHECK-NEXT: .size: 41213; CHECK-NEXT: .type_name: int1214; CHECK-NEXT: .value_kind: by_value1215; CHECK-NEXT: - .offset: 81216; CHECK-NEXT: .size: 81217; CHECK-NEXT: .value_kind: hidden_global_offset_x1218; CHECK-NEXT: - .offset: 161219; CHECK-NEXT: .size: 81220; CHECK-NEXT: .value_kind: hidden_global_offset_y1221; CHECK-NEXT: - .offset: 241222; CHECK-NEXT: .size: 81223; CHECK-NEXT: .value_kind: hidden_global_offset_z1224; CHECK-NEXT: - .offset: 321225; CHECK-NEXT: .size: 81226; CHECK-NEXT: .value_kind: hidden_printf_buffer1227; CHECK-NEXT: - .offset: 401228; CHECK-NEXT: .size: 81229; CHECK-NEXT: .value_kind: hidden_none1230; CHECK-NEXT: - .offset: 481231; CHECK-NEXT: .size: 81232; CHECK-NEXT: .value_kind: hidden_none1233; CHECK-NEXT: - .offset: 561234; CHECK-NEXT: .size: 81235; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1236; CHECK: .language: OpenCL C1237; CHECK-NEXT: .language_version:1238; CHECK-NEXT: - 21239; CHECK-NEXT: - 01240; CHECK: .name: test_reqd_wgs_vec_type_hint1241; CHECK: .reqd_workgroup_size:1242; CHECK-NEXT: - 11243; CHECK-NEXT: - 21244; CHECK-NEXT: - 41245; CHECK: .symbol: test_reqd_wgs_vec_type_hint.kd1246; CHECK: .vec_type_hint: int1247define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a) #01248 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31249 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !51250 !reqd_work_group_size !6 {1251 ret void1252}1253 1254; CHECK: - .args:1255; CHECK-NEXT: - .name: a1256; CHECK-NEXT: .offset: 01257; CHECK-NEXT: .size: 41258; CHECK-NEXT: .type_name: int1259; CHECK-NEXT: .value_kind: by_value1260; CHECK-NEXT: - .offset: 81261; CHECK-NEXT: .size: 81262; CHECK-NEXT: .value_kind: hidden_global_offset_x1263; CHECK-NEXT: - .offset: 161264; CHECK-NEXT: .size: 81265; CHECK-NEXT: .value_kind: hidden_global_offset_y1266; CHECK-NEXT: - .offset: 241267; CHECK-NEXT: .size: 81268; CHECK-NEXT: .value_kind: hidden_global_offset_z1269; CHECK-NEXT: - .offset: 321270; CHECK-NEXT: .size: 81271; CHECK-NEXT: .value_kind: hidden_printf_buffer1272; CHECK-NEXT: - .offset: 401273; CHECK-NEXT: .size: 81274; CHECK-NEXT: .value_kind: hidden_none1275; CHECK-NEXT: - .offset: 481276; CHECK-NEXT: .size: 81277; CHECK-NEXT: .value_kind: hidden_none1278; CHECK-NEXT: - .offset: 561279; CHECK-NEXT: .size: 81280; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1281; CHECK: .language: OpenCL C1282; CHECK-NEXT: .language_version:1283; CHECK-NEXT: - 21284; CHECK-NEXT: - 01285; CHECK: .name: test_wgs_hint_vec_type_hint1286; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd1287; CHECK: .vec_type_hint: uint41288; CHECK: .workgroup_size_hint:1289; CHECK-NEXT: - 81290; CHECK-NEXT: - 161291; CHECK-NEXT: - 321292define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a) #01293 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !31294 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !71295 !work_group_size_hint !8 {1296 ret void1297}1298 1299; CHECK: - .args:1300; CHECK-NEXT: - .address_space: global1301; CHECK-NEXT: .name: a1302; CHECK-NEXT: .offset: 01303; CHECK-NEXT: .size: 81304; CHECK-NEXT: .type_name: 'int addrspace(5)* addrspace(5)*'1305; CHECK-NEXT: .value_kind: global_buffer1306; CHECK-NEXT: - .offset: 81307; CHECK-NEXT: .size: 81308; CHECK-NEXT: .value_kind: hidden_global_offset_x1309; CHECK-NEXT: - .offset: 161310; CHECK-NEXT: .size: 81311; CHECK-NEXT: .value_kind: hidden_global_offset_y1312; CHECK-NEXT: - .offset: 241313; CHECK-NEXT: .size: 81314; CHECK-NEXT: .value_kind: hidden_global_offset_z1315; CHECK-NEXT: - .offset: 321316; CHECK-NEXT: .size: 81317; CHECK-NEXT: .value_kind: hidden_printf_buffer1318; CHECK-NEXT: - .offset: 401319; CHECK-NEXT: .size: 81320; CHECK-NEXT: .value_kind: hidden_none1321; CHECK-NEXT: - .offset: 481322; CHECK-NEXT: .size: 81323; CHECK-NEXT: .value_kind: hidden_none1324; CHECK-NEXT: - .offset: 561325; CHECK-NEXT: .size: 81326; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1327; CHECK: .language: OpenCL C1328; CHECK-NEXT: .language_version:1329; CHECK-NEXT: - 21330; CHECK-NEXT: - 01331; CHECK: .name: test_arg_ptr_to_ptr1332; CHECK: .symbol: test_arg_ptr_to_ptr.kd1333define amdgpu_kernel void @test_arg_ptr_to_ptr(ptr addrspace(1) %a) #01334 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !801335 !kernel_arg_base_type !80 !kernel_arg_type_qual !4 {1336 ret void1337}1338 1339; CHECK: - .args:1340; CHECK-NEXT: .name: a1341; CHECK-NEXT: .offset: 01342; CHECK-NEXT: .size: 81343; CHECK-NEXT: .type_name: struct B1344; CHECK-NEXT: .value_kind: by_value1345; CHECK-NEXT: - .offset: 81346; CHECK-NEXT: .size: 81347; CHECK-NEXT: .value_kind: hidden_global_offset_x1348; CHECK-NEXT: - .offset: 161349; CHECK-NEXT: .size: 81350; CHECK-NEXT: .value_kind: hidden_global_offset_y1351; CHECK-NEXT: - .offset: 241352; CHECK-NEXT: .size: 81353; CHECK-NEXT: .value_kind: hidden_global_offset_z1354; CHECK-NEXT: - .offset: 321355; CHECK-NEXT: .size: 81356; CHECK-NEXT: .value_kind: hidden_printf_buffer1357; CHECK-NEXT: - .offset: 401358; CHECK-NEXT: .size: 81359; CHECK-NEXT: .value_kind: hidden_none1360; CHECK-NEXT: - .offset: 481361; CHECK-NEXT: .size: 81362; CHECK-NEXT: .value_kind: hidden_none1363; CHECK-NEXT: - .offset: 561364; CHECK-NEXT: .size: 81365; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1366; CHECK: .language: OpenCL C1367; CHECK-NEXT: .language_version:1368; CHECK-NEXT: - 21369; CHECK-NEXT: - 01370; CHECK: .name: test_arg_struct_contains_ptr1371; CHECK: .symbol: test_arg_struct_contains_ptr.kd1372define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B %a) #01373 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !821374 !kernel_arg_base_type !82 !kernel_arg_type_qual !4 {1375 ret void1376}1377 1378; CHECK: - .args:1379; CHECK-NEXT: - .name: a1380; CHECK-NEXT: .offset: 01381; CHECK-NEXT: .size: 161382; CHECK-NEXT: .type_name: 'global int addrspace(5)* __attribute__((ext_vector_type(2)))'1383; CHECK-NEXT: .value_kind: by_value1384; CHECK-NEXT: - .offset: 161385; CHECK-NEXT: .size: 81386; CHECK-NEXT: .value_kind: hidden_global_offset_x1387; CHECK-NEXT: - .offset: 241388; CHECK-NEXT: .size: 81389; CHECK-NEXT: .value_kind: hidden_global_offset_y1390; CHECK-NEXT: - .offset: 321391; CHECK-NEXT: .size: 81392; CHECK-NEXT: .value_kind: hidden_global_offset_z1393; CHECK-NEXT: - .offset: 401394; CHECK-NEXT: .size: 81395; CHECK-NEXT: .value_kind: hidden_printf_buffer1396; CHECK-NEXT: - .offset: 481397; CHECK-NEXT: .size: 81398; CHECK-NEXT: .value_kind: hidden_none1399; CHECK-NEXT: - .offset: 561400; CHECK-NEXT: .size: 81401; CHECK-NEXT: .value_kind: hidden_none1402; CHECK-NEXT: - .offset: 641403; CHECK-NEXT: .size: 81404; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1405; CHECK: .language: OpenCL C1406; CHECK-NEXT: .language_version:1407; CHECK-NEXT: - 21408; CHECK-NEXT: - 01409; CHECK: .name: test_arg_vector_of_ptr1410; CHECK: .symbol: test_arg_vector_of_ptr.kd1411define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x ptr addrspace(1)> %a) #01412 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !831413 !kernel_arg_base_type !83 !kernel_arg_type_qual !4 {1414 ret void1415}1416 1417; CHECK: - .args:1418; CHECK-NEXT: - .address_space: global1419; CHECK-NEXT: .name: a1420; CHECK-NEXT: .offset: 01421; CHECK-NEXT: .size: 81422; CHECK-NEXT: .type_name: clk_event_t1423; CHECK-NEXT: .value_kind: global_buffer1424; CHECK-NEXT: - .offset: 81425; CHECK-NEXT: .size: 81426; CHECK-NEXT: .value_kind: hidden_global_offset_x1427; CHECK-NEXT: - .offset: 161428; CHECK-NEXT: .size: 81429; CHECK-NEXT: .value_kind: hidden_global_offset_y1430; CHECK-NEXT: - .offset: 241431; CHECK-NEXT: .size: 81432; CHECK-NEXT: .value_kind: hidden_global_offset_z1433; CHECK-NEXT: - .offset: 321434; CHECK-NEXT: .size: 81435; CHECK-NEXT: .value_kind: hidden_printf_buffer1436; CHECK-NEXT: - .offset: 401437; CHECK-NEXT: .size: 81438; CHECK-NEXT: .value_kind: hidden_none1439; CHECK-NEXT: - .offset: 481440; CHECK-NEXT: .size: 81441; CHECK-NEXT: .value_kind: hidden_none1442; CHECK-NEXT: - .offset: 561443; CHECK-NEXT: .size: 81444; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1445; CHECK: .language: OpenCL C1446; CHECK-NEXT: .language_version:1447; CHECK-NEXT: - 21448; CHECK-NEXT: - 01449; CHECK: .name: test_arg_unknown_builtin_type1450; CHECK: .symbol: test_arg_unknown_builtin_type.kd1451define amdgpu_kernel void @test_arg_unknown_builtin_type(1452 ptr addrspace(1) %a) #01453 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !841454 !kernel_arg_base_type !84 !kernel_arg_type_qual !4 {1455 ret void1456}1457 1458; CHECK: - .args:1459; CHECK-NEXT: - .address_space: global1460; CHECK-NEXT: .name: a1461; CHECK-NEXT: .offset: 01462; CHECK-NEXT: .size: 81463; CHECK-NEXT: .type_name: 'long addrspace(5)*'1464; CHECK-NEXT: .value_kind: global_buffer1465; CHECK-NEXT: - .address_space: local1466; CHECK-NEXT: .name: b1467; CHECK-NEXT: .offset: 81468; CHECK-NEXT: .pointee_align: 11469; CHECK-NEXT: .size: 41470; CHECK-NEXT: .type_name: 'char addrspace(5)*'1471; CHECK-NEXT: .value_kind: dynamic_shared_pointer1472; CHECK-NEXT: - .address_space: local1473; CHECK-NEXT: .name: c1474; CHECK-NEXT: .offset: 121475; CHECK-NEXT: .pointee_align: 21476; CHECK-NEXT: .size: 41477; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'1478; CHECK-NEXT: .value_kind: dynamic_shared_pointer1479; CHECK-NEXT: - .address_space: local1480; CHECK-NEXT: .name: d1481; CHECK-NEXT: .offset: 161482; CHECK-NEXT: .pointee_align: 41483; CHECK-NEXT: .size: 41484; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'1485; CHECK-NEXT: .value_kind: dynamic_shared_pointer1486; CHECK-NEXT: - .address_space: local1487; CHECK-NEXT: .name: e1488; CHECK-NEXT: .offset: 201489; CHECK-NEXT: .pointee_align: 41490; CHECK-NEXT: .size: 41491; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'1492; CHECK-NEXT: .value_kind: dynamic_shared_pointer1493; CHECK-NEXT: - .address_space: local1494; CHECK-NEXT: .name: f1495; CHECK-NEXT: .offset: 241496; CHECK-NEXT: .pointee_align: 81497; CHECK-NEXT: .size: 41498; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'1499; CHECK-NEXT: .value_kind: dynamic_shared_pointer1500; CHECK-NEXT: - .address_space: local1501; CHECK-NEXT: .name: g1502; CHECK-NEXT: .offset: 281503; CHECK-NEXT: .pointee_align: 161504; CHECK-NEXT: .size: 41505; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'1506; CHECK-NEXT: .value_kind: dynamic_shared_pointer1507; CHECK-NEXT: - .address_space: local1508; CHECK-NEXT: .name: h1509; CHECK-NEXT: .offset: 321510; CHECK-NEXT: .pointee_align: 11511; CHECK-NEXT: .size: 41512; CHECK-NEXT: .value_kind: dynamic_shared_pointer1513; CHECK-NEXT: - .offset: 401514; CHECK-NEXT: .size: 81515; CHECK-NEXT: .value_kind: hidden_global_offset_x1516; CHECK-NEXT: - .offset: 481517; CHECK-NEXT: .size: 81518; CHECK-NEXT: .value_kind: hidden_global_offset_y1519; CHECK-NEXT: - .offset: 561520; CHECK-NEXT: .size: 81521; CHECK-NEXT: .value_kind: hidden_global_offset_z1522; CHECK-NEXT: - .offset: 641523; CHECK-NEXT: .size: 81524; CHECK-NEXT: .value_kind: hidden_printf_buffer1525; CHECK-NEXT: - .offset: 721526; CHECK-NEXT: .size: 81527; CHECK-NEXT: .value_kind: hidden_none1528; CHECK-NEXT: - .offset: 801529; CHECK-NEXT: .size: 81530; CHECK-NEXT: .value_kind: hidden_none1531; CHECK-NEXT: - .offset: 881532; CHECK-NEXT: .size: 81533; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1534; CHECK: .language: OpenCL C1535; CHECK-NEXT: .language_version:1536; CHECK-NEXT: - 21537; CHECK-NEXT: - 01538; CHECK: .name: test_pointee_align1539; CHECK: .symbol: test_pointee_align.kd1540define amdgpu_kernel void @test_pointee_align(ptr addrspace(1) %a,1541 ptr addrspace(3) %b,1542 ptr addrspace(3) align 2 %c,1543 ptr addrspace(3) align 4 %d,1544 ptr addrspace(3) align 4 %e,1545 ptr addrspace(3) align 8 %f,1546 ptr addrspace(3) align 16 %g,1547 ptr addrspace(3) %h) #01548 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !931549 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {1550 ret void1551}1552 1553; CHECK: - .args:1554; CHECK-NEXT: - .address_space: global1555; CHECK-NEXT: .name: a1556; CHECK-NEXT: .offset: 01557; CHECK-NEXT: .size: 81558; CHECK-NEXT: .type_name: 'long addrspace(5)*'1559; CHECK-NEXT: .value_kind: global_buffer1560; CHECK-NEXT: - .address_space: local1561; CHECK-NEXT: .name: b1562; CHECK-NEXT: .offset: 81563; CHECK-NEXT: .pointee_align: 81564; CHECK-NEXT: .size: 41565; CHECK-NEXT: .type_name: 'char addrspace(5)*'1566; CHECK-NEXT: .value_kind: dynamic_shared_pointer1567; CHECK-NEXT: - .address_space: local1568; CHECK-NEXT: .name: c1569; CHECK-NEXT: .offset: 121570; CHECK-NEXT: .pointee_align: 321571; CHECK-NEXT: .size: 41572; CHECK-NEXT: .type_name: 'char2 addrspace(5)*'1573; CHECK-NEXT: .value_kind: dynamic_shared_pointer1574; CHECK-NEXT: - .address_space: local1575; CHECK-NEXT: .name: d1576; CHECK-NEXT: .offset: 161577; CHECK-NEXT: .pointee_align: 641578; CHECK-NEXT: .size: 41579; CHECK-NEXT: .type_name: 'char3 addrspace(5)*'1580; CHECK-NEXT: .value_kind: dynamic_shared_pointer1581; CHECK-NEXT: - .address_space: local1582; CHECK-NEXT: .name: e1583; CHECK-NEXT: .offset: 201584; CHECK-NEXT: .pointee_align: 2561585; CHECK-NEXT: .size: 41586; CHECK-NEXT: .type_name: 'char4 addrspace(5)*'1587; CHECK-NEXT: .value_kind: dynamic_shared_pointer1588; CHECK-NEXT: - .address_space: local1589; CHECK-NEXT: .name: f1590; CHECK-NEXT: .offset: 241591; CHECK-NEXT: .pointee_align: 1281592; CHECK-NEXT: .size: 41593; CHECK-NEXT: .type_name: 'char8 addrspace(5)*'1594; CHECK-NEXT: .value_kind: dynamic_shared_pointer1595; CHECK-NEXT: - .address_space: local1596; CHECK-NEXT: .name: g1597; CHECK-NEXT: .offset: 281598; CHECK-NEXT: .pointee_align: 10241599; CHECK-NEXT: .size: 41600; CHECK-NEXT: .type_name: 'char16 addrspace(5)*'1601; CHECK-NEXT: .value_kind: dynamic_shared_pointer1602; CHECK-NEXT: - .address_space: local1603; CHECK-NEXT: .name: h1604; CHECK-NEXT: .offset: 321605; CHECK-NEXT: .pointee_align: 161606; CHECK-NEXT: .size: 41607; CHECK-NEXT: .value_kind: dynamic_shared_pointer1608; CHECK-NEXT: - .offset: 401609; CHECK-NEXT: .size: 81610; CHECK-NEXT: .value_kind: hidden_global_offset_x1611; CHECK-NEXT: - .offset: 481612; CHECK-NEXT: .size: 81613; CHECK-NEXT: .value_kind: hidden_global_offset_y1614; CHECK-NEXT: - .offset: 561615; CHECK-NEXT: .size: 81616; CHECK-NEXT: .value_kind: hidden_global_offset_z1617; CHECK-NEXT: - .offset: 641618; CHECK-NEXT: .size: 81619; CHECK-NEXT: .value_kind: hidden_printf_buffer1620; CHECK-NEXT: - .offset: 721621; CHECK-NEXT: .size: 81622; CHECK-NEXT: .value_kind: hidden_none1623; CHECK-NEXT: - .offset: 801624; CHECK-NEXT: .size: 81625; CHECK-NEXT: .value_kind: hidden_none1626; CHECK-NEXT: - .offset: 881627; CHECK-NEXT: .size: 81628; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1629; CHECK: .language: OpenCL C1630; CHECK-NEXT: .language_version:1631; CHECK-NEXT: - 21632; CHECK-NEXT: - 01633; CHECK: .name: test_pointee_align_attribute1634; CHECK: .symbol: test_pointee_align_attribute.kd1635define amdgpu_kernel void @test_pointee_align_attribute(ptr addrspace(1) align 16 %a,1636 ptr addrspace(3) align 8 %b,1637 ptr addrspace(3) align 32 %c,1638 ptr addrspace(3) align 64 %d,1639 ptr addrspace(3) align 256 %e,1640 ptr addrspace(3) align 128 %f,1641 ptr addrspace(3) align 1024 %g,1642 ptr addrspace(3) align 16 %h) #01643 !kernel_arg_addr_space !91 !kernel_arg_access_qual !92 !kernel_arg_type !931644 !kernel_arg_base_type !93 !kernel_arg_type_qual !94 {1645 ret void1646}1647; CHECK: - .args:1648; CHECK-NEXT: - .name: arg1649; CHECK-NEXT: .offset: 01650; CHECK-NEXT: .size: 251651; CHECK-NEXT: .type_name: __block_literal1652; CHECK-NEXT: .value_kind: by_value1653; CHECK-NEXT: - .offset: 321654; CHECK-NEXT: .size: 81655; CHECK-NEXT: .value_kind: hidden_global_offset_x1656; CHECK-NEXT: - .offset: 401657; CHECK-NEXT: .size: 81658; CHECK-NEXT: .value_kind: hidden_global_offset_y1659; CHECK-NEXT: - .offset: 481660; CHECK-NEXT: .size: 81661; CHECK-NEXT: .value_kind: hidden_global_offset_z1662; CHECK-NEXT: - .offset: 561663; CHECK-NEXT: .size: 81664; CHECK-NEXT: .value_kind: hidden_printf_buffer1665; CHECK-NEXT: - .offset: 641666; CHECK-NEXT: .size: 81667; CHECK-NEXT: .value_kind: hidden_none1668; CHECK-NEXT: - .offset: 721669; CHECK-NEXT: .size: 81670; CHECK-NEXT: .value_kind: hidden_none1671; CHECK-NEXT: - .offset: 801672; CHECK-NEXT: .size: 81673; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1674; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle1675; CHECK: .language: OpenCL C1676; CHECK-NEXT: .language_version:1677; CHECK-NEXT: - 21678; CHECK-NEXT: - 01679; CHECK: .name: __test_block_invoke_kernel1680; CHECK: .symbol: __test_block_invoke_kernel.kd1681define amdgpu_kernel void @__test_block_invoke_kernel(1682 <{ i32, i32, ptr, ptr addrspace(1), i8 }> %arg) #1 !associated !1121683 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !1101684 !kernel_arg_base_type !110 !kernel_arg_type_qual !4 {1685 ret void1686}1687 1688; CHECK: - .args:1689; CHECK-NEXT: - .name: a1690; CHECK-NEXT: .offset: 01691; CHECK-NEXT: .size: 11692; CHECK-NEXT: .type_name: char1693; CHECK-NEXT: .value_kind: by_value1694; CHECK-NEXT: - .offset: 81695; CHECK-NEXT: .size: 81696; CHECK-NEXT: .value_kind: hidden_global_offset_x1697; CHECK-NEXT: - .offset: 161698; CHECK-NEXT: .size: 81699; CHECK-NEXT: .value_kind: hidden_global_offset_y1700; CHECK-NEXT: - .offset: 241701; CHECK-NEXT: .size: 81702; CHECK-NEXT: .value_kind: hidden_global_offset_z1703; CHECK-NEXT: - .offset: 321704; CHECK-NEXT: .size: 81705; CHECK-NEXT: .value_kind: hidden_printf_buffer1706; CHECK-NEXT: - .offset: 401707; CHECK-NEXT: .size: 81708; CHECK-NEXT: .value_kind: hidden_default_queue1709; CHECK-NEXT: - .offset: 481710; CHECK-NEXT: .size: 81711; CHECK-NEXT: .value_kind: hidden_completion_action1712; CHECK-NEXT: - .offset: 561713; CHECK-NEXT: .size: 81714; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg1715; CHECK: .language: OpenCL C1716; CHECK-NEXT: .language_version:1717; CHECK-NEXT: - 21718; CHECK-NEXT: - 01719; CHECK: .name: test_enqueue_kernel_caller1720; CHECK: .symbol: test_enqueue_kernel_caller.kd1721define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #21722 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !91723 !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {1724 ret void1725}1726 1727; CHECK: - .args:1728; CHECK-NEXT: - .name: ptr1729; CHECK-NEXT: .offset: 01730; CHECK-NEXT: .size: 81731; CHECK-NEXT: .value_kind: global_buffer1732; CHECK: .name: unknown_addrspace_kernarg1733; CHECK: .symbol: unknown_addrspace_kernarg.kd1734define amdgpu_kernel void @unknown_addrspace_kernarg(ptr addrspace(12345) %ptr) #0 {1735 ret void1736}1737 1738; Make sure the device_enqueue_symbol is not reported1739; CHECK: - .args: []1740; CHECK-NEXT: .group_segment_fixed_size: 01741; CHECK-NEXT: .kernarg_segment_align: 41742; CHECK-NEXT: .kernarg_segment_size: 01743; CHECK-NEXT: .language: OpenCL C1744; CHECK-NEXT: .language_version:1745; CHECK-NEXT: - 21746; CHECK-NEXT: - 01747; CHECK-NEXT: .max_flat_workgroup_size: 10241748; CHECK-NEXT: .name: associated_global_not_handle1749; CHECK-NEXT: .private_segment_fixed_size: 01750; CHECK-NEXT: .sgpr_count:1751; CHECK-NEXT: .sgpr_spill_count: 01752; CHECK-NEXT: .symbol: associated_global_not_handle.kd1753; CHECK-NEXT: .vgpr_count:1754; CHECK-NEXT: .vgpr_spill_count: 01755; CHECK-NEXT: .wavefront_size: 641756; CHECK-NOT: device_enqueue_symbol1757define amdgpu_kernel void @associated_global_not_handle() #3 !associated !113 {1758 ret void1759}1760 1761; CHECK: amdhsa.printf:1762; CHECK-NEXT: - '1:1:4:%d\n'1763; CHECK-NEXT: - '2:1:8:%g\n'1764; CHECK: amdhsa.version:1765; CHECK-NEXT: - 11766; CHECK-NEXT: - 11767 1768attributes #0 = { optnone noinline "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-implicitarg-num-bytes"="56" }1769attributes #1 = { optnone noinline "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-implicitarg-num-bytes"="56" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }1770attributes #2 = { optnone noinline "amdgpu-implicitarg-num-bytes"="56" }1771attributes #3 = { optnone noinline "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" }1772 1773!llvm.module.flags = !{!0}1774!0 = !{i32 1, !"amdhsa_code_object_version", i32 400}1775 1776!llvm.printf.fmts = !{!100, !101}1777 1778!1 = !{i32 0}1779!2 = !{!"none"}1780!3 = !{!"int"}1781!4 = !{!""}1782!5 = !{i32 poison, i32 1}1783!6 = !{i32 1, i32 2, i32 4}1784!7 = !{<4 x i32> poison, i32 0}1785!8 = !{i32 8, i32 16, i32 32}1786!9 = !{!"char"}1787!10 = !{!"ushort2"}1788!11 = !{!"int3"}1789!12 = !{!"ulong4"}1790!13 = !{!"half8"}1791!14 = !{!"float16"}1792!15 = !{!"double16"}1793!16 = !{!"int addrspace(5)*"}1794!17 = !{!"image2d_t"}1795!18 = !{!"sampler_t"}1796!19 = !{!"queue_t"}1797!20 = !{!"struct A"}1798!21 = !{!"i128"}1799!22 = !{i32 0, i32 0, i32 0}1800!23 = !{!"none", !"none", !"none"}1801!24 = !{!"int", !"short2", !"char3"}1802!25 = !{!"", !"", !""}1803!26 = !{half poison, i32 1}1804!27 = !{float poison, i32 1}1805!28 = !{double poison, i32 1}1806!29 = !{i8 poison, i32 1}1807!30 = !{i16 poison, i32 1}1808!31 = !{i64 poison, i32 1}1809!32 = !{ptr addrspace(5) poison, i32 1}1810!50 = !{i32 1, i32 2, i32 3}1811!51 = !{!"int addrspace(5)*", !"int addrspace(5)*", !"int addrspace(5)*"}1812!60 = !{i32 1, i32 1, i32 1}1813!61 = !{!"read_only", !"write_only", !"read_write"}1814!62 = !{!"image1d_t", !"image2d_t", !"image3d_t"}1815!70 = !{!"volatile", !"const restrict", !"pipe"}1816!80 = !{!"int addrspace(5)* addrspace(5)*"}1817!81 = !{i32 1}1818!82 = !{!"struct B"}1819!83 = !{!"global int addrspace(5)* __attribute__((ext_vector_type(2)))"}1820!84 = !{!"clk_event_t"}1821!opencl.ocl.version = !{!90}1822!90 = !{i32 2, i32 0}1823!91 = !{i32 0, i32 3, i32 3, i32 3, i32 3, i32 3, i32 3}1824!92 = !{!"none", !"none", !"none", !"none", !"none", !"none", !"none"}1825!93 = !{!"long addrspace(5)*", !"char addrspace(5)*", !"char2 addrspace(5)*", !"char3 addrspace(5)*", !"char4 addrspace(5)*", !"char8 addrspace(5)*", !"char16 addrspace(5)*"}1826!94 = !{!"", !"", !"", !"", !"", !"", !""}1827!100 = !{!"1:1:4:%d\5Cn"}1828!101 = !{!"2:1:8:%g\5Cn"}1829!110 = !{!"__block_literal"}1830!111 = !{!"char", !"char"}1831!112 = !{ptr addrspace(1) @__test_block_invoke_kernel_runtime_handle }1832!113 = !{ptr addrspace(1) @not.a.handle }1833 1834; PARSER: AMDGPU HSA Metadata Parser Test: PASS1835