Skip to content

Commit b01235c

Browse files
steffenlarsensteffenlarsen
authored
[SYCL][sycl-post-link] Generate global_id_mapping for device_global name mapping (#7799)
Some backends do not natively support the HostAccessINTEL decoration and instead need to know the name of the device_global in the device program. However, the SYCL runtime will only be informed about the unique name, so this commit adds program metadata for making the mapping between unique names and actual names in the programs, allowing the backends to make the mappings internally. Split from #7797. Signed-off-by: Larsen, Steffen <[email protected]>
1 parent b423e1d commit b01235c

File tree

3 files changed

+89
-0
lines changed

3 files changed

+89
-0
lines changed

‎llvm/include/llvm/Support/PropertySetIO.h

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -85,6 +85,9 @@ class PropertyValue {
8585
PropertyValue(const std::vector<T> &Data)
8686
: PropertyValue(reinterpret_cast<const byte *>(Data.data()),
8787
Data.size() * sizeof(T) * /* bits in one byte */ 8) {}
88+
PropertyValue(const llvm::StringRef &Str)
89+
: PropertyValue(reinterpret_cast<const byte *>(Str.data()),
90+
Str.size() * sizeof(char) * /* bits in one byte */ 8) {}
8891
PropertyValue(const PropertyValue &P);
8992
PropertyValue(PropertyValue &&P);
9093

‎llvm/test/tools/sycl-post-link/device-globals/test_global_variable_name_mapping_metadata.ll

Lines changed: 75 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,75 @@
1+
; RUN: sycl-post-link --device-globals --emit-program-metadata -S %s -o %t.files.table
2+
; RUN: FileCheck %s -input-file=%t.files_0.prop --check-prefix CHECK-PROP
3+
4+
; This test is intended to check that the global_id_mapping program metadata properties are
5+
; generated for device_global variables in a program.
6+
7+
source_filename = "test_global_variable.cpp"
8+
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
9+
target triple = "spir64-unknown-unknown"
10+
11+
%"class.cl::sycl::ext::oneapi::device_global.0" = type { i32 addrspace(4)* }
12+
%"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 }
13+
%class.anon.0 = type { i8 }
14+
15+
@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #0
16+
@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #1
17+
@_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #2
18+
@_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #3
19+
@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #6
20+
21+
define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 {
22+
entry:
23+
%this.addr = alloca %class.anon.0 addrspace(4)*, align 8
24+
%this.addr.ascast = addrspacecast %class.anon.0 addrspace(4)** %this.addr to %class.anon.0 addrspace(4)* addrspace(4)*
25+
store %class.anon.0 addrspace(4)* %this, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
26+
%this1 = load %class.anon.0 addrspace(4)*, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
27+
%call1 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int1 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5
28+
%call2 = call spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(1)* @_ZL7dg_int2 to %"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)*)) #5
29+
%call3 = call spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool3 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5
30+
%call4 = call spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) addrspacecast (%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(1)* @_ZL8dg_bool4 to %"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)*)) #5
31+
ret void
32+
}
33+
34+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
35+
declare spir_func align 4 dereferenceable(4) i32 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIiJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.0" addrspace(4)* align 8 dereferenceable_or_null(8) %this) #4 align 2
36+
37+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
38+
declare spir_func align 1 dereferenceable(1) i8 addrspace(4)* @_ZNK2cl4sycl3ext6oneapi13device_globalIbJNS2_8PropertyIXadsoKcL_ZL5Name1EEEXadsoS5_L_ZL6Value1EEEEENS4_IXadsoS5_L_ZL5Name2EEEXadsoS5_L_ZL6Value2EEEEENS4_IXadsoS5_L_ZL5Name3EEEXadsoS5_L_ZL6Value3EEEEENS4_IXadsoS5_L_ZL5Name4EEEXadsoS5_L_ZL6Value4EEEEEEE3getEv(%"class.cl::sycl::ext::oneapi::device_global.1" addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2
39+
40+
attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" "sycl-device-global-size"="4" }
41+
attributes #1 = { "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" "sycl-implement-in-csr"="false" "sycl-init-mode"="1" "sycl-device-global-size"="4" }
42+
attributes #2 = { "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" "sycl-device-image-scope"="true" "sycl-host-access"="0" "sycl-implement-in-csr" "sycl-init-mode"="0" "sycl-device-global-size"="1" }
43+
attributes #3 = { "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" "sycl-device-image-scope" "sycl-host-access"="2" "sycl-device-global-size"="1" }
44+
attributes #4 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
45+
attributes #5 = { convergent nounwind }
46+
; no sycl-device-global-size attribute, this is not a device global variable but it contains compile-time properties,
47+
; a metadata node will be generated.
48+
attributes #6 = { "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" "sycl-device-image-scope"="false" "sycl-host-access"="1" "sycl-implement-in-csr"="true" "sycl-init-mode"="0" }
49+
50+
!llvm.dependent-libraries = !{!0}
51+
!llvm.module.flags = !{!1, !2}
52+
!opencl.spir.version = !{!3}
53+
!spirv.Source = !{!4}
54+
!llvm.ident = !{!5}
55+
56+
!0 = !{!"libcpmt"}
57+
!1 = !{i32 1, !"wchar_size", i32 2}
58+
!2 = !{i32 7, !"frame-pointer", i32 2}
59+
!3 = !{i32 1, i32 2}
60+
!4 = !{i32 4, i32 100000}
61+
!5 = !{!"clang version 14.0.0"}
62+
63+
; The encoding of global_id_mapping program metadata entries contain:
64+
; 1. 8 bytes denoting the bit-size of the byte array, here 64 bits or 8 bytes.
65+
; 2. N bytes with the name of the corresponding device_global variable, where
66+
; N is the value in the previous 8 bytes.
67+
;
68+
; CHECK-PROP: [SYCL/program metadata]
69+
; CHECK-PROP-NEXT: 6da74a122db9f35d____ZL7dg_int1@global_id_mapping=2|YBAAAAAAAAwXax0Nkd2Xp5GdxA
70+
; CHECK-PROP-NEXT: 7da74a1187b9f35d____ZL7dg_int2@global_id_mapping=2|YBAAAAAAAAwXax0Nkd2Xp5GdyA
71+
; CHECK-PROP-NEXT: 9d329ad59055e972____ZL8dg_bool3@global_id_mapping=2|gBAAAAAAAAwXaxEOkd2Xi92bsND
72+
; CHECK-PROP-NEXT: dda2bad52c45c432____ZL8dg_bool4@global_id_mapping=2|gBAAAAAAAAwXaxEOkd2Xi92bsRD
73+
;
74+
; The variable is not a device global one and must be ignored.
75+
; CHECK-PROP-NOT: 6da74a122db9f35d____ZL7no_dg_int1@global_id_mapping

‎llvm/tools/sycl-post-link/sycl-post-link.cpp

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -437,6 +437,17 @@ std::string saveModuleProperties(module_split::ModuleDesc &MD,
437437
MetadataNames.push_back(Func.getName().str() + "@reqd_work_group_size");
438438
ProgramMetadata.insert({MetadataNames.back(), KernelReqdWorkGroupSize});
439439
}
440+
441+
// Add global_id_mapping information with mapping between device-global
442+
// unique identifiers and the variable's name in the IR.
443+
for (auto &GV : M.globals()) {
444+
if (!isDeviceGlobalVariable(GV))
445+
continue;
446+
447+
StringRef GlobalID = getGlobalVariableUniqueId(GV);
448+
MetadataNames.push_back(GlobalID.str() + "@global_id_mapping");
449+
ProgramMetadata.insert({MetadataNames.back(), GV.getName()});
450+
}
440451
}
441452
if (MD.isESIMD()) {
442453
PropSet[PropSetRegTy::SYCL_MISC_PROP].insert({"isEsimdImage", true});

0 commit comments

Comments
 (0)