[SYCL] Make SYCL device compiler less dependent on SYCL accessor

class implementation. Main changes are: 1. In SYCL kernel entry point we initialize accessors through special accessor method __set_pointer. 2. Minor changes in lit tests. 3. Added hack to emit used top-level decls without CallGraph building. Signed-off-by: Vladimir Lazarev <vladimir.lazarev@intel.com>
intel · Jan 22, 2019 · f29e1cd · f29e1cd
1 parent 9a282fb
commit f29e1cd
Show file tree

Hide file tree

Showing 6 changed files with 74 additions and 38 deletions.
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -2876,7 +2876,11 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
       // Look for a declaration that's lexically in a record.
       for (const auto *FD = cast<FunctionDecl>(D)->getMostRecentDecl(); FD;
            FD = FD->getPreviousDecl()) {
-        if (isa<CXXRecordDecl>(FD->getLexicalDeclContext())) {
+        // For SYCL we also should emit a definition for a function because all
+        // top-level declarations without OpenCL kernel attribute are ignored
+        // now.
+        // TODO: fix this hack
+        if (isa<CXXRecordDecl>(FD->getLexicalDeclContext()) || LangOpts.SYCL) {
           if (FD->doesThisDeclarationHaveABody()) {
             addDeferredDeclToEmit(GD.getWithDecl(FD));
             break;

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
@@ -62,14 +62,15 @@ CompoundStmt *CreateSYCLKernelBody(Sema &S, CXXMemberCallExpr *e,
   if (LE) {
     // Create Lambda object
     CXXRecordDecl *LC = LE->getLambdaClass();
-    auto Lambda_VD = VarDecl::Create(
+    auto LambdaVD = VarDecl::Create(
         S.Context, DC, SourceLocation(), SourceLocation(), LC->getIdentifier(),
         QualType(LC->getTypeForDecl(), 0), LC->getLambdaTypeInfo(), SC_None);
+
     Stmt *DS = new (S.Context)
-        DeclStmt(DeclGroupRef(Lambda_VD), SourceLocation(), SourceLocation());
+        DeclStmt(DeclGroupRef(LambdaVD), SourceLocation(), SourceLocation());
     BodyStmts.push_back(DS);
-    auto Lambda_DRE = DeclRefExpr::Create(
-        S.Context, NestedNameSpecifierLoc(), SourceLocation(), Lambda_VD, false,
+    auto LambdaDRE = DeclRefExpr::Create(
+        S.Context, NestedNameSpecifierLoc(), SourceLocation(), LambdaVD, false,
         DeclarationNameInfo(), QualType(LC->getTypeForDecl(), 0), VK_LValue);
 
     // Init Lambda fields
@@ -78,37 +79,66 @@ CompoundStmt *CreateSYCLKernelBody(Sema &S, CXXMemberCallExpr *e,
     auto TargetFunc = dyn_cast<FunctionDecl>(DC);
     auto TargetFuncParam =
         TargetFunc->param_begin(); // Iterator to ParamVarDecl (VarDecl)
-    for (auto CaptureField : LE->captures()) {
-      VarDecl *CapturedVar =
-          CaptureField
-              .getCapturedVar(); // accessor, need to do setInit for this
+    for (auto Field : LC->fields()) {
       QualType ParamType = (*TargetFuncParam)->getOriginalType();
       auto DRE = DeclRefExpr::Create(
           S.Context, NestedNameSpecifierLoc(), SourceLocation(),
           *TargetFuncParam, false, DeclarationNameInfo(), ParamType, VK_LValue);
 
-      Expr *Res = ImplicitCastExpr::Create(
-          S.Context, ParamType, CK_LValueToRValue, DRE, nullptr, VK_RValue);
+      CXXRecordDecl *CRD = Field->getType()->getAsCXXRecordDecl();
+      if (CRD) {
+        llvm::SmallVector<Expr *, 16> ParamStmts;
+        DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none);
+        auto AccessorME = MemberExpr::Create(
+            S.Context, LambdaDRE, false, SourceLocation(),
+            NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP,
+            DeclarationNameInfo(Field->getDeclName(), SourceLocation()),
+            nullptr, Field->getType(), VK_LValue, OK_Ordinary);
+
+        for (auto Method : CRD->methods()) {
+          if (Method->getNameInfo().getName().getAsString() ==
+              "__set_pointer") {
+            DeclAccessPair MethodDAP = DeclAccessPair::make(Method, AS_none);
+            auto ME = MemberExpr::Create(
+                S.Context, AccessorME, false, SourceLocation(),
+                NestedNameSpecifierLoc(), SourceLocation(), Method, MethodDAP,
+                Method->getNameInfo(), nullptr, Method->getType(), VK_LValue,
+                OK_Ordinary);
+
+            // Not referenced -> not emitted
+            S.MarkFunctionReferenced(SourceLocation(), Method, true);
+
+            QualType ResultTy = Method->getReturnType();
+            ExprValueKind VK = Expr::getValueKindForType(ResultTy);
+            ResultTy = ResultTy.getNonLValueExprType(S.Context);
+
+            // __set_pointer needs one parameter
+            QualType paramTy = (*(Method->param_begin()))->getOriginalType();
 
-      Expr *InitCapture = new (S.Context) InitListExpr(
-          S.Context, SourceLocation(), /*initExprs*/ Res, SourceLocation());
-      CapturedVar->setInit(InitCapture);
-      InitCapture->setType(CapturedVar->getType());
-      InitCaptures.push_back(InitCapture);
+            // C++ address space attribute != opencl address space attribute
+            Expr *qualifiersCast = ImplicitCastExpr::Create(
+                S.Context, paramTy, CK_NoOp, DRE, nullptr, VK_LValue);
+            Expr *Res =
+                ImplicitCastExpr::Create(S.Context, paramTy, CK_LValueToRValue,
+                                         qualifiersCast, nullptr, VK_RValue);
+
+            ParamStmts.push_back(Res);
+
+            // lambda.accessor.__set_pointer(kernel_parameter)
+            CXXMemberCallExpr *Call = CXXMemberCallExpr::Create(
+                S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation());
+            BodyStmts.push_back(Call);
+          }
+        }
+      }
       TargetFuncParam++;
     }
 
-    Expr *InitLambdaCaptures = new (S.Context)
-        InitListExpr(S.Context, SourceLocation(), /*initExprs*/ InitCaptures,
-                     SourceLocation());
-    InitLambdaCaptures->setType(Lambda_VD->getType());
-    Lambda_VD->setInit(InitLambdaCaptures);
-
     // Create Lambda operator () call
     FunctionDecl *LO = LE->getCallOperator();
     ArrayRef<ParmVarDecl *> Args = LO->parameters();
     llvm::SmallVector<Expr *, 16> ParamStmts(1);
-    ParamStmts[0] = dyn_cast<Expr>(Lambda_DRE);
+    ParamStmts[0] = dyn_cast<Expr>(LambdaDRE);
 
     // Collect arguments for () operator
     for (auto Arg : Args) {

diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp
@@ -1,8 +1,9 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang -cc1 -DCL_TARGET_OPENCL_VERSION=220 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
 
 // CHECK: define {{.*}}spir_kernel void @kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] {
 // CHECK: ![[MD]] = !{}
-// XFAIL: *
+
+// XFAIL:*
 
 #include <CL/sycl.hpp>
 

diff --git a/clang/test/CodeGenSYCL/kernel-with-id.cpp b/clang/test/CodeGenSYCL/kernel-with-id.cpp
@@ -1,11 +1,11 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
-// XFAIL: *
-#include <CL/sycl.hpp>
+// RUN: %clang -cc1 -DCL_TARGET_OPENCL_VERSION=220 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
+
+// XFAIL:*
 
+#include <CL/sycl.hpp>
 #include <array>
 
-constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read;
-constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write;
+
 
 int main() {
   const size_t array_size = 1;
@@ -15,10 +15,11 @@ int main() {
   cl::sycl::buffer<cl::sycl::cl_int, 1> bufferA(A.data(), numOfItems);
 
   deviceQueue.submit([&](cl::sycl::handler &cgh) {
-    auto accessorA = bufferA.template get_access<sycl_read>(cgh);
-    // CHECK:  %wiID = alloca %"struct.cl::sycl::id", align 8
-    // CHECK:  call spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_2idILm1EEEE_clES5_"(%class.anon* %0, %"struct.cl::sycl::id"* byval align 8 %wiID)
-    // CHECK: %call = call spir_func i64 @_Z13get_global_idj(i32 0)
+    auto accessorA = bufferA.template get_access<cl::sycl::access::mode::read_write>(cgh);
+// CHECK: %wiID = alloca %"struct.cl::sycl::id", align 8
+// CHECK: call spir_func void @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0EE13__set_pointerEPU3AS1i(%"class.cl::sycl::accessor"* %1, i32 addrspace(1)* %2)
+// CHECK: call spir_func void @"_ZZZ4mainENK3$_0clERN2cl4sycl7handlerEENKUlNS1_2idILm1EEEE_clES5_"(%class.anon* %0, %"struct.cl::sycl::id"* byval align 8 %wiID)
+// CHECK: %call = call spir_func i64 @_Z13get_global_idj(i32 0)
     cgh.parallel_for<class kernel_function>(numOfItems,
       [=](cl::sycl::id<1> wiID) {
         accessorA[wiID] = accessorA[wiID] * accessorA[wiID];

diff --git a/clang/test/CodeGenSYCL/spir-calling-conv.cpp b/clang/test/CodeGenSYCL/spir-calling-conv.cpp
@@ -1,6 +1,6 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include  -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang -cc1 -DCL_TARGET_OPENCL_VERSION=220 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
 
-// XFAIL: *
+// XFAIL:*
 
 #include <CL/sycl.hpp>
 

diff --git a/clang/test/CodeGenSYCL/spir-opencl-version.cpp b/clang/test/CodeGenSYCL/spir-opencl-version.cpp
@@ -1,6 +1,6 @@
-// RUN: %clang -cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
+// RUN: %clang -cc1 -DCL_TARGET_OPENCL_VERSION=220 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -S -I /sycl_include_path -I /opencl_include_path -I /usr/include/c++/4.8.5 -I /usr/include/c++/4.8.5/x86_64-redhat-linux -I /usr/include/c++/4.8.5/backward -I /include -I /usr/include -fcxx-exceptions -fexceptions -emit-llvm -x c++ %s -o - | FileCheck %s
 
-// XFAIL: *
+// XFAIL:*
 
 #include <CL/sycl.hpp>