Skip to content

Commit f82ddf4

Browse files
author
Artem Gindinson
authored
[SYCL] Always generate specialization constants buffer (#4591)
The discrepancy of kernel signatures between SPIR-V images (for which SYCL2020 spec constants are supported natively) and AOT-compiled images (which actually require the buffer to emulate SYCL2020 spec constants) has blocked the multi-target compilation & execution when the feature is used. Align the kernel signatures by generating a (nullified) buffer for the SPIR-V target as well and make sure it is handled correctly by the DPC++ RT. Signed-off-by: Artem Gindinson <[email protected]>
1 parent a346c08 commit f82ddf4

File tree

6 files changed

+38
-38
lines changed

6 files changed

+38
-38
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 2 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -2220,14 +2220,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
22202220
return true;
22212221
}
22222222

2223-
// Generate kernel argument to intialize specialization constants. This
2224-
// argument is only generated when the target has no native support for
2225-
// specialization constants
2223+
// Generate kernel argument to initialize specialization constants.
22262224
void handleSyclKernelHandlerType() {
22272225
ASTContext &Context = SemaRef.getASTContext();
2228-
if (isDefaultSPIRArch(Context))
2229-
return;
2230-
22312226
StringRef Name = "_arg__specialization_constants_buffer";
22322227
addParam(Name, Context.getPointerType(Context.getAddrSpaceQualType(
22332228
Context.CharTy, LangAS::sycl_global)));
@@ -2470,9 +2465,6 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
24702465
}
24712466

24722467
void handleSyclKernelHandlerType() {
2473-
ASTContext &Context = SemaRef.getASTContext();
2474-
if (isDefaultSPIRArch(Context))
2475-
return;
24762468
addParam(DC.getParamVarDeclsForCurrentField()[0]->getType(),
24772469
"SYCL2020 specialization constant");
24782470
}
@@ -3321,13 +3313,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
33213313
void handleSyclKernelHandlerType(QualType Ty) {
33223314
// The compiler generated kernel argument used to initialize SYCL 2020
33233315
// specialization constants, `specialization_constants_buffer`, should
3324-
// have corresponding entry in integration header. This argument is
3325-
// only generated when target has no native support for specialization
3326-
// constants.
3316+
// have corresponding entry in integration header.
33273317
ASTContext &Context = SemaRef.getASTContext();
3328-
if (isDefaultSPIRArch(Context))
3329-
return;
3330-
33313318
// Offset is zero since kernel_handler argument is not part of
33323319
// kernel object (i.e. it is not captured)
33333320
addParam(Context.getPointerType(Context.CharTy),

‎clang/test/CodeGenSYCL/int_header_sycl2020_spec_const.cpp

Lines changed: 16 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,18 @@
1-
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
2-
// RUN: FileCheck -input-file=%t.h %s --check-prefix=NONATIVESUPPORT --check-prefix=ALL
1+
// Generic SPIR-V target
32
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
4-
// RUN: FileCheck -input-file=%t.h %s --check-prefix=NATIVESUPPORT --check-prefix=ALL
3+
// RUN: FileCheck -input-file=%t.h %s
4+
//
5+
// SPIR-V AOT targets
6+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64_gen-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
7+
// RUN: FileCheck -input-file=%t.h %s
8+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64_x86_64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
9+
// RUN: FileCheck -input-file=%t.h %s
10+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64_fpga-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
11+
// RUN: FileCheck -input-file=%t.h %s
12+
//
13+
// Non-SPIR target
14+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -fsycl-int-header=%t.h %s -o %t.out %s -o %t.out
15+
// RUN: FileCheck -input-file=%t.h %s
516

617
// This test checks that the compiler generates required information
718
// in integration header for kernel_handler type (SYCL 2020 specialization
@@ -24,6 +35,5 @@ int main() {
2435
kh);
2536
});
2637
}
27-
// ALL: const kernel_param_desc_t kernel_signatures[] = {
28-
// NONATIVESUPPORT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 }
29-
// NATIVESUPPORT-NOT: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 }
38+
// CHECK: const kernel_param_desc_t kernel_signatures[] = {
39+
// CHECK: { kernel_param_kind_t::kind_specialization_constants_buffer, 8, 0 }

‎clang/test/CodeGenSYCL/kernel-handler.cpp

Lines changed: 13 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NONATIVESUPPORT
2-
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefix=NATIVESUPPORT
1+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=ALL,NONATIVESUPPORT
2+
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=ALL,NATIVESUPPORT
33

44
// This test checks IR generated when kernel_handler argument
55
// (used to handle SYCL 2020 specialization constants) is passed
@@ -22,18 +22,17 @@ void test(int val) {
2222
});
2323
}
2424

25-
// NONATIVESUPPORT: define dso_local void @{{.*}}test_kernel_handler{{[^(]*}}
26-
// NONATIVESUPPORT-SAME: (i32 %_arg_, i8 addrspace(1)* %_arg__specialization_constants_buffer)
27-
// NONATIVESUPPORT: %kh = alloca %"class.cl::sycl::kernel_handler", align 1
25+
// ALL: define dso_local{{ spir_kernel | }}void @{{.*}}test_kernel_handler{{[^(]*}}
26+
// ALL-SAME: (i32 %_arg_, i8 addrspace(1)* %_arg__specialization_constants_buffer)
27+
// ALL: %kh = alloca %"class.cl::sycl::kernel_handler", align 1
28+
2829
// NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)** %_arg__specialization_constants_buffer.addr, align 8
2930
// NONATIVESUPPORT: %[[ADDRSPACECAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[KH]] to i8*
3031
// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.cl::sycl::kernel_handler"* nonnull align 1 dereferenceable(1) %kh, i8* %[[ADDRSPACECAST]])
31-
// NONATIVESUPPORT: void @[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]
32-
// NONATIVESUPPORT-SAME: byval(%"class.cl::sycl::kernel_handler")
33-
34-
// NATIVESUPPORT: define dso_local spir_kernel void @{{.*}}test_kernel_handler{{[^(]*}}
35-
// NATIVESUPPORT-SAME: (i32 %_arg_)
36-
// NATIVESUPPORT: %kh = alloca %"class.cl::sycl::kernel_handler"
37-
// NATIVESUPPORT-NOT: __init_specialization_constants_buffer
38-
// NATIVE-SUPPORT: call spir_func void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]"
39-
// NATIVE-SUPPORT-SAME: byval(%"class.cl::sycl::kernel_handler")
32+
33+
// NATIVESUPPORT-NOT: load i8 addrspace(1)*, i8 addrspace(1)** %_arg__specialization_constants_buffer.addr, align 8
34+
// NATIVESUPPORT-NOT: addrspacecast i8 addrspace(1)* %{{[0-9]+}} to i8*
35+
// NATIVESUPPORT-NOT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.cl::sycl::kernel_handler"* nonnull align 1 dereferenceable(1) %kh, i8* %{{[0-9]+}})
36+
37+
// ALL: call{{ spir_func | }}void @{{[a-zA-Z0-9_$]+}}kernel_handler{{[a-zA-Z0-9_$]+}}
38+
// ALL-SAME: byval(%"class.cl::sycl::kernel_handler")

‎clang/test/SemaSYCL/kernel-handler.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -107,8 +107,9 @@ int main() {
107107
// Test AST for default SPIR architecture
108108

109109
// Check test_kernel_handler parameters
110-
// NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int)'
110+
// NATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, __global char *)'
111111
// NATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int'
112+
// NATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__specialization_constants_buffer '__global char *'
112113

113114
// Check declaration and initialization of kernel object local clone
114115
// NATIVESUPPORT-NEXT: CompoundStmt

‎sycl/source/detail/device_image_impl.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -183,7 +183,7 @@ class device_image_impl {
183183

184184
RT::PiMem &get_spec_const_buffer_ref() noexcept {
185185
std::lock_guard<std::mutex> Lock{MSpecConstAccessMtx};
186-
if (nullptr == MSpecConstsBuffer) {
186+
if (nullptr == MSpecConstsBuffer && !MSpecConstsBlob.empty()) {
187187
const detail::plugin &Plugin = getSyclObjImpl(MContext)->getPlugin();
188188
Plugin.call<PiApiKind::piMemBufferCreate>(
189189
detail::getSyclObjImpl(MContext)->getHandleRef(),

‎sycl/source/detail/scheduler/commands.cpp

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1797,8 +1797,11 @@ static pi_result SetKernelParamsAndLaunch(
17971797
}
17981798
assert(DeviceImageImpl != nullptr);
17991799
RT::PiMem SpecConstsBuffer = DeviceImageImpl->get_spec_const_buffer_ref();
1800+
// Avoid taking an address of nullptr
1801+
RT::PiMem *SpecConstsBufferArg =
1802+
SpecConstsBuffer ? &SpecConstsBuffer : nullptr;
18001803
Plugin.call<PiApiKind::piextKernelSetArgMemObj>(Kernel, NextTrueIndex,
1801-
&SpecConstsBuffer);
1804+
SpecConstsBufferArg);
18021805
break;
18031806
}
18041807
case kernel_param_kind_t::kind_invalid:

0 commit comments

Comments
 (0)