This is an archive of the discontinued LLVM Phabricator instance.

Differential D63886

[AMDGPU] Added a new metadata for multi grid sync implicit argument.
ClosedPublic

Authored by cdevadas on Jun 27 2019, 9:29 AM.

Details

Reviewers
yaxunl
scott.linder
Commits
rGa62413526d80: [AMDGPU] Added a new metadata for multi grid sync implicit argument
rL365217: [AMDGPU] Added a new metadata for multi grid sync implicit argument
Summary

The metadata has been supported for both V2 and V3 code object.

Diff Detail

Event Timeline

cdevadas created this revision.Jun 27 2019, 9:29 AM
Herald added a project: Restricted Project. · View Herald TranscriptJun 27 2019, 9:29 AM
Herald added subscribers: llvm-commits, kristina, t-tye and 7 others. · View Herald Transcript
cdevadas mentioned this in D63756: [AMDGPU] Increased the number of implicit argument bytes for both OpenCL and HIP (CLANG)..Jun 27 2019, 9:50 AM
scott.linder added a comment.Jul 1 2019, 1:09 PM

We should update the v2 and v3 metadata portions of docs/AMDGPUUsage.rst, see https://llvm.org/docs/AMDGPUUsage.html#code-object-metadata

cdevadas updated this revision to Diff 207478.Jul 1 2019, 11:30 PM

Updated the document, AMDGPUUsage.rst with the new metadata information.

scott.linder accepted this revision.Jul 3 2019, 10:04 AM
This revision is now accepted and ready to land.Jul 3 2019, 10:04 AM
yaxunl accepted this revision.Jul 3 2019, 1:39 PM

LGTM

Closed by commit rL365217: [AMDGPU] Added a new metadata for multi grid sync implicit argument (authored by yaxunl). · Explain WhyJul 5 2019, 9:09 AM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
docs/
12 lines
include/
llvm/
Support/
1 line
lib/
BinaryFormat/
1 line
Support/
2 lines
Target/
AMDGPU/
8 lines
test/
CodeGen/
AMDGPU/
488 lines
479 lines
66 lines
67 lines

Diff 207478

docs/AMDGPUUsage.rst

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 1,385 Lines • ▼ Show 20 Lines "ValueKind" string Required Kernel argument kind that
passed in the kernarg. passed in the kernarg.
"HiddenCompletionAction" "HiddenCompletionAction"
A global address space pointer A global address space pointer
to help link enqueued kernels into to help link enqueued kernels into
the ancestor tree for determining the ancestor tree for determining
when the parent kernel has finished. when the parent kernel has finished.
"HiddenMultiGridSyncArg"
A global address space pointer for
multi-grid synchronization is
passed in the kernarg.
"ValueType" string Required Kernel argument value type. Only "ValueType" string Required Kernel argument value type. Only
present if "ValueKind" is present if "ValueKind" is
"ByValue". For vector data "ByValue". For vector data
types, the value is for the types, the value is for the
element type. Values include: element type. Values include:
- "Struct" - "Struct"
- "I8" - "I8"
▲ Show 20 Lines • Show All 473 Lines • ▼ Show 20 Lines ".value_kind" string Required Kernel argument kind that
passed in the kernarg. passed in the kernarg.
"hidden_completion_action" "hidden_completion_action"
A global address space pointer A global address space pointer
to help link enqueued kernels into to help link enqueued kernels into
the ancestor tree for determining the ancestor tree for determining
when the parent kernel has finished. when the parent kernel has finished.
"hidden_multigrid_sync_arg"
A global address space pointer for
multi-grid synchronization is
passed in the kernarg.
".value_type" string Required Kernel argument value type. Only ".value_type" string Required Kernel argument value type. Only
present if ".value_kind" is present if ".value_kind" is
"by_value". For vector data "by_value". For vector data
types, the value is for the types, the value is for the
element type. Values include: element type. Values include:
- "struct" - "struct"
- "i8" - "i8"
▲ Show 20 Lines • Show All 3,800 Lines • ▼ Show 20 Lines .. table:: OpenCL kernel implicit arguments appended for AMDHSA OS
1 8 8 OpenCL Global Offset X 1 8 8 OpenCL Global Offset X
2 8 8 OpenCL Global Offset Y 2 8 8 OpenCL Global Offset Y
3 8 8 OpenCL Global Offset Z 3 8 8 OpenCL Global Offset Z
4 8 8 OpenCL address of printf buffer 4 8 8 OpenCL address of printf buffer
5 8 8 OpenCL address of virtual queue used by 5 8 8 OpenCL address of virtual queue used by
enqueue_kernel. enqueue_kernel.
6 8 8 OpenCL address of AqlWrap struct used by 6 8 8 OpenCL address of AqlWrap struct used by
enqueue_kernel. enqueue_kernel.
7 8 8 Pointer argument used for Multi-gird
synchronization.
======== ==== ========= =========================================== ======== ==== ========= ===========================================
.. _amdgpu-hcc: .. _amdgpu-hcc:
HCC HCC
--- ---
When the language is HCC the following differences occur: When the language is HCC the following differences occur:
▲ Show 20 Lines • Show All 783 Lines • Show Last 20 Lines

include/llvm/Support/AMDGPUMetadata.h

Show First 20 Lines • Show All 68 Lines • ▼ Show 20 Linesenum class ValueKind : uint8_t {
Queue = 6, Queue = 6,
HiddenGlobalOffsetX = 7, HiddenGlobalOffsetX = 7,
HiddenGlobalOffsetY = 8, HiddenGlobalOffsetY = 8,
HiddenGlobalOffsetZ = 9, HiddenGlobalOffsetZ = 9,
HiddenNone = 10, HiddenNone = 10,
HiddenPrintfBuffer = 11, HiddenPrintfBuffer = 11,
HiddenDefaultQueue = 12, HiddenDefaultQueue = 12,
HiddenCompletionAction = 13, HiddenCompletionAction = 13,
HiddenMultiGridSyncArg = 14,
Unknown = 0xff Unknown = 0xff
}; };
/// Value types. /// Value types.
enum class ValueType : uint8_t { enum class ValueType : uint8_t {
Struct = 0, Struct = 0,
I8 = 1, I8 = 1,
U8 = 2, U8 = 2,
▲ Show 20 Lines • Show All 428 Lines • Show Last 20 Lines

lib/BinaryFormat/AMDGPUMetadataVerifier.cpp

Show First 20 Lines • Show All 115 Lines • ▼ Show 20 Lines if (!verifyScalarEntry(ArgsMap, ".value_kind", true,
.Case("queue", true) .Case("queue", true)
.Case("hidden_global_offset_x", true) .Case("hidden_global_offset_x", true)
.Case("hidden_global_offset_y", true) .Case("hidden_global_offset_y", true)
.Case("hidden_global_offset_z", true) .Case("hidden_global_offset_z", true)
.Case("hidden_none", true) .Case("hidden_none", true)
.Case("hidden_printf_buffer", true) .Case("hidden_printf_buffer", true)
.Case("hidden_default_queue", true) .Case("hidden_default_queue", true)
.Case("hidden_completion_action", true) .Case("hidden_completion_action", true)
.Case("hidden_multigrid_sync_arg", true)
.Default(false); .Default(false);
})) }))
return false; return false;
if (!verifyScalarEntry(ArgsMap, ".value_type", true, if (!verifyScalarEntry(ArgsMap, ".value_type", true,
msgpack::Type::String, msgpack::Type::String,
[](msgpack::DocNode &SNode) { [](msgpack::DocNode &SNode) {
return StringSwitch<bool>(SNode.getString()) return StringSwitch<bool>(SNode.getString())
.Case("struct", true) .Case("struct", true)
▲ Show 20 Lines • Show All 184 Lines • Show Last 20 Lines

lib/Support/AMDGPUMetadata.cpp

Show First 20 Lines • Show All 59 Lines • ▼ Show 20 Lines static void enumeration(IO &YIO, ValueKind &EN) {
YIO.enumCase(EN, "HiddenGlobalOffsetX", ValueKind::HiddenGlobalOffsetX); YIO.enumCase(EN, "HiddenGlobalOffsetX", ValueKind::HiddenGlobalOffsetX);
YIO.enumCase(EN, "HiddenGlobalOffsetY", ValueKind::HiddenGlobalOffsetY); YIO.enumCase(EN, "HiddenGlobalOffsetY", ValueKind::HiddenGlobalOffsetY);
YIO.enumCase(EN, "HiddenGlobalOffsetZ", ValueKind::HiddenGlobalOffsetZ); YIO.enumCase(EN, "HiddenGlobalOffsetZ", ValueKind::HiddenGlobalOffsetZ);
YIO.enumCase(EN, "HiddenNone", ValueKind::HiddenNone); YIO.enumCase(EN, "HiddenNone", ValueKind::HiddenNone);
YIO.enumCase(EN, "HiddenPrintfBuffer", ValueKind::HiddenPrintfBuffer); YIO.enumCase(EN, "HiddenPrintfBuffer", ValueKind::HiddenPrintfBuffer);
YIO.enumCase(EN, "HiddenDefaultQueue", ValueKind::HiddenDefaultQueue); YIO.enumCase(EN, "HiddenDefaultQueue", ValueKind::HiddenDefaultQueue);
YIO.enumCase(EN, "HiddenCompletionAction", YIO.enumCase(EN, "HiddenCompletionAction",
ValueKind::HiddenCompletionAction); ValueKind::HiddenCompletionAction);
YIO.enumCase(EN, "HiddenMultiGridSyncArg",
ValueKind::HiddenMultiGridSyncArg);
} }
}; };
template <> template <>
struct ScalarEnumerationTraits<ValueType> { struct ScalarEnumerationTraits<ValueType> {
static void enumeration(IO &YIO, ValueType &EN) { static void enumeration(IO &YIO, ValueType &EN) {
YIO.enumCase(EN, "Struct", ValueType::Struct); YIO.enumCase(EN, "Struct", ValueType::Struct);
YIO.enumCase(EN, "I8", ValueType::I8); YIO.enumCase(EN, "I8", ValueType::I8);
▲ Show 20 Lines • Show All 147 Lines • Show Last 20 Lines

lib/Target/AMDGPU/AMDGPUHSAMetadataStreamer.cpp

Show First 20 Lines • Show All 429 Lines • ▼ Show 20 Lines if (HiddenArgNumBytes >= 48) {
if (Func.hasFnAttribute("calls-enqueue-kernel")) { if (Func.hasFnAttribute("calls-enqueue-kernel")) {
emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenDefaultQueue); emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenDefaultQueue);
emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenCompletionAction); emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenCompletionAction);
} else { } else {
emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone); emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone); emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenNone);
} }
} }
// Emit the pointer argument for multi-grid object.
if (HiddenArgNumBytes >= 56)
emitKernelArg(DL, Int8PtrTy, ValueKind::HiddenMultiGridSyncArg);
} }
bool MetadataStreamerV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) { bool MetadataStreamerV2::emitTo(AMDGPUTargetStreamer &TargetStreamer) {
return TargetStreamer.EmitHSAMetadata(getHSAMetadata()); return TargetStreamer.EmitHSAMetadata(getHSAMetadata());
} }
void MetadataStreamerV2::begin(const Module &Mod) { void MetadataStreamerV2::begin(const Module &Mod) {
emitVersion(); emitVersion();
▲ Show 20 Lines • Show All 413 Lines • ▼ Show 20 Lines if (HiddenArgNumBytes >= 48) {
if (Func.hasFnAttribute("calls-enqueue-kernel")) { if (Func.hasFnAttribute("calls-enqueue-kernel")) {
emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, Args); emitKernelArg(DL, Int8PtrTy, "hidden_default_queue", Offset, Args);
emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args); emitKernelArg(DL, Int8PtrTy, "hidden_completion_action", Offset, Args);
} else { } else {
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args); emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args); emitKernelArg(DL, Int8PtrTy, "hidden_none", Offset, Args);
} }
} }
// Emit the pointer argument for multi-grid object.
if (HiddenArgNumBytes >= 56)
emitKernelArg(DL, Int8PtrTy, "hidden_multigrid_sync_arg", Offset, Args);
} }
msgpack::MapDocNode msgpack::MapDocNode
MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF, MetadataStreamerV3::getHSAKernelProps(const MachineFunction &MF,
const SIProgramInfo &ProgramInfo) const { const SIProgramInfo &ProgramInfo) const {
const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>(); const GCNSubtarget &STM = MF.getSubtarget<GCNSubtarget>();
const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>(); const SIMachineFunctionInfo &MFI = *MF.getInfo<SIMachineFunctionInfo>();
const Function &F = MF.getFunction(); const Function &F = MF.getFunction();
▲ Show 20 Lines • Show All 73 Lines • Show Last 20 Lines

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

Show All 38 Lines
; CHECK-NEXT: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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_char ; CHECK: .name: test_char
; CHECK: .symbol: test_char.kd ; CHECK: .symbol: test_char.kd
define amdgpu_kernel void @test_char(i8 %a) #0 define amdgpu_kernel void @test_char(i8 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
Show All 25 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_printf_buffer ; CHECK-NEXT: .value_kind: hidden_printf_buffer
; CHECK-NEXT: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0 define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_int3(<3 x i32> %a) #0 define amdgpu_kernel void @test_int3(<3 x i32> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0 define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_half8(<8 x half> %a) #0 define amdgpu_kernel void @test_half8(<8 x half> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 96
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 104
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 112
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_float16(<16 x float> %a) #0 define amdgpu_kernel void @test_float16(<16 x float> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 160
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 168
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 176
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_double16(<16 x double> %a) #0 define amdgpu_kernel void @test_double16(<16 x double> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a) #0 define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a) #0 define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a) #0 define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_i128(i128 %a) #0 define amdgpu_kernel void @test_i128(i128 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21
Show All 32 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0 define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0
!kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24
Show All 36 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g, define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g,
i32 addrspace(4)* %c, i32 addrspace(4)* %c,
▲ Show 20 Lines • Show All 41 Lines • ▼ Show 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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,
Show All 40 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_access_qual(%opencl.image1d_t addrspace(1)* %ro, define amdgpu_kernel void @test_access_qual(%opencl.image1d_t addrspace(1)* %ro,
%opencl.image2d_t addrspace(1)* %wo, %opencl.image2d_t addrspace(1)* %wo,
Show All 22 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: half ; CHECK: .vec_type_hint: half
define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: float ; CHECK: .vec_type_hint: float
define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #0
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: double ; CHECK: .vec_type_hint: double
define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #0
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: char ; CHECK: .vec_type_hint: char
define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #0
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: short ; CHECK: .vec_type_hint: short
define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #0
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: long ; CHECK: .vec_type_hint: long
define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #0
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: unknown ; CHECK: .vec_type_hint: unknown
define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #0
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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:
; CHECK-NEXT: - 1 ; CHECK-NEXT: - 1
; CHECK-NEXT: - 2 ; CHECK-NEXT: - 2
Show All 26 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .vec_type_hint: uint4 ; CHECK: .vec_type_hint: uint4
; CHECK: .workgroup_size_hint: ; CHECK: .workgroup_size_hint:
Show All 27 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a) #0 define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a) #0
!kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B addrspace(5)* byval %a) #0 define amdgpu_kernel void @test_arg_struct_contains_ptr(%struct.B addrspace(5)* byval %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !82 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !82
Show All 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a) #0 define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83
Show All 21 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_arg_unknown_builtin_type( define amdgpu_kernel void @test_arg_unknown_builtin_type(
%opencl.clk_event_t addrspace(1)* %a) #0 %opencl.clk_event_t addrspace(1)* %a) #0
▲ Show 20 Lines • Show All 77 Lines • ▼ Show 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 88
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a, define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a,
i8 addrspace(3)* %b, i8 addrspace(3)* %b,
▲ Show 20 Lines • Show All 83 Lines • ▼ Show 20 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 88
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_pointee_align_attribute(i64 addrspace(1)* align 16 %a, define amdgpu_kernel void @test_pointee_align_attribute(i64 addrspace(1)* align 16 %a,
i8 addrspace(3)* align 8 %b, i8 addrspace(3)* align 8 %b,
Show All 26 Lines
; 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: .value_type: i64 ; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 80
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
; CHECK: .symbol: __test_block_invoke_kernel.kd ; CHECK: .symbol: __test_block_invoke_kernel.kd
define amdgpu_kernel void @__test_block_invoke_kernel( define amdgpu_kernel void @__test_block_invoke_kernel(
Show All 32 Lines
; CHECK-NEXT: .size: 8 ; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_default_queue ; CHECK-NEXT: .value_kind: hidden_default_queue
; CHECK-NEXT: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global ; 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: .value_type: i8 ; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; 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
define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
Show All 15 Lines
; CHECK: amdhsa.printf: ; CHECK: amdhsa.printf:
; CHECK-NEXT: - '1:1:4:%d\n' ; CHECK-NEXT: - '1:1:4:%d\n'
; CHECK-NEXT: - '2:1:8:%g\n' ; CHECK-NEXT: - '2:1:8:%g\n'
; CHECK: amdhsa.version: ; CHECK: amdhsa.version:
; CHECK-NEXT: - 1 ; CHECK-NEXT: - 1
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
attributes #0 = { "amdgpu-implicitarg-num-bytes"="48" } attributes #0 = { "amdgpu-implicitarg-num-bytes"="56" }
attributes #1 = { "amdgpu-implicitarg-num-bytes"="48" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" } attributes #1 = { "amdgpu-implicitarg-num-bytes"="56" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
attributes #2 = { "amdgpu-implicitarg-num-bytes"="48" "calls-enqueue-kernel" } attributes #2 = { "amdgpu-implicitarg-num-bytes"="56" "calls-enqueue-kernel" }
!llvm.printf.fmts = !{!100, !101} !llvm.printf.fmts = !{!100, !101}
!1 = !{i32 0} !1 = !{i32 0}
!2 = !{!"none"} !2 = !{!"none"}
!3 = !{!"int"} !3 = !{!"int"}
!4 = !{!""} !4 = !{!""}
!5 = !{i32 undef, i32 1} !5 = !{i32 undef, i32 1}
▲ Show 20 Lines • Show All 49 Lines • Show Last 20 Lines

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

Show First 20 Lines • Show All 48 Lines • ▼ Show 20 Lines
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NOT: ValueKind: HiddenDefaultQueue ; CHECK-NOT: ValueKind: HiddenDefaultQueue
; CHECK-NOT: ValueKind: HiddenCompletionAction ; CHECK-NOT: ValueKind: HiddenCompletionAction
; CHECK: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_char(i8 %a) #0 define amdgpu_kernel void @test_char(i8 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
!kernel_arg_base_type !9 !kernel_arg_type_qual !4 { !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_ushort2 ; CHECK: - Name: test_ushort2
; CHECK-NEXT: SymbolName: 'test_ushort2@kd' ; CHECK-NEXT: SymbolName: 'test_ushort2@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0 define amdgpu_kernel void @test_ushort2(<2 x i16> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !10
!kernel_arg_base_type !10 !kernel_arg_type_qual !4 { !kernel_arg_base_type !10 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_int3 ; CHECK: - Name: test_int3
; CHECK-NEXT: SymbolName: 'test_int3@kd' ; CHECK-NEXT: SymbolName: 'test_int3@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_int3(<3 x i32> %a) #0 define amdgpu_kernel void @test_int3(<3 x i32> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !11
!kernel_arg_base_type !11 !kernel_arg_type_qual !4 { !kernel_arg_base_type !11 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_ulong4 ; CHECK: - Name: test_ulong4
; CHECK-NEXT: SymbolName: 'test_ulong4@kd' ; CHECK-NEXT: SymbolName: 'test_ulong4@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0 define amdgpu_kernel void @test_ulong4(<4 x i64> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !12
!kernel_arg_base_type !12 !kernel_arg_type_qual !4 { !kernel_arg_base_type !12 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_half8 ; CHECK: - Name: test_half8
; CHECK-NEXT: SymbolName: 'test_half8@kd' ; CHECK-NEXT: SymbolName: 'test_half8@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_half8(<8 x half> %a) #0 define amdgpu_kernel void @test_half8(<8 x half> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !13
!kernel_arg_base_type !13 !kernel_arg_type_qual !4 { !kernel_arg_base_type !13 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_float16 ; CHECK: - Name: test_float16
; CHECK-NEXT: SymbolName: 'test_float16@kd' ; CHECK-NEXT: SymbolName: 'test_float16@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_float16(<16 x float> %a) #0 define amdgpu_kernel void @test_float16(<16 x float> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !14
!kernel_arg_base_type !14 !kernel_arg_type_qual !4 { !kernel_arg_base_type !14 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_double16 ; CHECK: - Name: test_double16
; CHECK-NEXT: SymbolName: 'test_double16@kd' ; CHECK-NEXT: SymbolName: 'test_double16@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_double16(<16 x double> %a) #0 define amdgpu_kernel void @test_double16(<16 x double> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !15
!kernel_arg_base_type !15 !kernel_arg_type_qual !4 { !kernel_arg_base_type !15 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_pointer ; CHECK: - Name: test_pointer
; CHECK-NEXT: SymbolName: 'test_pointer@kd' ; CHECK-NEXT: SymbolName: 'test_pointer@kd'
Show All 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
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: - Name: test_image ; CHECK: - Name: test_image
; CHECK-NEXT: SymbolName: 'test_image@kd' ; CHECK-NEXT: SymbolName: 'test_image@kd'
Show All 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a) #0 define amdgpu_kernel void @test_image(%opencl.image2d_t addrspace(1)* %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !17
!kernel_arg_base_type !17 !kernel_arg_type_qual !4 { !kernel_arg_base_type !17 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_sampler ; CHECK: - Name: test_sampler
; CHECK-NEXT: SymbolName: 'test_sampler@kd' ; CHECK-NEXT: SymbolName: 'test_sampler@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
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: - Name: test_queue ; CHECK: - Name: test_queue
; CHECK-NEXT: SymbolName: 'test_queue@kd' ; CHECK-NEXT: SymbolName: 'test_queue@kd'
Show All 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a) #0 define amdgpu_kernel void @test_queue(%opencl.queue_t addrspace(1)* %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !19
!kernel_arg_base_type !19 !kernel_arg_type_qual !4 { !kernel_arg_base_type !19 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_struct ; CHECK: - Name: test_struct
; CHECK-NEXT: SymbolName: 'test_struct@kd' ; CHECK-NEXT: SymbolName: 'test_struct@kd'
Show All 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a) #0 define amdgpu_kernel void @test_struct(%struct.A addrspace(5)* byval %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !20
!kernel_arg_base_type !20 !kernel_arg_type_qual !4 { !kernel_arg_base_type !20 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_i128 ; CHECK: - Name: test_i128
; CHECK-NEXT: SymbolName: 'test_i128@kd' ; CHECK-NEXT: SymbolName: 'test_i128@kd'
Show All 19 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_i128(i128 %a) #0 define amdgpu_kernel void @test_i128(i128 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !21
!kernel_arg_base_type !21 !kernel_arg_type_qual !4 { !kernel_arg_base_type !21 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_multi_arg ; CHECK: - Name: test_multi_arg
; CHECK-NEXT: SymbolName: 'test_multi_arg@kd' ; CHECK-NEXT: SymbolName: 'test_multi_arg@kd'
Show All 33 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0 define amdgpu_kernel void @test_multi_arg(i32 %a, <2 x i16> %b, <3 x i8> %c) #0
!kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24 !kernel_arg_addr_space !22 !kernel_arg_access_qual !23 !kernel_arg_type !24
!kernel_arg_base_type !24 !kernel_arg_type_qual !25 { !kernel_arg_base_type !24 !kernel_arg_type_qual !25 {
ret void ret void
} }
; CHECK: - Name: test_addr_space ; CHECK: - Name: test_addr_space
; CHECK-NEXT: SymbolName: 'test_addr_space@kd' ; CHECK-NEXT: SymbolName: 'test_addr_space@kd'
Show All 37 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g, define amdgpu_kernel void @test_addr_space(i32 addrspace(1)* %g,
i32 addrspace(4)* %c, i32 addrspace(4)* %c,
i32 addrspace(3)* %l) #0 i32 addrspace(3)* %l) #0
!kernel_arg_addr_space !50 !kernel_arg_access_qual !23 !kernel_arg_type !51 !kernel_arg_addr_space !50 !kernel_arg_access_qual !23 !kernel_arg_type !51
!kernel_arg_base_type !51 !kernel_arg_type_qual !25 { !kernel_arg_base_type !51 !kernel_arg_type_qual !25 {
ret void ret void
} }
▲ Show 20 Lines • Show All 42 Lines • ▼ Show 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
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
} }
Show All 38 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_access_qual(%opencl.image1d_t addrspace(1)* %ro, define amdgpu_kernel void @test_access_qual(%opencl.image1d_t addrspace(1)* %ro,
%opencl.image2d_t addrspace(1)* %wo, %opencl.image2d_t addrspace(1)* %wo,
%opencl.image3d_t addrspace(1)* %rw) #0 %opencl.image3d_t addrspace(1)* %rw) #0
!kernel_arg_addr_space !60 !kernel_arg_access_qual !61 !kernel_arg_type !62 !kernel_arg_addr_space !60 !kernel_arg_access_qual !61 !kernel_arg_type !62
!kernel_arg_base_type !62 !kernel_arg_type_qual !25 { !kernel_arg_base_type !62 !kernel_arg_type_qual !25 {
ret void ret void
} }
Show All 23 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_half(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !26 { !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !26 {
ret void ret void
} }
; CHECK: - Name: test_vec_type_hint_float ; CHECK: - Name: test_vec_type_hint_float
; CHECK-NEXT: SymbolName: 'test_vec_type_hint_float@kd' ; CHECK-NEXT: SymbolName: 'test_vec_type_hint_float@kd'
Show All 21 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_float(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !27 { !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !27 {
ret void ret void
} }
; CHECK: - Name: test_vec_type_hint_double ; CHECK: - Name: test_vec_type_hint_double
; CHECK-NEXT: SymbolName: 'test_vec_type_hint_double@kd' ; CHECK-NEXT: SymbolName: 'test_vec_type_hint_double@kd'
Show All 21 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_double(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !28 { !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !28 {
ret void ret void
} }
; CHECK: - Name: test_vec_type_hint_char ; CHECK: - Name: test_vec_type_hint_char
; CHECK-NEXT: SymbolName: 'test_vec_type_hint_char@kd' ; CHECK-NEXT: SymbolName: 'test_vec_type_hint_char@kd'
Show All 21 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_char(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !29 { !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !29 {
ret void ret void
} }
; CHECK: - Name: test_vec_type_hint_short ; CHECK: - Name: test_vec_type_hint_short
; CHECK-NEXT: SymbolName: 'test_vec_type_hint_short@kd' ; CHECK-NEXT: SymbolName: 'test_vec_type_hint_short@kd'
Show All 21 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_short(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !30 { !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !30 {
ret void ret void
} }
; CHECK: - Name: test_vec_type_hint_long ; CHECK: - Name: test_vec_type_hint_long
; CHECK-NEXT: SymbolName: 'test_vec_type_hint_long@kd' ; CHECK-NEXT: SymbolName: 'test_vec_type_hint_long@kd'
Show All 21 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_long(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !31 { !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !31 {
ret void ret void
} }
; CHECK: - Name: test_vec_type_hint_unknown ; CHECK: - Name: test_vec_type_hint_unknown
; CHECK-NEXT: SymbolName: 'test_vec_type_hint_unknown@kd' ; CHECK-NEXT: SymbolName: 'test_vec_type_hint_unknown@kd'
Show All 21 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #0 define amdgpu_kernel void @test_vec_type_hint_unknown(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !32 { !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !32 {
ret void ret void
} }
; CHECK: - Name: test_reqd_wgs_vec_type_hint ; CHECK: - Name: test_reqd_wgs_vec_type_hint
; CHECK-NEXT: SymbolName: 'test_reqd_wgs_vec_type_hint@kd' ; CHECK-NEXT: SymbolName: 'test_reqd_wgs_vec_type_hint@kd'
Show All 22 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a) #0 define amdgpu_kernel void @test_reqd_wgs_vec_type_hint(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !5 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !5
!reqd_work_group_size !6 { !reqd_work_group_size !6 {
ret void ret void
} }
; CHECK: - Name: test_wgs_hint_vec_type_hint ; CHECK: - Name: test_wgs_hint_vec_type_hint
Show All 23 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a) #0 define amdgpu_kernel void @test_wgs_hint_vec_type_hint(i32 %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !3
!kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !7 !kernel_arg_base_type !3 !kernel_arg_type_qual !4 !vec_type_hint !7
!work_group_size_hint !8 { !work_group_size_hint !8 {
ret void ret void
} }
; CHECK: - Name: test_arg_ptr_to_ptr ; CHECK: - Name: test_arg_ptr_to_ptr
Show All 21 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a) #0 define amdgpu_kernel void @test_arg_ptr_to_ptr(i32 addrspace(5)* addrspace(1)* %a) #0
!kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !80
!kernel_arg_base_type !80 !kernel_arg_type_qual !4 { !kernel_arg_base_type !80 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_arg_struct_contains_ptr ; CHECK: - Name: test_arg_struct_contains_ptr
; CHECK-NEXT: SymbolName: 'test_arg_struct_contains_ptr@kd' ; CHECK-NEXT: SymbolName: 'test_arg_struct_contains_ptr@kd'
▲ Show 20 Lines • Show All 55 Lines • ▼ Show 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a) #0 define amdgpu_kernel void @test_arg_vector_of_ptr(<2 x i32 addrspace(1)*> %a) #0
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !83
!kernel_arg_base_type !83 !kernel_arg_type_qual !4 { !kernel_arg_base_type !83 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_arg_unknown_builtin_type ; CHECK: - Name: test_arg_unknown_builtin_type
; CHECK-NEXT: SymbolName: 'test_arg_unknown_builtin_type@kd' ; CHECK-NEXT: SymbolName: 'test_arg_unknown_builtin_type@kd'
Show All 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_arg_unknown_builtin_type( define amdgpu_kernel void @test_arg_unknown_builtin_type(
%opencl.clk_event_t addrspace(1)* %a) #0 %opencl.clk_event_t addrspace(1)* %a) #0
!kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !84 !kernel_arg_addr_space !81 !kernel_arg_access_qual !2 !kernel_arg_type !84
!kernel_arg_base_type !84 !kernel_arg_type_qual !4 { !kernel_arg_base_type !84 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_pointee_align ; CHECK: - Name: test_pointee_align
▲ Show 20 Lines • Show All 82 Lines • ▼ Show 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a, define amdgpu_kernel void @test_pointee_align(i64 addrspace(1)* %a,
i8 addrspace(3)* %b, i8 addrspace(3)* %b,
<2 x i8> addrspace(3)* %c, <2 x i8> addrspace(3)* %c,
<3 x i8> addrspace(3)* %d, <3 x i8> addrspace(3)* %d,
<4 x i8> addrspace(3)* %e, <4 x i8> addrspace(3)* %e,
<8 x i8> addrspace(3)* %f, <8 x i8> addrspace(3)* %f,
<16 x i8> addrspace(3)* %g, <16 x i8> addrspace(3)* %g,
{} addrspace(3)* %h) #0 {} addrspace(3)* %h) #0
▲ Show 20 Lines • Show All 88 Lines • ▼ Show 20 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_pointee_align_attribute(i64 addrspace(1)* align 16 %a, define amdgpu_kernel void @test_pointee_align_attribute(i64 addrspace(1)* align 16 %a,
i8 addrspace(3)* align 8 %b, i8 addrspace(3)* align 8 %b,
<2 x i8> addrspace(3)* align 32 %c, <2 x i8> addrspace(3)* align 32 %c,
<3 x i8> addrspace(3)* align 64 %d, <3 x i8> addrspace(3)* align 64 %d,
<4 x i8> addrspace(3)* align 256 %e, <4 x i8> addrspace(3)* align 256 %e,
<8 x i8> addrspace(3)* align 128 %f, <8 x i8> addrspace(3)* align 128 %f,
<16 x i8> addrspace(3)* align 1024 %g, <16 x i8> addrspace(3)* align 1024 %g,
{} addrspace(3)* align 16 %h) #0 {} addrspace(3)* align 16 %h) #0
Show All 29 Lines
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ ; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64 ; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenPrintfBuffer ; CHECK-NEXT: ValueKind: HiddenPrintfBuffer
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @__test_block_invoke_kernel( define amdgpu_kernel void @__test_block_invoke_kernel(
<{ i32, i32, i8*, i8 addrspace(1)*, i8 }> %arg) #1 <{ i32, i32, i8*, i8 addrspace(1)*, i8 }> %arg) #1
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !110
!kernel_arg_base_type !110 !kernel_arg_type_qual !4 { !kernel_arg_base_type !110 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: test_enqueue_kernel_caller ; CHECK: - Name: test_enqueue_kernel_caller
Show All 30 Lines
; CHECK-NEXT: ValueKind: HiddenDefaultQueue ; CHECK-NEXT: ValueKind: HiddenDefaultQueue
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8 ; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenCompletionAction ; CHECK-NEXT: ValueKind: HiddenCompletionAction
; CHECK-NEXT: ValueType: I8 ; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global ; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2 define amdgpu_kernel void @test_enqueue_kernel_caller(i8 %a) #2
!kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9 !kernel_arg_addr_space !1 !kernel_arg_access_qual !2 !kernel_arg_type !9
!kernel_arg_base_type !9 !kernel_arg_type_qual !4 { !kernel_arg_base_type !9 !kernel_arg_type_qual !4 {
ret void ret void
} }
; CHECK: - Name: unknown_addrspace_kernarg ; CHECK: - Name: unknown_addrspace_kernarg
; CHECK: Args: ; CHECK: Args:
; CHECK-NEXT: - Name: ptr ; CHECK-NEXT: - Name: ptr
; CHECK-NEXT: Size: 8 ; CHECK-NEXT: Size: 8
; CHECK-NEXT: Align: 8 ; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: GlobalBuffer ; CHECK-NEXT: ValueKind: GlobalBuffer
; CHECK-NEXT: ValueType: I32 ; CHECK-NEXT: ValueType: I32
define amdgpu_kernel void @unknown_addrspace_kernarg(i32 addrspace(12345)* %ptr) #0 { define amdgpu_kernel void @unknown_addrspace_kernarg(i32 addrspace(12345)* %ptr) #0 {
ret void ret void
} }
attributes #0 = { "amdgpu-implicitarg-num-bytes"="48" } attributes #0 = { "amdgpu-implicitarg-num-bytes"="56" }
attributes #1 = { "amdgpu-implicitarg-num-bytes"="48" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" } attributes #1 = { "amdgpu-implicitarg-num-bytes"="56" "runtime-handle"="__test_block_invoke_kernel_runtime_handle" }
attributes #2 = { "amdgpu-implicitarg-num-bytes"="48" "calls-enqueue-kernel" } attributes #2 = { "amdgpu-implicitarg-num-bytes"="56" "calls-enqueue-kernel" }
!llvm.printf.fmts = !{!100, !101} !llvm.printf.fmts = !{!100, !101}
!1 = !{i32 0} !1 = !{i32 0}
!2 = !{!"none"} !2 = !{!"none"}
!3 = !{!"int"} !3 = !{!"int"}
!4 = !{!""} !4 = !{!""}
!5 = !{i32 undef, i32 1} !5 = !{i32 undef, i32 1}
▲ Show 20 Lines • Show All 49 Lines • Show Last 20 Lines

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

Show First 20 Lines • Show All 265 Lines • ▼ Show 20 Lines
entry: entry:
%a.val = load half, half addrspace(1)* %a %a.val = load half, half addrspace(1)* %a
%b.val = load half, half addrspace(1)* %b %b.val = load half, half addrspace(1)* %b
%r.val = fadd half %a.val, %b.val %r.val = fadd half %a.val, %b.val
store half %r.val, half addrspace(1)* %r store half %r.val, half addrspace(1)* %r
ret void ret void
} }
; CHECK: - .args:
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .name: r
; CHECK-NEXT: .offset: 0
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .name: a
; CHECK-NEXT: .offset: 8
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .name: b
; CHECK-NEXT: .offset: 16
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: global_buffer
; CHECK-NEXT: .value_type: f16
; CHECK-NEXT: - .offset: 24
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_x
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 32
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_y
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .offset: 40
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_global_offset_z
; CHECK-NEXT: .value_type: i64
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 48
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 56
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 64
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_none
; CHECK-NEXT: .value_type: i8
; CHECK-NEXT: - .address_space: global
; CHECK-NEXT: .offset: 72
; CHECK-NEXT: .size: 8
; CHECK-NEXT: .value_kind: hidden_multigrid_sync_arg
; CHECK-NEXT: .value_type: i8
; CHECK: .name: test56
; CHECK: .symbol: test56.kd
define amdgpu_kernel void @test56(
half addrspace(1)* %r,
half addrspace(1)* %a,
half addrspace(1)* %b) #5 {
entry:
%a.val = load half, half addrspace(1)* %a
%b.val = load half, half addrspace(1)* %b
%r.val = fadd half %a.val, %b.val
store half %r.val, half addrspace(1)* %r
ret void
}
; CHECK: amdhsa.version: ; CHECK: amdhsa.version:
; CHECK-NEXT: - 1 ; CHECK-NEXT: - 1
; CHECK-NEXT: - 0 ; CHECK-NEXT: - 0
attributes #0 = { "amdgpu-implicitarg-num-bytes"="8" } attributes #0 = { "amdgpu-implicitarg-num-bytes"="8" }
attributes #1 = { "amdgpu-implicitarg-num-bytes"="16" } attributes #1 = { "amdgpu-implicitarg-num-bytes"="16" }
attributes #2 = { "amdgpu-implicitarg-num-bytes"="24" } attributes #2 = { "amdgpu-implicitarg-num-bytes"="24" }
attributes #3 = { "amdgpu-implicitarg-num-bytes"="32" } attributes #3 = { "amdgpu-implicitarg-num-bytes"="32" }
attributes #4 = { "amdgpu-implicitarg-num-bytes"="48" } attributes #4 = { "amdgpu-implicitarg-num-bytes"="48" }
attributes #5 = { "amdgpu-implicitarg-num-bytes"="56" }

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

Show First 20 Lines • Show All 272 Lines • ▼ Show 20 Lines
entry: entry:
%a.val = load half, half addrspace(1)* %a %a.val = load half, half addrspace(1)* %a
%b.val = load half, half addrspace(1)* %b %b.val = load half, half addrspace(1)* %b
%r.val = fadd half %a.val, %b.val %r.val = fadd half %a.val, %b.val
store half %r.val, half addrspace(1)* %r store half %r.val, half addrspace(1)* %r
ret void ret void
} }
; CHECK: - Name: test56
; CHECK: SymbolName: 'test56@kd'
; CHECK: Args:
; CHECK-NEXT: - Name: r
; CHECK-NEXT: Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: GlobalBuffer
; CHECK-NEXT: ValueType: F16
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Name: a
; CHECK-NEXT: Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: GlobalBuffer
; CHECK-NEXT: ValueType: F16
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Name: b
; CHECK-NEXT: Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: GlobalBuffer
; CHECK-NEXT: ValueType: F16
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetX
; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetY
; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenGlobalOffsetZ
; CHECK-NEXT: ValueType: I64
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenNone
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: - Size: 8
; CHECK-NEXT: Align: 8
; CHECK-NEXT: ValueKind: HiddenMultiGridSyncArg
; CHECK-NEXT: ValueType: I8
; CHECK-NEXT: AddrSpaceQual: Global
; CHECK-NEXT: CodeProps:
define amdgpu_kernel void @test56(
half addrspace(1)* %r,
half addrspace(1)* %a,
half addrspace(1)* %b) #5 {
entry:
%a.val = load half, half addrspace(1)* %a
%b.val = load half, half addrspace(1)* %b
%r.val = fadd half %a.val, %b.val
store half %r.val, half addrspace(1)* %r
ret void
}
attributes #0 = { "amdgpu-implicitarg-num-bytes"="8" } attributes #0 = { "amdgpu-implicitarg-num-bytes"="8" }
attributes #1 = { "amdgpu-implicitarg-num-bytes"="16" } attributes #1 = { "amdgpu-implicitarg-num-bytes"="16" }
attributes #2 = { "amdgpu-implicitarg-num-bytes"="24" } attributes #2 = { "amdgpu-implicitarg-num-bytes"="24" }
attributes #3 = { "amdgpu-implicitarg-num-bytes"="32" } attributes #3 = { "amdgpu-implicitarg-num-bytes"="32" }
attributes #4 = { "amdgpu-implicitarg-num-bytes"="48" } attributes #4 = { "amdgpu-implicitarg-num-bytes"="48" }
attributes #5 = { "amdgpu-implicitarg-num-bytes"="56" }