This is an archive of the discontinued LLVM Phabricator instance.

Differential D132145

AMDGPU/MetaData: Restrict address space key to only be emitted for "global_buffer" and "dynamic_shared_pointer"
ClosedPublic

Authored by raghavmedicherla on Aug 18 2022, 10:00 AM.

Details

Reviewers
scott.linder
kzhuravl
Commits
rG79d2529c10bf: AMDGPU/MetaData: Restrict address space key to only be emitted for…
Summary

This matches .address_space docs at https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v3.

Diff Detail

Event Timeline

raghavmedicherla created this revision.Aug 18 2022, 10:00 AM
Herald added a project: Restricted Project. · View Herald TranscriptAug 18 2022, 10:00 AM
Herald added subscribers: kosarev, foad, kerbowa and 8 others. · View Herald Transcript
raghavmedicherla requested review of this revision.Aug 18 2022, 10:00 AM
Herald added a project: Restricted Project. · View Herald TranscriptAug 18 2022, 10:00 AM
Herald added a subscriber: wdng. · View Herald Transcript
kzhuravl added a comment.Aug 18 2022, 10:04 AM

Please upload the full diff

raghavmedicherla updated this revision to Diff 453696.Aug 18 2022, 10:10 AM
raghavmedicherla removed projects: Restricted Project, Restricted Project.

Uploaded full diff.

Herald added projects: Restricted Project, Restricted Project. · View Herald TranscriptAug 18 2022, 10:10 AM
raghavmedicherla removed subscribers: wdng, arsenm, jvesely and 9 others.Aug 18 2022, 10:12 AM
Harbormaster completed remote builds in B182027: Diff 453696.Aug 18 2022, 11:02 AM
scott.linder accepted this revision.Aug 18 2022, 12:31 PM
This revision is now accepted and ready to land.Aug 18 2022, 12:31 PM
Closed by commit rG79d2529c10bf: AMDGPU/MetaData: Restrict address space key to only be emitted for… (authored by raghavmedicherla). · Explain WhyAug 23 2022, 11:01 AM
This revision was automatically updated to reflect the committed changes.
raghavmedicherla added a commit: rG79d2529c10bf: AMDGPU/MetaData: Restrict address space key to only be emitted for….
foad added inline comments.Aug 24 2022, 1:51 AM
llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp
771–774

Comment seems redundant (see https://llvm.org/docs/CodingStandards.html#commenting).

Revision Contents

Diff 454899

llvm/lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

Show First 20 Lines • Show All 762 Lines • ▼ Show 20 Linesvoid MetadataStreamerV3::emitKernelArg(
Arg[".offset"] = Arg.getDocument()->getNode(Offset); Arg[".offset"] = Arg.getDocument()->getNode(Offset);
Offset += Size; Offset += Size;
Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true); Arg[".value_kind"] = Arg.getDocument()->getNode(ValueKind, /*Copy=*/true);
if (PointeeAlign) if (PointeeAlign)
Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value()); Arg[".pointee_align"] = Arg.getDocument()->getNode(PointeeAlign->value());
if (auto PtrTy = dyn_cast<PointerType>(Ty)) if (auto PtrTy = dyn_cast<PointerType>(Ty))
if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace())) if (auto Qualifier = getAddressSpaceQualifier(PtrTy->getAddressSpace()))
Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier, /*Copy=*/true); // Limiting address space to emit only for a certain ValueKind.
if (ValueKind == "global_buffer" || ValueKind == "dynamic_shared_pointer")
Arg[".address_space"] = Arg.getDocument()->getNode(*Qualifier,
/*Copy=*/true);
foadUnsubmitted
Not Done
ReplyInline Actions

Comment seems redundant (see https://llvm.org/docs/CodingStandards.html#commenting).

foad: Comment seems redundant (see https://llvm.org/docs/CodingStandards.html#commenting).
if (auto AQ = getAccessQualifier(AccQual)) if (auto AQ = getAccessQualifier(AccQual))
Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true); Arg[".access"] = Arg.getDocument()->getNode(*AQ, /*Copy=*/true);
// TODO: Emit Arg[".actual_access"]. // TODO: Emit Arg[".actual_access"].
SmallVector<StringRef, 1> SplitTypeQuals; SmallVector<StringRef, 1> SplitTypeQuals;
TypeQual.split(SplitTypeQuals, " ", -1, false); TypeQual.split(SplitTypeQuals, " ", -1, false);
▲ Show 20 Lines • Show All 303 Lines • Show Last 20 Lines

llvm/test/CodeGen/AMDGPU/hsa-metadata-enqueue-kernel-v3.ll

Show First 20 Lines • Show All 41 Lines • ▼ Show 20 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_hostcall_buffer ; CHECK-NEXT: .value_kind: hidden_hostcall_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_default_queue ; CHECK-NEXT: .value_kind: hidden_default_queue
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_completion_action ; CHECK-NEXT: .value_kind: hidden_completion_action
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_enqueue_kernel_caller ; CHECK: .name: test_enqueue_kernel_caller
; CHECK: .symbol: test_enqueue_kernel_caller.kd ; CHECK: .symbol: test_enqueue_kernel_caller.kd
Show All 23 Lines

llvm/test/CodeGen/AMDGPU/hsa-metadata-from-llvm-ir-full-v3.ll

Show All 27 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NOT: .value_kind: hidden_default_queue ; CHECK-NOT: .value_kind: hidden_default_queue
; CHECK-NOT: .value_kind: hidden_completion_action ; CHECK-NOT: .value_kind: hidden_completion_action
; CHECK-NOT: .value_kind: hidden_hostcall_buffer ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK: .value_kind: hidden_multigrid_sync_arg ; CHECK: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
Show All 17 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NOT: .value_kind: hidden_default_queue ; CHECK-NOT: .value_kind: hidden_default_queue
; CHECK-NOT: .value_kind: hidden_completion_action ; CHECK-NOT: .value_kind: hidden_completion_action
; CHECK-NOT: .value_kind: hidden_hostcall_buffer ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK: .value_kind: hidden_multigrid_sync_arg ; CHECK: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
Show All 21 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 528 ; CHECK-NEXT: - .offset: 528
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 536 ; CHECK-NEXT: - .offset: 536
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 544
; CHECK-NEXT: .offset: 544
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NOT: .value_kind: hidden_default_queue ; CHECK-NOT: .value_kind: hidden_default_queue
; CHECK-NOT: .value_kind: hidden_completion_action ; CHECK-NOT: .value_kind: hidden_completion_action
; CHECK-NOT: .value_kind: hidden_hostcall_buffer ; CHECK-NOT: .value_kind: hidden_hostcall_buffer
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK: .value_kind: hidden_multigrid_sync_arg ; CHECK: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
Show All 17 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_ushort2 ; CHECK: .name: test_ushort2
; CHECK: .symbol: test_ushort2.kd ; CHECK: .symbol: test_ushort2.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_int3 ; CHECK: .name: test_int3
; CHECK: .symbol: test_int3.kd ; CHECK: .symbol: test_int3.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 48 ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_ulong4 ; CHECK: .name: test_ulong4
; CHECK: .symbol: test_ulong4.kd ; CHECK: .symbol: test_ulong4.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_half8 ; CHECK: .name: test_half8
; CHECK: .symbol: test_half8.kd ; CHECK: .symbol: test_half8.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 72 ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 80 ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 88
; CHECK-NEXT: .offset: 88
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 96
; CHECK-NEXT: .offset: 96
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 104
; CHECK-NEXT: .offset: 104
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 112
; CHECK-NEXT: .offset: 112
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_float16 ; CHECK: .name: test_float16
; CHECK: .symbol: test_float16.kd ; CHECK: .symbol: test_float16.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 136 ; CHECK-NEXT: - .offset: 136
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 144 ; CHECK-NEXT: - .offset: 144
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 152
; CHECK-NEXT: .offset: 152
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 160
; CHECK-NEXT: .offset: 160
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 168
; CHECK-NEXT: .offset: 168
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 176
; CHECK-NEXT: .offset: 176
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_double16 ; CHECK: .name: test_double16
; CHECK: .symbol: test_double16.kd ; CHECK: .symbol: test_double16.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_pointer ; CHECK: .name: test_pointer
; CHECK: .symbol: test_pointer.kd ; CHECK: .symbol: test_pointer.kd
define amdgpu_kernel void @test_pointer(i32 addrspace(1)* %a) #0 define amdgpu_kernel void @test_pointer(i32 addrspace(1)* %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !16
!kernel_arg_base_type !16 !kernel_arg_type_qual !4 { !kernel_arg_base_type !16 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - .args: ; CHECK: - .args:
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .name: a
; CHECK-NEXT: .name: a
; CHECK-NEXT: .offset: 0 ; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: image2d_t ; CHECK-NEXT: .type_name: image2d_t
; CHECK-NEXT: .value_kind: image ; CHECK-NEXT: .value_kind: image
; CHECK-NEXT: - .offset: 8 ; CHECK-NEXT: - .offset: 8
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_image ; CHECK: .name: test_image
; CHECK: .symbol: test_image.kd ; CHECK: .symbol: test_image.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_sampler ; CHECK: .name: test_sampler
; CHECK: .symbol: test_sampler.kd ; CHECK: .symbol: test_sampler.kd
define amdgpu_kernel void @test_sampler(i32 %a) #0 define amdgpu_kernel void @test_sampler(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !18
!kernel_arg_base_type !18 !kernel_arg_type_qual !4 { !kernel_arg_base_type !18 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - .args: ; CHECK: - .args:
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .name: a
; CHECK-NEXT: .name: a
; CHECK-NEXT: .offset: 0 ; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: queue_t ; CHECK-NEXT: .type_name: queue_t
; CHECK-NEXT: .value_kind: queue ; CHECK-NEXT: .value_kind: queue
; CHECK-NEXT: - .offset: 8 ; CHECK-NEXT: - .offset: 8
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_queue ; CHECK: .name: test_queue
; CHECK: .symbol: test_queue.kd ; CHECK: .symbol: test_queue.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_struct ; CHECK: .name: test_struct
; CHECK: .symbol: test_struct.kd ; CHECK: .symbol: test_struct.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_struct_byref_constant ; CHECK: .name: test_struct_byref_constant
; CHECK: .symbol: test_struct_byref_constant.kd ; CHECK: .symbol: test_struct_byref_constant.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 48 ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_array ; CHECK: .name: test_array
; CHECK: .symbol: test_array.kd ; CHECK: .symbol: test_array.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 48 ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_array_byref_constant ; CHECK: .name: test_array_byref_constant
; CHECK: .symbol: test_array_byref_constant.kd ; CHECK: .symbol: test_array_byref_constant.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_i128 ; CHECK: .name: test_i128
; CHECK: .symbol: test_i128.kd ; CHECK: .symbol: test_i128.kd
Show All 23 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_multi_arg ; CHECK: .name: test_multi_arg
; CHECK: .symbol: test_multi_arg.kd ; CHECK: .symbol: test_multi_arg.kd
Show All 27 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_addr_space ; CHECK: .name: test_addr_space
; CHECK: .symbol: test_addr_space.kd ; CHECK: .symbol: test_addr_space.kd
Show All 16 Lines
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .is_const: true ; CHECK-NEXT: .is_const: true
; CHECK-NEXT: .is_restrict: true ; CHECK-NEXT: .is_restrict: true
; CHECK-NEXT: .name: b ; CHECK-NEXT: .name: b
; CHECK-NEXT: .offset: 8 ; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: 'int addrspace(5)*' ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
; CHECK-NEXT: .value_kind: global_buffer ; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .is_pipe: true
; CHECK-NEXT: .is_pipe: true
; CHECK-NEXT: .name: c ; CHECK-NEXT: .name: c
; CHECK-NEXT: .offset: 16 ; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: 'int addrspace(5)*' ; CHECK-NEXT: .type_name: 'int addrspace(5)*'
; CHECK-NEXT: .value_kind: pipe ; CHECK-NEXT: .value_kind: pipe
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_type_qual ; CHECK: .name: test_type_qual
; CHECK: .symbol: test_type_qual.kd ; CHECK: .symbol: test_type_qual.kd
define amdgpu_kernel void @test_type_qual(i32 addrspace(1)* %a, define amdgpu_kernel void @test_type_qual(i32 addrspace(1)* %a,
i32 addrspace(1)* %b, i32 addrspace(1)* %b,
%opencl.pipe_t addrspace(1)* %c) #0 %opencl.pipe_t addrspace(1)* %c) #0
!kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !51 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !51
!kernel_arg_base_type !51 !kernel_arg_type_qual !70 { !kernel_arg_base_type !51 !kernel_arg_type_qual !70 {
ret void ret void
} }
; CHECK: - .args: ; CHECK: - .args:
; CHECK-NEXT: - .access: read_only ; CHECK-NEXT: - .access: read_only
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: .name: ro ; CHECK-NEXT: .name: ro
; CHECK-NEXT: .offset: 0 ; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: image1d_t ; CHECK-NEXT: .type_name: image1d_t
; CHECK-NEXT: .value_kind: image ; CHECK-NEXT: .value_kind: image
; CHECK-NEXT: - .access: write_only ; CHECK-NEXT: - .access: write_only
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: .name: wo ; CHECK-NEXT: .name: wo
; CHECK-NEXT: .offset: 8 ; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: image2d_t ; CHECK-NEXT: .type_name: image2d_t
; CHECK-NEXT: .value_kind: image ; CHECK-NEXT: .value_kind: image
; CHECK-NEXT: - .access: read_write ; CHECK-NEXT: - .access: read_write
; CHECK-NEXT: .address_space: global
; CHECK-NEXT: .name: rw ; CHECK-NEXT: .name: rw
; CHECK-NEXT: .offset: 16 ; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .type_name: image3d_t ; CHECK-NEXT: .type_name: image3d_t
; CHECK-NEXT: .value_kind: image ; CHECK-NEXT: .value_kind: image
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_access_qual ; CHECK: .name: test_access_qual
; CHECK: .symbol: test_access_qual.kd ; CHECK: .symbol: test_access_qual.kd
Show All 15 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_vec_type_hint_half ; CHECK: .name: test_vec_type_hint_half
; CHECK: .symbol: test_vec_type_hint_half.kd ; CHECK: .symbol: test_vec_type_hint_half.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_vec_type_hint_float ; CHECK: .name: test_vec_type_hint_float
; CHECK: .symbol: test_vec_type_hint_float.kd ; CHECK: .symbol: test_vec_type_hint_float.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_vec_type_hint_double ; CHECK: .name: test_vec_type_hint_double
; CHECK: .symbol: test_vec_type_hint_double.kd ; CHECK: .symbol: test_vec_type_hint_double.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_vec_type_hint_char ; CHECK: .name: test_vec_type_hint_char
; CHECK: .symbol: test_vec_type_hint_char.kd ; CHECK: .symbol: test_vec_type_hint_char.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_vec_type_hint_short ; CHECK: .name: test_vec_type_hint_short
; CHECK: .symbol: test_vec_type_hint_short.kd ; CHECK: .symbol: test_vec_type_hint_short.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_vec_type_hint_long ; CHECK: .name: test_vec_type_hint_long
; CHECK: .symbol: test_vec_type_hint_long.kd ; CHECK: .symbol: test_vec_type_hint_long.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_vec_type_hint_unknown ; CHECK: .name: test_vec_type_hint_unknown
; CHECK: .symbol: test_vec_type_hint_unknown.kd ; CHECK: .symbol: test_vec_type_hint_unknown.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_reqd_wgs_vec_type_hint ; CHECK: .name: test_reqd_wgs_vec_type_hint
; CHECK: .reqd_workgroup_size: ; CHECK: .reqd_workgroup_size:
Show All 19 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_wgs_hint_vec_type_hint ; CHECK: .name: test_wgs_hint_vec_type_hint
; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd ; CHECK: .symbol: test_wgs_hint_vec_type_hint.kd
Show All 20 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_arg_ptr_to_ptr ; CHECK: .name: test_arg_ptr_to_ptr
; CHECK: .symbol: test_arg_ptr_to_ptr.kd ; CHECK: .symbol: test_arg_ptr_to_ptr.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_arg_struct_contains_ptr ; CHECK: .name: test_arg_struct_contains_ptr
; CHECK: .symbol: test_arg_struct_contains_ptr.kd ; CHECK: .symbol: test_arg_struct_contains_ptr.kd
Show All 13 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_arg_vector_of_ptr ; CHECK: .name: test_arg_vector_of_ptr
; CHECK: .symbol: test_arg_vector_of_ptr.kd ; CHECK: .symbol: test_arg_vector_of_ptr.kd
Show All 14 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_arg_unknown_builtin_type ; CHECK: .name: test_arg_unknown_builtin_type
; CHECK: .symbol: test_arg_unknown_builtin_type.kd ; CHECK: .symbol: test_arg_unknown_builtin_type.kd
▲ Show 20 Lines • Show All 63 Lines • ▼ Show 20 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 48 ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 56 ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 88
; CHECK-NEXT: .offset: 88
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_pointee_align ; CHECK: .name: test_pointee_align
; CHECK: .symbol: test_pointee_align.kd ; CHECK: .symbol: test_pointee_align.kd
▲ Show 20 Lines • Show All 69 Lines • ▼ Show 20 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 48 ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 56 ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 88
; CHECK-NEXT: .offset: 88
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_pointee_align_attribute ; CHECK: .name: test_pointee_align_attribute
; CHECK: .symbol: test_pointee_align_attribute.kd ; CHECK: .symbol: test_pointee_align_attribute.kd
Show All 19 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 48 ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle ; CHECK: .device_enqueue_symbol: __test_block_invoke_kernel_runtime_handle
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: __test_block_invoke_kernel ; CHECK: .name: __test_block_invoke_kernel
Show All 15 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_default_queue ; CHECK-NEXT: .value_kind: hidden_default_queue
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_completion_action ; CHECK-NEXT: .value_kind: hidden_completion_action
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_enqueue_kernel_caller ; CHECK: .name: test_enqueue_kernel_caller
; CHECK: .symbol: test_enqueue_kernel_caller.kd ; CHECK: .symbol: test_enqueue_kernel_caller.kd
▲ Show 20 Lines • Show All 85 Lines • Show Last 20 Lines

llvm/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v3.ll

Show First 20 Lines • Show All 162 Lines • ▼ Show 20 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_hostcall_buffer ; CHECK-NEXT: .value_kind: hidden_hostcall_buffer
; CHECK: .name: test32 ; CHECK: .name: test32
; CHECK: .symbol: test32.kd ; CHECK: .symbol: test32.kd
define amdgpu_kernel void @test32( define amdgpu_kernel void @test32(
half addrspace(1)* %r, half addrspace(1)* %r,
half addrspace(1)* %a, half addrspace(1)* %a,
half addrspace(1)* %b) #3 { half addrspace(1)* %b) #3 {
Show All 25 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_hostcall_buffer ; CHECK-NEXT: .value_kind: hidden_hostcall_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK: .name: test48 ; CHECK: .name: test48
; CHECK: .symbol: test48.kd ; CHECK: .symbol: test48.kd
define amdgpu_kernel void @test48( define amdgpu_kernel void @test48(
half addrspace(1)* %r, half addrspace(1)* %r,
half addrspace(1)* %a, half addrspace(1)* %a,
half addrspace(1)* %b) #4 { half addrspace(1)* %b) #4 {
Show All 25 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 32 ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 40 ; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 48
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_hostcall_buffer ; CHECK-NEXT: .value_kind: hidden_hostcall_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 56
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 64
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none ; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 72
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK: .name: test56 ; CHECK: .name: test56
; CHECK: .symbol: test56.kd ; CHECK: .symbol: test56.kd
define amdgpu_kernel void @test56( define amdgpu_kernel void @test56(
half addrspace(1)* %r, half addrspace(1)* %r,
half addrspace(1)* %a, half addrspace(1)* %a,
half addrspace(1)* %b) #5 { half addrspace(1)* %b) #5 {
Show All 20 Lines

llvm/test/CodeGen/AMDGPU/hsa-metadata-hidden-args-v5.ll

Show First 20 Lines • Show All 57 Lines • ▼ Show 20 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 80 ; CHECK-NEXT: - .offset: 80
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .offset: 88 ; CHECK-NEXT: - .offset: 88
; CHECK-NEXT: .size: 2 ; CHECK-NEXT: .size: 2
; CHECK-NEXT: .value_kind: hidden_grid_dims ; CHECK-NEXT: .value_kind: hidden_grid_dims
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 96
; CHECK-NEXT: .offset: 96
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 104
; CHECK-NEXT: .offset: 104
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_hostcall_buffer ; CHECK-NEXT: .value_kind: hidden_hostcall_buffer
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 112
; CHECK-NEXT: .offset: 112
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg ; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 120
; CHECK-NEXT: .offset: 120
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_heap_v1 ; CHECK-NEXT: .value_kind: hidden_heap_v1
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 128
; CHECK-NEXT: .offset: 128
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_default_queue ; CHECK-NEXT: .value_kind: hidden_default_queue
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 136
; CHECK-NEXT: .offset: 136
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_completion_action ; CHECK-NEXT: .value_kind: hidden_completion_action
; GFX8-NEXT: - .offset: 216 ; GFX8-NEXT: - .offset: 216
; GFX8-NEXT: .size: 4 ; GFX8-NEXT: .size: 4
; GFX8-NEXT: .value_kind: hidden_private_base ; GFX8-NEXT: .value_kind: hidden_private_base
; GFX8-NEXT: - .offset: 220 ; GFX8-NEXT: - .offset: 220
; GFX8-NEXT: .size: 4 ; GFX8-NEXT: .size: 4
; GFX8-NEXT: .value_kind: hidden_shared_base ; GFX8-NEXT: .value_kind: hidden_shared_base
; CHECK: - .address_space: global ; CHECK: - .offset: 224
; CHECK-NEXT: .offset: 224
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_queue_ptr ; CHECK-NEXT: .value_kind: hidden_queue_ptr
; CHECK: .name: test_v5 ; CHECK: .name: test_v5
; CHECK: .symbol: test_v5.kd ; CHECK: .symbol: test_v5.kd
; CHECK: amdhsa.version: ; CHECK: amdhsa.version:
; CHECK-NEXT: - 1 ; CHECK-NEXT: - 1
Show All 20 Lines

llvm/test/CodeGen/AMDGPU/hsa-metadata-hostcall-present-v3-asan.ll

Show All 11 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x ; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: - .offset: 16 ; CHECK-NEXT: - .offset: 16
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y ; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: - .offset: 24 ; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z ; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: - .address_space: global ; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .offset: 32
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_hostcall_buffer ; CHECK-NEXT: .value_kind: hidden_hostcall_buffer
; CHECK: .language: OpenCL C ; CHECK: .language: OpenCL C
; CHECK-NEXT: .language_version: ; CHECK-NEXT: .language_version:
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
; CHECK: .name: test_kernel ; CHECK: .name: test_kernel
; CHECK: .symbol: test_kernel.kd ; CHECK: .symbol: test_kernel.kd
Show All 22 Lines

llvm/test/CodeGen/AMDGPU/hsa-metadata-images-v3.ll

Show All 12 Lines
%opencl.image2d_depth_t = type opaque %opencl.image2d_depth_t = type opaque
%opencl.image2d_msaa_t = type opaque %opencl.image2d_msaa_t = type opaque
%opencl.image2d_msaa_depth_t = type opaque %opencl.image2d_msaa_depth_t = type opaque
%opencl.image3d_t = type opaque %opencl.image3d_t = type opaque
; CHECK: --- ; CHECK: ---
; CHECK: amdhsa.kernels: ; CHECK: amdhsa.kernels:
; CHECK: - .args: ; CHECK: - .args:
; CHECK: - .address_space: global ; CHECK: - .name: a
; CHECK: .name: a
; CHECK: .offset: 0 ; CHECK: .offset: 0
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image1d_t ; CHECK: .type_name: image1d_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: b
; CHECK: .name: b
; CHECK: .offset: 8 ; CHECK: .offset: 8
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image1d_array_t ; CHECK: .type_name: image1d_array_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: c
; CHECK: .name: c
; CHECK: .offset: 16 ; CHECK: .offset: 16
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image1d_buffer_t ; CHECK: .type_name: image1d_buffer_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: d
; CHECK: .name: d
; CHECK: .offset: 24 ; CHECK: .offset: 24
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_t ; CHECK: .type_name: image2d_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: e
; CHECK: .name: e
; CHECK: .offset: 32 ; CHECK: .offset: 32
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_array_t ; CHECK: .type_name: image2d_array_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: f
; CHECK: .name: f
; CHECK: .offset: 40 ; CHECK: .offset: 40
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_array_depth_t ; CHECK: .type_name: image2d_array_depth_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: g
; CHECK: .name: g
; CHECK: .offset: 48 ; CHECK: .offset: 48
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_array_msaa_t ; CHECK: .type_name: image2d_array_msaa_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: h
; CHECK: .name: h
; CHECK: .offset: 56 ; CHECK: .offset: 56
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_array_msaa_depth_t ; CHECK: .type_name: image2d_array_msaa_depth_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: i
; CHECK: .name: i
; CHECK: .offset: 64 ; CHECK: .offset: 64
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_depth_t ; CHECK: .type_name: image2d_depth_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: j
; CHECK: .name: j
; CHECK: .offset: 72 ; CHECK: .offset: 72
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_msaa_t ; CHECK: .type_name: image2d_msaa_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: k
; CHECK: .name: k
; CHECK: .offset: 80 ; CHECK: .offset: 80
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image2d_msaa_depth_t ; CHECK: .type_name: image2d_msaa_depth_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
; CHECK: - .address_space: global ; CHECK: - .name: l
; CHECK: .name: l
; CHECK: .offset: 88 ; CHECK: .offset: 88
; CHECK: .size: 8 ; CHECK: .size: 8
; CHECK: .type_name: image3d_t ; CHECK: .type_name: image3d_t
; CHECK: .value_kind: image ; CHECK: .value_kind: image
define amdgpu_kernel void @test(%opencl.image1d_t addrspace(1)* %a, define amdgpu_kernel void @test(%opencl.image1d_t addrspace(1)* %a,
%opencl.image1d_array_t addrspace(1)* %b, %opencl.image1d_array_t addrspace(1)* %b,
%opencl.image1d_buffer_t addrspace(1)* %c, %opencl.image1d_buffer_t addrspace(1)* %c,
%opencl.image2d_t addrspace(1)* %d, %opencl.image2d_t addrspace(1)* %d,
Show All 21 Lines