From 83a91d19c4e56e74318cecf28e07febf638c7892 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Fri, 9 Apr 2021 17:28:32 +0300 Subject: [PATCH 1/2] [SYCL] Fix address space for spec constants buffer This patch sets OpenCL global address space for specialization constants buffer instead of generic one. It fixes the crash which happens during passing nullptr to piKernelSetArg PI API func Co-authored-by: Vlad Romanov --- clang/lib/Sema/SemaSYCL.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a76d5c3bc7adb..8eddf1fadea0e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2018,7 +2018,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { return; StringRef Name = "_arg__specialization_constants_buffer"; - addParam(Name, Context.getPointerType(Context.CharTy)); + addParam(Name, Context.getPointerType(Context.getAddrSpaceQualType( + Context.CharTy, LangAS::opencl_global))); } void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } From 32c98d9c984d76661a5eb9f41d63a9fadb46e241 Mon Sep 17 00:00:00 2001 From: Dmitry Vodopyanov Date: Mon, 12 Apr 2021 20:30:26 +0300 Subject: [PATCH 2/2] Fix tests --- clang/test/CodeGenSYCL/kernel-handler.cpp | 7 ++++--- clang/test/SemaSYCL/kernel-handler.cpp | 18 ++++++++++-------- 2 files changed, 14 insertions(+), 11 deletions(-) diff --git a/clang/test/CodeGenSYCL/kernel-handler.cpp b/clang/test/CodeGenSYCL/kernel-handler.cpp index 9c24acc31a74c..735b76028d13d 100644 --- a/clang/test/CodeGenSYCL/kernel-handler.cpp +++ b/clang/test/CodeGenSYCL/kernel-handler.cpp @@ -23,10 +23,11 @@ void test(int val) { } // NONATIVESUPPORT: define dso_local void @"{{.*}}test_kernel_handler{{.*}}" -// NONATIVESUPPORT-SAME: (i32 %_arg_, i8* %_arg__specialization_constants_buffer) +// NONATIVESUPPORT-SAME: (i32 %_arg_, i8 addrspace(1)* %_arg__specialization_constants_buffer) // NONATIVESUPPORT: %kh = alloca %"class.[[MANGLEDCLASS:[a-zA-Z0-9_]+]].cl::sycl::kernel_handler", align 1 -// NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8*, i8** %_arg__specialization_constants_buffer.addr, align 8 -// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler"* nonnull dereferenceable(1) %kh, i8* %[[KH]]) +// NONATIVESUPPORT: %[[KH:[0-9]+]] = load i8 addrspace(1)*, i8 addrspace(1)** %_arg__specialization_constants_buffer.addr, align 8 +// NONATIVESUPPORT: %[[ADDRSPACECAST:[0-9]+]] = addrspacecast i8 addrspace(1)* %[[KH]] to i8* +// NONATIVESUPPORT: call void @{{.*}}__init_specialization_constants_buffer{{.*}}(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler"* nonnull dereferenceable(1) %kh, i8* %[[ADDRSPACECAST]]) // NONATIVESUPPORT: void @"[[MANGLEDKERNELCALL:[a-zA-Z0-9_$]+]]" // NONATIVESUPPORT-SAME: byval(%"class.[[MANGLEDCLASS]].cl::sycl::kernel_handler") diff --git a/clang/test/SemaSYCL/kernel-handler.cpp b/clang/test/SemaSYCL/kernel-handler.cpp index a5df0e186e10b..a2ee5170d039f 100644 --- a/clang/test/SemaSYCL/kernel-handler.cpp +++ b/clang/test/SemaSYCL/kernel-handler.cpp @@ -28,9 +28,9 @@ int main() { } // Check test_kernel_handler parameters -// NONATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, char *)' +// NONATIVESUPPORT: FunctionDecl {{.*}}test_kernel_handler{{.*}} 'void (int, __global char *)' // NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__specialization_constants_buffer 'char *' +// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__specialization_constants_buffer '__global char *' // Check declaration and initialization of kernel object local clone // NONATIVESUPPORT-NEXT: CompoundStmt @@ -49,8 +49,9 @@ int main() { // NONATIVESUPPORT-NEXT: CXXMemberCallExpr {{.*}} 'void' // NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' -// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' -// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' 'char *' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} '__global char *' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' '__global char *' // NONATIVESUPPORT-NEXT: CompoundStmt // NONATIVESUPPORT-NEXT: CXXOperatorCallExpr // NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'void (*)(sycl::kernel_handler) const' @@ -63,9 +64,9 @@ int main() { // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' 'sycl::kernel_handler' // Check test_pfwg_kernel_handler parameters -// NONATIVESUPPORT: FunctionDecl {{.*}}test_pfwg_kernel_handler{{.*}} 'void (int, char *)' +// NONATIVESUPPORT: FunctionDecl {{.*}}test_pfwg_kernel_handler{{.*}} 'void (int, __global char *)' // NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg_ 'int' -// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__specialization_constants_buffer 'char *' +// NONATIVESUPPORT-NEXT: ParmVarDecl {{.*}} used _arg__specialization_constants_buffer '__global char *' // Check declaration and initialization of kernel object local clone // NONATIVESUPPORT-NEXT: CompoundStmt @@ -84,8 +85,9 @@ int main() { // NONATIVESUPPORT-NEXT: CXXMemberCallExpr {{.*}} 'void' // NONATIVESUPPORT-NEXT: MemberExpr {{.*}} 'void (char *)' lvalue .__init_specialization_constants_buffer // NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'sycl::kernel_handler' lvalue Var {{.*}} 'kh' -// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' -// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} 'char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' 'char *' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} 'char *' +// NONATIVESUPPORT-NEXT: ImplicitCastExpr {{.*}} '__global char *' +// NONATIVESUPPORT-NEXT: DeclRefExpr {{.*}} '__global char *' lvalue ParmVar {{.*}} '_arg__specialization_constants_buffer' '__global char *' // NONATIVESUPPORT-NEXT: CompoundStmt // NONATIVESUPPORT-NEXT: ExprWithCleanups // NONATIVESUPPORT-NEXT: CXXOperatorCallExpr