Skip to content

Commit 105a151

Browse files
vmaksimovmaksimo
committed
Merge remote-tracking branch 'intel_llvm/sycl' into llvmspirv_pulldown
2 parents 0e952c4 + ae89341 commit 105a151

38 files changed

+1429
-289
lines changed

‎.github/workflows/gh_pages.yml

Lines changed: 3 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -41,12 +41,9 @@ jobs:
4141
git rm -rf .
4242
touch .nojekyll
4343
yes | \cp -rf $GITHUB_WORKSPACE/build/tools/sycl/doc/html/* .
44-
mkdir doxygen
45-
yes | \cp -rf $GITHUB_WORKSPACE/build/tools/sycl/doc/doxygen/html/* doxygen/
46-
mkdir clang
47-
yes | \cp -rf $GITHUB_WORKSPACE/build/tools/clang/docs/html/* clang/
48-
mkdir clang_doxygen
49-
yes | \cp -rf $GITHUB_WORKSPACE/build/tools/clang/docs/doxygen/html/* clang_doxygen/
44+
mv $GITHUB_WORKSPACE/build/tools/sycl/doc/doxygen/html doxygen/
45+
mv $GITHUB_WORKSPACE/build/tools/clang/docs/html clang/
46+
mv $GITHUB_WORKSPACE/build/tools/clang/docs/doxygen/html clang_doxygen/
5047
git config --global user.name "iclsrc"
5148
git config --global user.email "[email protected]"
5249
git add .

‎clang/docs/LanguageExtensions.rst

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2453,7 +2453,7 @@ their usual pattern without any special treatment.
24532453
constexpr const char * __builtin_sycl_unique_stable_name( type-id );
24542454
24552455
``__builtin_sycl_unique_stable_id``
2456-
----------------------------------
2456+
-----------------------------------
24572457
24582458
Like ``__builtin_sycl_unique_stable_name``, this builtin generates a unique and
24592459
stable name as a string literal to support sharing it across split compliations.

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -4331,6 +4331,15 @@ void Clang::ConstructHostCompilerJob(Compilation &C, const JobAction &JA,
43314331
}
43324332
} else
43334333
HostCompileArgs.push_back("-E");
4334+
4335+
// Add the integration header.
4336+
StringRef Header =
4337+
TC.getDriver().getIntegrationHeader(InputFile.getBaseInput());
4338+
if (types::getPreprocessedType(InputFile.getType()) != types::TY_INVALID &&
4339+
!Header.empty()) {
4340+
HostCompileArgs.push_back(IsMSVCHostCompiler ? "-FI" : "-include");
4341+
HostCompileArgs.push_back(TCArgs.MakeArgString(Header));
4342+
}
43344343
} else if (isa<AssembleJobAction>(JA)) {
43354344
HostCompileArgs.push_back("-c");
43364345
if (IsMSVCHostCompiler)
@@ -4375,15 +4384,6 @@ void Clang::ConstructHostCompilerJob(Compilation &C, const JobAction &JA,
43754384
HostCompileArgs.push_back(Output.getFilename());
43764385
}
43774386

4378-
// Add the integration header.
4379-
StringRef Header =
4380-
TC.getDriver().getIntegrationHeader(InputFile.getBaseInput());
4381-
if (types::getPreprocessedType(InputFile.getType()) != types::TY_INVALID &&
4382-
!Header.empty()) {
4383-
HostCompileArgs.push_back(IsMSVCHostCompiler ? "-FI" : "-include");
4384-
HostCompileArgs.push_back(TCArgs.MakeArgString(Header));
4385-
}
4386-
43874387
SmallString<128> ExecPath;
43884388
if (HostCompilerDefArg) {
43894389
ExecPath = HostCompilerDefArg->getValue();

‎clang/test/Driver/sycl-host-compiler.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,15 +5,15 @@
55
// RUN: %clangxx -fsycl-use-footer -fsycl -fsycl-host-compiler=/some/dir/g++ %s -### 2>&1 \
66
// RUN: | FileCheck -check-prefix=HOST_COMPILER %s
77
// HOST_COMPILER: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]" "-fsycl-int-footer={{.*}}"
8-
// HOST_COMPILER: g++{{.*}} "-E" "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-o" "[[TMPII:.+\.ii]]"{{.*}} "-include" "[[INTHEADER]]"
9-
// HOST_COMPILER: g++{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-o" "[[HOSTOBJ:.+\.o]]"{{.*}}
8+
// HOST_COMPILER: g++{{.*}} "-E" "-include" "[[INTHEADER]]" "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-o" "[[TMPII:.+\.ii]]"
9+
// HOST_COMPILER: g++{{.*}} "-c" "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-o" "[[HOSTOBJ:.+\.o]]"{{.*}}
1010
// HOST_COMPILER: ld{{.*}} "[[HOSTOBJ]]"
1111

1212
// RUN: %clang_cl -fsycl-use-footer -fsycl -fsycl-host-compiler=/some/dir/cl %s -### 2>&1 \
1313
// RUN: | FileCheck -check-prefix=HOST_COMPILER_CL %s
1414
// HOST_COMPILER_CL: clang{{.*}} "-fsycl-is-device"{{.*}} "-fsycl-int-header=[[INTHEADER:.+\.h]]" "-fsycl-int-footer={{.*}}"
15-
// HOST_COMPILER_CL: cl{{.*}} "-P" "-Fi[[TMPII:.+\.ii]]"{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-FI" "[[INTHEADER]]"
16-
// HOST_COMPILER_CL: cl{{.*}} "-Fo[[HOSTOBJ:.+\.obj]]"{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}}
15+
// HOST_COMPILER_CL: cl{{.*}} "-P" "-Fi[[TMPII:.+\.ii]]" "-FI" "[[INTHEADER]]"{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"
16+
// HOST_COMPILER_CL: cl{{.*}} "-c" "-Fo[[HOSTOBJ:.+\.obj]]"{{.*}} "-I" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}}
1717
// HOST_COMPILER_CL: link{{.*}} "[[HOSTOBJ]]"
1818

1919
/// check for additional host options

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -189,6 +189,7 @@ class PropertySetRegistry {
189189
static constexpr char SYCL_DEVICELIB_REQ_MASK[] = "SYCL/devicelib req mask";
190190
static constexpr char SYCL_KERNEL_PARAM_OPT_INFO[] = "SYCL/kernel param opt";
191191
static constexpr char SYCL_MISC_PROP[] = "SYCL/misc properties";
192+
static constexpr char SYCL_ASSERT_USED[] = "SYCL/assert used";
192193

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

‎llvm/lib/Support/PropertySetIO.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -198,6 +198,7 @@ constexpr char PropertySetRegistry::SYCL_DEVICELIB_REQ_MASK[];
198198
constexpr char PropertySetRegistry::SYCL_SPEC_CONSTANTS_DEFAULT_VALUES[];
199199
constexpr char PropertySetRegistry::SYCL_KERNEL_PARAM_OPT_INFO[];
200200
constexpr char PropertySetRegistry::SYCL_MISC_PROP[];
201+
constexpr char PropertySetRegistry::SYCL_ASSERT_USED[];
201202

202203
} // namespace util
203204
} // namespace llvm

‎llvm/test/tools/sycl-post-link/assert-property-1.ll

Lines changed: 159 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,159 @@
1+
; This test checks that the post-link tool properly generates "assert used"
2+
; property - it should include only kernels that call assertions in their call
3+
; graph.
4+
5+
; RUN: sycl-post-link -split=auto -symbols -S %s -o %t.table
6+
; RUN: FileCheck %s -input-file=%t_0.prop
7+
8+
; SYCL source:
9+
; void foo() {
10+
; assert(0);
11+
; }
12+
; void bar() {
13+
; assert(1);
14+
; }
15+
; void baz() {
16+
; foo();
17+
; }
18+
;
19+
; int main() {
20+
; queue Q;
21+
; Q.submit([&] (handler& CGH) {
22+
; CGH.parallel_for<class TheKernel>(range<2>{2, 10}, [=](item<2> It) {
23+
; foo();
24+
; });
25+
; CGH.parallel_for<class TheKernel2>(range<2>{2, 10}, [=](item<2> It) {
26+
; bar();
27+
; });
28+
; CGH.parallel_for<class TheKernel3>(range<2>{2, 10}, [=](item<2> It) {
29+
; baz();
30+
; bar();
31+
; });
32+
; });
33+
; Q.wait();
34+
; return 0;
35+
; }
36+
37+
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"
38+
target triple = "spir64_x86_64-unknown-unknown-sycldevice"
39+
40+
%"class._ZTSN2cl4sycl5rangeILi2EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" }
41+
%"class._ZTSN2cl4sycl6detail5arrayILi2EEE.cl::sycl::detail::array" = type { [2 x i64] }
42+
%"class._ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEEUlNS1_4itemILi2ELb1EEEE_.anon" = type { i8 }
43+
44+
@.str = private unnamed_addr addrspace(1) constant [2 x i8] c"0\00", align 1
45+
@.str.1 = private unnamed_addr addrspace(1) constant [11 x i8] c"assert.cpp\00", align 1
46+
@__PRETTY_FUNCTION__._Z3foov = private unnamed_addr addrspace(1) constant [11 x i8] c"void foo()\00", align 1
47+
@__spirv_BuiltInGlobalInvocationId = external dso_local addrspace(1) constant <3 x i64>, align 32
48+
@__spirv_BuiltInLocalInvocationId = external dso_local addrspace(1) constant <3 x i64>, align 32
49+
@_ZL10assert_fmt = internal addrspace(2) constant [85 x i8] c"%s:%d: %s: global id: [%lu,%lu,%lu], local id: [%lu,%lu,%lu] Assertion `%s` failed.\0A\00", align 1
50+
51+
; CHECK: [SYCL/assert used]
52+
53+
; Function Attrs: convergent norecurse nounwind mustprogress
54+
define dso_local spir_func void @_Z3foov() {
55+
entry:
56+
tail call spir_func void @__assert_fail(i8 addrspace(4)* getelementptr inbounds ([2 x i8], [2 x i8] addrspace(4)* addrspacecast ([2 x i8] addrspace(1)* @.str to [2 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* getelementptr inbounds ([11 x i8], [11 x i8] addrspace(4)* addrspacecast ([11 x i8] addrspace(1)* @.str.1 to [11 x i8] addrspace(4)*), i64 0, i64 0), i32 8, i8 addrspace(4)* getelementptr inbounds ([11 x i8], [11 x i8] addrspace(4)* addrspacecast ([11 x i8] addrspace(1)* @__PRETTY_FUNCTION__._Z3foov to [11 x i8] addrspace(4)*), i64 0, i64 0))
57+
ret void
58+
}
59+
60+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel
61+
; Function Attrs: convergent norecurse
62+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE9TheKernel"() #0 {
63+
entry:
64+
call spir_func void @_Z3foov()
65+
ret void
66+
}
67+
68+
; Function Attrs: nofree norecurse nosync nounwind readnone willreturn mustprogress
69+
define dso_local spir_func void @_Z3barv() {
70+
entry:
71+
ret void
72+
}
73+
74+
; CHECK-NOT: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2
75+
; Function Attrs: norecurse
76+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel2"() #1 {
77+
entry:
78+
call spir_func void @_Z3barv()
79+
ret void
80+
}
81+
82+
; Function Attrs: convergent inlinehint norecurse nounwind mustprogress
83+
define internal spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_4itemILi2ELb1EEEE1_clES5_"() unnamed_addr #8 align 2 {
84+
entry:
85+
call spir_func void @_Z3bazv()
86+
call spir_func void @_Z3barv()
87+
ret void
88+
}
89+
90+
; Function Attrs: convergent norecurse nounwind mustprogress
91+
define dso_local spir_func void @_Z3bazv() {
92+
entry:
93+
tail call spir_func void @__assert_fail(i8 addrspace(4)* getelementptr inbounds ([2 x i8], [2 x i8] addrspace(4)* addrspacecast ([2 x i8] addrspace(1)* @.str to [2 x i8] addrspace(4)*), i64 0, i64 0), i8 addrspace(4)* getelementptr inbounds ([11 x i8], [11 x i8] addrspace(4)* addrspacecast ([11 x i8] addrspace(1)* @.str.1 to [11 x i8] addrspace(4)*), i64 0, i64 0), i32 8, i8 addrspace(4)* getelementptr inbounds ([11 x i8], [11 x i8] addrspace(4)* addrspacecast ([11 x i8] addrspace(1)* @__PRETTY_FUNCTION__._Z3foov to [11 x i8] addrspace(4)*), i64 0, i64 0))
94+
ret void
95+
}
96+
97+
; CHECK: _ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3
98+
; Function Attrs: convergent norecurse
99+
define weak_odr dso_local spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE10TheKernel3"() #0 {
100+
entry:
101+
call spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_4itemILi2ELb1EEEE1_clES5_"()
102+
ret void
103+
}
104+
105+
; Function Attrs: convergent norecurse mustprogress
106+
define weak dso_local spir_func void @__assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func) {
107+
entry:
108+
%call = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_xv()
109+
%call1 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_yv()
110+
%call2 = tail call spir_func i64 @_Z28__spirv_GlobalInvocationId_zv()
111+
%call3 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_xv()
112+
%call4 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_yv()
113+
%call5 = tail call spir_func i64 @_Z27__spirv_LocalInvocationId_zv()
114+
tail call spir_func void @__devicelib_assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func, i64 %call, i64 %call1, i64 %call2, i64 %call3, i64 %call4, i64 %call5)
115+
ret void
116+
}
117+
118+
; Function Attrs: inlinehint norecurse mustprogress
119+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_xv() local_unnamed_addr
120+
121+
; Function Attrs: inlinehint norecurse mustprogress
122+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_yv() local_unnamed_addr
123+
124+
; Function Attrs: inlinehint norecurse mustprogress
125+
declare dso_local spir_func i64 @_Z28__spirv_GlobalInvocationId_zv() local_unnamed_addr
126+
127+
; Function Attrs: inlinehint norecurse mustprogress
128+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_xv() local_unnamed_addr
129+
130+
; Function Attrs: inlinehint norecurse mustprogress
131+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_yv() local_unnamed_addr
132+
133+
; Function Attrs: inlinehint norecurse mustprogress
134+
declare dso_local spir_func i64 @_Z27__spirv_LocalInvocationId_zv() local_unnamed_addr
135+
136+
; Function Attrs: convergent norecurse mustprogress
137+
define weak dso_local spir_func void @__devicelib_assert_fail(i8 addrspace(4)* %expr, i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func, i64 %gid0, i64 %gid1, i64 %gid2, i64 %lid0, i64 %lid1, i64 %lid2) {
138+
entry:
139+
%call = tail call spir_func i32 (i8 addrspace(2)*, ...) @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)* getelementptr inbounds ([85 x i8], [85 x i8] addrspace(2)* @_ZL10assert_fmt, i64 0, i64 0), i8 addrspace(4)* %file, i32 %line, i8 addrspace(4)* %func, i64 %gid0, i64 %gid1, i64 %gid2, i64 %lid0, i64 %lid1, i64 %lid2, i8 addrspace(4)* %expr)
140+
ret void
141+
}
142+
143+
; Function Attrs: convergent
144+
declare dso_local spir_func i32 @_Z18__spirv_ocl_printfPU3AS2Kcz(i8 addrspace(2)*, ...)
145+
146+
attributes #0 = { convergent norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="assert.cpp" "uniform-work-group-size"="true" }
147+
attributes #1 = { norecurse "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="assert.cpp" "uniform-work-group-size"="true" }
148+
149+
!opencl.spir.version = !{!0, !0, !0, !0, !0, !0, !0, !0, !0, !0, !0}
150+
!spirv.Source = !{!1, !1, !1, !1, !1, !1, !1, !1, !1, !1, !1}
151+
!llvm.ident = !{!2, !2, !2, !2, !2, !2, !2, !2, !2, !2, !2}
152+
!llvm.module.flags = !{!3, !4}
153+
154+
!0 = !{i32 1, i32 2}
155+
!1 = !{i32 4, i32 100000}
156+
!2 = !{!"clang version 13.0.0 (https://github.com/intel/llvm)"}
157+
!3 = !{i32 1, !"wchar_size", i32 4}
158+
!4 = !{i32 7, !"frame-pointer", i32 2}
159+
!5 = !{i32 -1, i32 -1}

0 commit comments

Comments
 (0)