Skip to content

Commit 88123c1

Browse files
author
Pavel Samolysov
authored
[sycl-post-link] Implement property set generation for device globals (#5372)
The implementation is in accordance to the New "SYCL/device globals" property set chapter of the design document [1]. The --device-globals command-line argument to run the processing of device global variables has also been added. [1] https://github.com/intel/llvm/blob/sycl/sycl/doc/DeviceGlobal.md#new-sycldevice-globals-property-set
1 parent af6b96c commit 88123c1

File tree

11 files changed

+365
-11
lines changed

11 files changed

+365
-11
lines changed

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -192,6 +192,7 @@ class PropertySetRegistry {
192192
static constexpr char SYCL_MISC_PROP[] = "SYCL/misc properties";
193193
static constexpr char SYCL_ASSERT_USED[] = "SYCL/assert used";
194194
static constexpr char SYCL_EXPORTED_SYMBOLS[] = "SYCL/exported symbols";
195+
static constexpr char SYCL_DEVICE_GLOBALS[] = "SYCL/device globals";
195196

196197
// Function for bulk addition of an entire property set under given category
197198
// (property set name).

‎llvm/lib/Support/PropertySetIO.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -201,6 +201,7 @@ constexpr char PropertySetRegistry::SYCL_PROGRAM_METADATA[];
201201
constexpr char PropertySetRegistry::SYCL_MISC_PROP[];
202202
constexpr char PropertySetRegistry::SYCL_ASSERT_USED[];
203203
constexpr char PropertySetRegistry::SYCL_EXPORTED_SYMBOLS[];
204+
constexpr char PropertySetRegistry::SYCL_DEVICE_GLOBALS[];
204205

205206
} // namespace util
206207
} // namespace llvm

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

Lines changed: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,90 @@
1+
; RUN: sycl-post-link --device-globals -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 DeviceGlobalPass adds all the required
5+
; properties in the 'SYCL/device globals' property set and handles the
6+
; 'device_image_scope' attribute written in any allowed form.
7+
8+
source_filename = "test_global_variable.cpp"
9+
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"
10+
target triple = "spir64-unknown-unknown"
11+
12+
%"class.cl::sycl::ext::oneapi::device_global.0" = type { i32 addrspace(4)* }
13+
%"class.cl::sycl::ext::oneapi::device_global.1" = type { i8 }
14+
%class.anon.0 = type { i8 }
15+
16+
@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #0
17+
@_ZL7dg_int2 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #1
18+
@_ZL8dg_bool3 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #2
19+
@_ZL8dg_bool4 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.1" zeroinitializer, align 1 #3
20+
@_ZL7no_dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #6
21+
22+
define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #4 align 2 {
23+
entry:
24+
%this.addr = alloca %class.anon.0 addrspace(4)*, align 8
25+
%this.addr.ascast = addrspacecast %class.anon.0 addrspace(4)** %this.addr to %class.anon.0 addrspace(4)* addrspace(4)*
26+
store %class.anon.0 addrspace(4)* %this, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
27+
%this1 = load %class.anon.0 addrspace(4)*, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
28+
%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
29+
%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
30+
%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
31+
%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
32+
ret void
33+
}
34+
35+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
36+
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
37+
38+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
39+
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
40+
41+
attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "device_image_scope"="false" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" }
42+
attributes #1 = { "sycl-unique-id"="7da74a1187b9f35d____ZL7dg_int2" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="4" }
43+
attributes #2 = { "sycl-unique-id"="9d329ad59055e972____ZL8dg_bool3" "device_image_scope"="true" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="1" }
44+
attributes #3 = { "sycl-unique-id"="dda2bad52c45c432____ZL8dg_bool4" "device_image_scope" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" "sycl-device-global-size"="1" }
45+
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" }
46+
attributes #5 = { convergent nounwind }
47+
; no the sycl-device-global-size attribute, this is not a device global variable
48+
attributes #6 = { "sycl-unique-id"="6da74a122db9f35d____ZL7no_dg_int1" "device_image_scope"="false" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" }
49+
!llvm.dependent-libraries = !{!0}
50+
!llvm.module.flags = !{!1, !2}
51+
!opencl.spir.version = !{!3}
52+
!spirv.Source = !{!4}
53+
!llvm.ident = !{!5}
54+
55+
!0 = !{!"libcpmt"}
56+
!1 = !{i32 1, !"wchar_size", i32 2}
57+
!2 = !{i32 7, !"frame-pointer", i32 2}
58+
!3 = !{i32 1, i32 2}
59+
!4 = !{i32 4, i32 100000}
60+
!5 = !{!"clang version 14.0.0"}
61+
62+
; Ensure that the default values are correct.
63+
; ABAAAAAAAAABAAAAAxxxxx is decoded to
64+
; "40 00 00 00 00 00 00 00 | 04 00 00 00 | 00 | xx xx xx" which consists of:
65+
; 1. 8 bytes denoting the bit-size of the byte array, here 64 bits or 8 bytes.
66+
; 2. 4 bytes with the value of the 32-bit uint32_t integer with the size of the
67+
; underlying type of the device global variable. Its value being 4.
68+
; 3. 1 byte with the value of the 8-bit uint8_t integer with the flag that
69+
; the device global variable has the 'device_image_scope' property.
70+
; Its value being 0, no property.
71+
; 4. Any 3 bytes used as padding to align the structure to 8 bytes.
72+
;
73+
; ABAAAAAAAAQAAAAABxxxxx is decoded to
74+
; "40 00 00 00 00 00 00 00 | 01 00 00 00 | 01 | xx xx xx" which consists of:
75+
; 1. 8 bytes denoting the bit-size of the byte array, here 64 bits or 8 bytes.
76+
; 2. 4 bytes with the value of the 32-bit uint32_t integer with the size of the
77+
; underlying type of the device global variable. Its value being 1.
78+
; 3. 1 byte with the value of the 8-bit uint8_t integer with the flag that
79+
; the device global variable has the 'device_image_scope' property.
80+
; Its value being 1, property is present.
81+
; 4. Any 3 bytes used as padding to align the structure to 8 bytes.
82+
;
83+
; CHECK-PROP: [SYCL/device globals]
84+
; CHECK-PROP-NEXT: 6da74a122db9f35d____ZL7dg_int1=2|ABAAAAAAAAABAAAAA
85+
; CHECK-PROP-NEXT: 7da74a1187b9f35d____ZL7dg_int2=2|ABAAAAAAAAABAAAAA
86+
; CHECK-PROP-NEXT: 9d329ad59055e972____ZL8dg_bool3=2|ABAAAAAAAAQAAAAAB
87+
; CHECK-PROP-NEXT: dda2bad52c45c432____ZL8dg_bool4=2|ABAAAAAAAAQAAAAAB
88+
;
89+
; The variable is not a device global one and must be ignored
90+
; CHECK-PROP-NOT: 6da74a122db9f35d____ZL7no_dg_int1

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

Lines changed: 47 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,47 @@
1+
; RUN: sycl-post-link --device-globals -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 sycl-post-link doesn't add the header for
5+
; the empty SYCL/device globals property set.
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.anon.0 = type { i8 }
13+
14+
@_ZL7dg_int1 = internal addrspace(1) constant %"class.cl::sycl::ext::oneapi::device_global.0" zeroinitializer, align 8 #0
15+
16+
define internal spir_func void @_ZZ4mainENKUlvE_clEv(%class.anon.0 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #1 align 2 {
17+
entry:
18+
%this.addr = alloca %class.anon.0 addrspace(4)*, align 8
19+
%this.addr.ascast = addrspacecast %class.anon.0 addrspace(4)** %this.addr to %class.anon.0 addrspace(4)* addrspace(4)*
20+
store %class.anon.0 addrspace(4)* %this, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
21+
%this1 = load %class.anon.0 addrspace(4)*, %class.anon.0 addrspace(4)* addrspace(4)* %this.addr.ascast, align 8
22+
%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)*)) #2
23+
ret void
24+
}
25+
26+
; Function Attrs: convergent mustprogress noinline norecurse nounwind optnone
27+
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) #1 align 2
28+
29+
; no the sycl-device-global-size attribute, this is not a device global variable
30+
attributes #0 = { "sycl-unique-id"="6da74a122db9f35d____ZL7dg_int1" "device_image_scope"="false" "host_access"="1" "implement_in_csr"="true" "init_mode"="0" }
31+
attributes #1 = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
32+
attributes #2 = { convergent nounwind }
33+
34+
!llvm.dependent-libraries = !{!0}
35+
!llvm.module.flags = !{!1, !2}
36+
!opencl.spir.version = !{!3}
37+
!spirv.Source = !{!4}
38+
!llvm.ident = !{!5}
39+
40+
!0 = !{!"libcpmt"}
41+
!1 = !{i32 1, !"wchar_size", i32 2}
42+
!2 = !{i32 7, !"frame-pointer", i32 2}
43+
!3 = !{i32 1, i32 2}
44+
!4 = !{i32 4, i32 100000}
45+
!5 = !{!"clang version 14.0.0"}
46+
47+
; CHECK-PROP-NOT: [SYCL/device globals]

‎llvm/test/tools/sycl-post-link/help.test

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -42,6 +42,7 @@ CHECK: --version - Display the version of this program
4242

4343
CHECK: sycl-post-link options:
4444

45+
CHECK: --device-globals - Lower and generate information about device global variables
4546
CHECK: -f - Enable binary output on terminals
4647
CHECK: --ir-output-only - Output single IR file
4748
CHECK: -o=<filename> - Output filename

‎llvm/tools/sycl-post-link/CMakeLists.txt

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,9 +19,10 @@ include_directories(
1919

2020
add_llvm_tool(sycl-post-link
2121
sycl-post-link.cpp
22-
SYCLKernelParamOptInfo.cpp
22+
DeviceGlobals.cpp
2323
SpecConstants.cpp
2424
SYCLDeviceLibReqMask.cpp
25+
SYCLKernelParamOptInfo.cpp
2526
ADDITIONAL_HEADER_DIRS
2627
${LLVMGenXIntrinsics_SOURCE_DIR}/GenXIntrinsics/include
2728
${LLVMGenXIntrinsics_BINARY_DIR}/GenXIntrinsics/include

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

Lines changed: 118 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,118 @@
1+
//===----- DeviceGlobals.cpp - SYCL Device Globals Pass -------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
// See comments in the header.
9+
//===----------------------------------------------------------------------===//
10+
11+
#include "DeviceGlobals.h"
12+
13+
#include "llvm/ADT/STLExtras.h"
14+
#include "llvm/ADT/StringRef.h"
15+
#include "llvm/IR/Module.h"
16+
17+
#include <cassert>
18+
19+
using namespace llvm;
20+
21+
namespace {
22+
23+
constexpr StringRef SYCL_DEVICE_GLOBAL_SIZE_ATTR = "sycl-device-global-size";
24+
constexpr StringRef SYCL_UNIQUE_ID_ATTR = "sycl-unique-id";
25+
constexpr StringRef SYCL_DEVICE_IMAGE_SCOPE_ATTR = "device_image_scope";
26+
27+
/// Converts the string into a boolean value. If the string is equal to "false"
28+
/// we consider its value as /c false, /true otherwise.
29+
///
30+
/// @param Value [in] "boolean as string" value.
31+
///
32+
/// @returns \c false if the value of \c Value equals to "false", \c true
33+
/// otherwise.
34+
bool toBool(StringRef Value) { return !Value.equals("false"); }
35+
36+
/// Checks whether the device global variable has the \c device_image_scope
37+
/// property. The variable has the property if the \c sycl-device-image-scope
38+
/// attribute is defined for the variable and the attribute value is not
39+
/// represented as \c false.
40+
///
41+
/// @param GV [in] Device Global variable.
42+
///
43+
/// @returns \c true if variable \c GV has the \c device_image_scope property,
44+
/// \c false otherwise.
45+
bool hasDeviceImageScope(const GlobalVariable &GV) {
46+
return GV.hasAttribute(SYCL_DEVICE_IMAGE_SCOPE_ATTR) &&
47+
toBool(
48+
GV.getAttribute(SYCL_DEVICE_IMAGE_SCOPE_ATTR).getValueAsString());
49+
}
50+
51+
/// Returns the size (in bytes) of the underlying type \c T of the device
52+
/// global variable.
53+
///
54+
/// The function gets this value from the LLVM IR attribute \c
55+
/// sycl-device-global-size.
56+
///
57+
/// @param GV [in] Device Global variable.
58+
///
59+
/// @returns the size (int bytes) of the underlying type \c T of the
60+
/// device global variable represented in the LLVM IR by \c GV.
61+
uint32_t getUnderlyingTypeSize(const GlobalVariable &GV) {
62+
assert(GV.hasAttribute(SYCL_DEVICE_GLOBAL_SIZE_ATTR) &&
63+
"The device global variable must have the 'sycl-device-global-size' "
64+
"attribute");
65+
uint32_t value;
66+
bool error = GV.getAttribute(SYCL_DEVICE_GLOBAL_SIZE_ATTR)
67+
.getValueAsString()
68+
.getAsInteger(10, value);
69+
assert(!error &&
70+
"The 'sycl-device-global-size' attribute must contain a number"
71+
" representing the size of the underlying type T of the device"
72+
" global variable");
73+
(void)error;
74+
return value;
75+
}
76+
77+
/// Returns the unique id for the device global variable.
78+
///
79+
/// The function gets this value from the LLVM IR attribute \c
80+
/// sycl-unique-id. If the attribute is not found for the variable
81+
/// an error should occur even in the release build.
82+
///
83+
/// @param GV [in] Device Global variable.
84+
///
85+
/// @returns the unique id of the device global variable represented
86+
/// in the LLVM IR by \c GV.
87+
StringRef getUniqueId(const GlobalVariable &GV) {
88+
assert(GV.hasAttribute(SYCL_UNIQUE_ID_ATTR) &&
89+
"a 'sycl-unique-id' string must be associated with every device "
90+
"global variable");
91+
return GV.getAttribute(SYCL_UNIQUE_ID_ATTR).getValueAsString();
92+
}
93+
94+
} // namespace
95+
96+
DeviceGlobalPropertyMapTy
97+
DeviceGlobalsPass::collectDeviceGlobalProperties(const Module &M) {
98+
DeviceGlobalPropertyMapTy DGM;
99+
auto DevGlobalFilter = [](auto &GV) {
100+
return GV.hasAttribute(SYCL_DEVICE_GLOBAL_SIZE_ATTR);
101+
};
102+
103+
auto DevGlobalNum = count_if(M.globals(), DevGlobalFilter);
104+
if (DevGlobalNum == 0)
105+
return DGM;
106+
107+
DGM.reserve(DevGlobalNum);
108+
109+
for (auto &GV : M.globals()) {
110+
if (!DevGlobalFilter(GV))
111+
continue;
112+
113+
DGM[getUniqueId(GV)] = {
114+
{{getUnderlyingTypeSize(GV), hasDeviceImageScope(GV)}}};
115+
}
116+
117+
return DGM;
118+
}

‎llvm/tools/sycl-post-link/DeviceGlobals.h

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
//===----- DeviceGlobals.h - SYCL Device Globals Pass ---------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// A transformation pass which converts symbolic device globals attributes
10+
// to integer id-based ones to later map to SPIRV device globals. The class
11+
// also contains a number of static methods to extract corresponding
12+
// attributes of the device global variables and save them as a property set
13+
// for the runtime.
14+
//===----------------------------------------------------------------------===//
15+
16+
#pragma once
17+
18+
#include "llvm/ADT/MapVector.h"
19+
20+
#include <cstdint>
21+
#include <vector>
22+
23+
namespace llvm {
24+
25+
class Module;
26+
class StringRef;
27+
28+
// Represents a device global variable - at SYCL RT level device global
29+
// variables are being represented as a byte-array.
30+
struct DeviceGlobalProperty {
31+
DeviceGlobalProperty(uint32_t Size, uint8_t DeviceImageScope)
32+
: Size(Size), DeviceImageScope(DeviceImageScope) {}
33+
34+
// Encodes size of the underlying type T of the device global variable.
35+
uint32_t Size;
36+
37+
// Either 1 (true) or 0 (false), telling whether the device global variable
38+
// was declared with the device_image_scope property.
39+
uint8_t DeviceImageScope;
40+
};
41+
42+
using DeviceGlobalPropertyMapTy =
43+
MapVector<StringRef, std::vector<DeviceGlobalProperty>>;
44+
45+
class DeviceGlobalsPass {
46+
public:
47+
// Searches given module for occurrences of device global variable-specific
48+
// metadata and builds "device global variable name" ->
49+
// vector<"variable properties"> map.
50+
static DeviceGlobalPropertyMapTy
51+
collectDeviceGlobalProperties(const Module &M);
52+
};
53+
54+
} // end namespace llvm

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

Lines changed: 1 addition & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -9,6 +9,7 @@
99
//===----------------------------------------------------------------------===//
1010

1111
#include "SpecConstants.h"
12+
#include "Support.h"
1213

1314
#include "llvm/ADT/APInt.h"
1415
#include "llvm/ADT/StringMap.h"
@@ -17,7 +18,6 @@
1718
#include "llvm/IR/Instruction.h"
1819
#include "llvm/IR/Instructions.h"
1920
#include "llvm/IR/Operator.h"
20-
#include "llvm/Support/ErrorHandling.h"
2121

2222
using namespace llvm;
2323

@@ -51,11 +51,6 @@ constexpr char SPEC_CONST_MD_STRING[] = "sycl.specialization-constants";
5151
constexpr char SPEC_CONST_DEFAULT_VAL_MD_STRING[] =
5252
"sycl.specialization-constants-default-values";
5353

54-
void AssertRelease(bool Cond, const char *Msg) {
55-
if (!Cond)
56-
report_fatal_error((Twine("SpecConstants.cpp: ") + Msg).str().c_str());
57-
}
58-
5954
StringRef getStringLiteralArg(const CallInst *CI, unsigned ArgNo,
6055
SmallVectorImpl<Instruction *> &DelInsts) {
6156
Value *V = CI->getArgOperand(ArgNo)->stripPointerCasts();

0 commit comments

Comments
 (0)