From 3b008831b7046a12d69635c373685aa34b37695f Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 8 Apr 2020 08:02:19 -0700 Subject: [PATCH 01/38] WIP on 'visitor' model for the SemaSYCL refactor Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 240 ++++++++++++++++++++++++++++-------- 1 file changed, 186 insertions(+), 54 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e795947add373..d86982980fba7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1256,16 +1256,6 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, } } -// Creates a mangled kernel name for given kernel name type -static std::string constructKernelName(QualType KernelNameType, - MangleContext &MC) { - SmallString<256> Result; - llvm::raw_svector_ostream Out(Result); - - MC.mangleTypeName(KernelNameType, Out); - return std::string(Out.str()); -} - static FunctionDecl * CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, ArrayRef ParamDescs) { @@ -1307,6 +1297,153 @@ CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, return OpenCLKernel; } +// The first template argument to the kernel function is used to identify the +// kernel itself. +static QualType calculateKernelNameType(ASTContext &Ctx, + FunctionDecl *KernelCallerFunc) { + // TODO: Not sure what the 'fully qualified type's purpose is here, the type + // itself should have its full qualified name, so figure out what the purpose + // is. + const TemplateArgumentList *TAL = + KernelCallerFunc->getTemplateSpecializationArgs(); + return TypeName::getFullyQualifiedType(TAL->get(0).getAsType(), Ctx, + /*WithGlobalNSPrefix=*/true); +} + +// Gets a name for the kernel caller func, calculated from the first template argument. +static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, + MangleContext &MC) { + QualType KernelNameType = + calculateKernelNameType(S.getASTContext(), KernelCallerFunc); + + if (S.getLangOpts().SYCLUnnamedLambda) + return PredefinedExpr::ComputeName( + S.getASTContext(), PredefinedExpr::UniqueStableNameType, KernelNameType); + + SmallString<256> Result; + llvm::raw_svector_ostream Out(Result); + + MC.mangleTypeName(KernelNameType, Out); + return std::string(Out.str()); +} + +// anonymous namespace so these don't get linkage. +namespace { +// A base type that the SYCL OpenCL Kernel construction task uses to implement +// individual tasks. +template +class SyclKernelFieldHandler { +protected: + Sema &SemaRef; + SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} + +public: + // TODO: Should the default behavior DO anything? Also, do these need any + // more params? + void handleSyclAccessorType(const FieldDecl *, QualType){} + void handleSyclSamplerType(const FieldDecl *, QualType){} + void handleSyclSpecConstantType(const FieldDecl *, QualType){} + void handleStructType(const FieldDecl *, QualType){} + void handleReferenceType(const FieldDecl *, QualType){} + void handlePointerType(const FieldDecl *, QualType){} + void handleArrayType(const FieldDecl *, QualType){} + void handleScalarType(const FieldDecl *, QualType){} +}; + +// A type to check the valididty of all of the argument types. +class SYCLKernelFieldChecker : public SyclKernelFieldHandler { + bool AllArgsValid = true; + DiagnosticsEngine &Diag; +public: + SYCLKernelFieldChecker(Sema &S) + : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} + bool isValid() { return AllArgsValid; } + + void handleSyclReferenceType(const FieldDecl *FD, QualType ArgTy) { + AllArgsValid = false; + Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) << ArgTy; + } + void handleSyclStructType(const FieldDecl *FD, QualType ArgTy) { + if (SemaRef.getASTContext().getLangOpts().SYCLStdLayoutKernelParams && + !ArgTy->isStandardLayoutType()) { + AllArgsValid = false; + Diag.Report(FD->getLocation(), diag::err_sycl_non_std_layout_type) + << ArgTy; + } else { + CXXRecordDecl *RD = ArgTy->getAsCXXRecordDecl(); + if (!RD->hasTrivialCopyConstructor()) { + AllArgsValid = false; + Diag.Report(FD->getLocation(), + diag::err_sycl_non_trivially_copy_ctor_dtor_type) + << 0 << ArgTy; + } else if (!RD->hasTrivialDestructor()) { + AllArgsValid = false; + Diag.Report(FD->getLocation(), + diag::err_sycl_non_trivially_copy_ctor_dtor_type) + << 1 << ArgTy; + } + } + } +}; + +// A type to Create and own the FunctionDecl for the kernel. +class SYCLKernelHeader : public SyclKernelFieldHandler { + SYCLKernelFieldChecker &ArgChecker; + FunctionDecl *KernelObj = nullptr; + +public: + // TODO: More params necessary for the kernel obj header. + SYCLKernelHeader(Sema &S, SYCLKernelFieldChecker &ArgChecker) + : SyclKernelFieldHandler(S), ArgChecker(ArgChecker) {} + ~SYCLKernelHeader() { + // TODO: Should 'body' do this? Does it need to do stuff in its destructor? + if (KernelObj && ArgChecker.isValid()) + SemaRef.addSyclDeviceDecl(KernelObj); + } +}; + +class SYCLKernelBody : public SyclKernelFieldHandler { + SYCLKernelHeader &Header; + +public: + SYCLKernelBody(Sema &S, SYCLKernelHeader &H) + : SyclKernelFieldHandler(S), Header(H) {} +}; +} + +template +void VisitKernelFields(RecordDecl::field_range Fields, Handlers& ... handlers) { + + // Implements the 'for-each-visitor' pattern. +#define KF_FOR_EACH(FUNC) \ + (void)std::initializer_list{(handlers.FUNC(Field, ArgTy), 0)...} + + for (const auto &Field : Fields){ + QualType ArgTy = Field->getType(); + + if (Util::isSyclAccessorType(ArgTy)) + KF_FOR_EACH(handleSyclAccessorType); + else if (Util::isSyclSamplerType(ArgTy)) + KF_FOR_EACH(handleSyclSamplerType); + else if (Util::isSyclSpecConstantType(ArgTy)) + KF_FOR_EACH(handleSyclSpecConstantType); + else if (ArgTy->isStructureOrClassType()) + KF_FOR_EACH(handleStructType); + else if (ArgTy->isReferenceType()) + KF_FOR_EACH(handleReferenceType); + else if (ArgTy->isPointerType()) + KF_FOR_EACH(handlePointerType); + else if (ArgTy->isArrayType()) + KF_FOR_EACH(handleArrayType); + else if (ArgTy->isScalarType()) + KF_FOR_EACH(handleScalarType); + else + llvm_unreachable("Unsupported kernel parameter type"); + } +#undef KF_FOR_EACH +} + + // Generates the OpenCL kernel using KernelCallerFunc (kernel caller // function) defined is SYCL headers. // Generated OpenCL kernel contains the body of the kernel caller function, @@ -1331,52 +1468,47 @@ CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, // void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC) { - CXXRecordDecl *LE = getKernelObjectType(KernelCallerFunc); - assert(LE && "invalid kernel caller"); + // The first argument to the KernelCallerFunc is the lambda object. + CXXRecordDecl *KernelLambda = getKernelObjectType(KernelCallerFunc); + assert(KernelLambda && "invalid kernel caller"); + std::string KernelName = constructKernelName(*this, KernelCallerFunc, MC); - // Build list of kernel arguments - llvm::SmallVector ParamDescs; - if (!buildArgTys(getASTContext(), LE, ParamDescs)) - return; + SYCLKernelFieldChecker checker(*this); + SYCLKernelHeader header(*this, checker); + SYCLKernelBody body(*this, header); - // Extract name from kernel caller parameters and mangle it. - const TemplateArgumentList *TemplateArgs = - KernelCallerFunc->getTemplateSpecializationArgs(); - assert(TemplateArgs && "No template argument info"); - QualType KernelNameType = TypeName::getFullyQualifiedType( - TemplateArgs->get(0).getAsType(), getASTContext(), true); - - std::string Name; - // TODO SYCLIntegrationHeader also computes a unique stable name. It should - // probably lose this responsibility and only use the name provided here. - if (getLangOpts().SYCLUnnamedLambda) - Name = PredefinedExpr::ComputeName( - getASTContext(), PredefinedExpr::UniqueStableNameExpr, KernelNameType); - else - Name = constructKernelName(KernelNameType, MC); - - // TODO Maybe don't emit integration header inside the Sema? - populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); - - FunctionDecl *OpenCLKernel = - CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); - - ContextRAII FuncContext(*this, OpenCLKernel); - - // Let's copy source location of a functor/lambda to emit nicer diagnostics - OpenCLKernel->setLocation(LE->getLocation()); - - // If the source function is implicitly inline, the kernel should be marked - // such as well. This allows the kernel to be ODR'd if there are multiple uses - // in different translation units. - OpenCLKernel->setImplicitlyInline(KernelCallerFunc->isInlined()); - - ConstructingOpenCLKernel = true; - CompoundStmt *OpenCLKernelBody = - CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); - ConstructingOpenCLKernel = false; - OpenCLKernel->setBody(OpenCLKernelBody); - addSyclDeviceDecl(OpenCLKernel); + VisitKernelFields(KernelLambda->fields(), checker, header, body); + + /* + // Build list of kernel arguments + llvm::SmallVector ParamDescs; + if (!buildArgTys(getASTContext(), LE, ParamDescs)) + return; + + // TODO Maybe don't emit integration header inside the Sema? + populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); + + FunctionDecl *OpenCLKernel = + CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); + + ContextRAII FuncContext(*this, OpenCLKernel); + + // Let's copy source location of a functor/lambda to emit nicer diagnostics + OpenCLKernel->setLocation(LE->getLocation()); + + // If the source function is implicitly inline, the kernel should be marked + // such as well. This allows the kernel to be ODR'd if there are multiple + uses + // in different translation units. + OpenCLKernel->setImplicitlyInline(KernelCallerFunc->isInlined()); + + ConstructingOpenCLKernel = true; + CompoundStmt *OpenCLKernelBody = + CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); + ConstructingOpenCLKernel = false; + OpenCLKernel->setBody(OpenCLKernelBody); + addSyclDeviceDecl(OpenCLKernel); + */ } void Sema::MarkDevice(void) { From 2ea64c80d1641b8ad49050955d000124832dc541 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Wed, 8 Apr 2020 08:48:11 -0700 Subject: [PATCH 02/38] clang format Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d86982980fba7..f5550c5e83df7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1310,15 +1310,17 @@ static QualType calculateKernelNameType(ASTContext &Ctx, /*WithGlobalNSPrefix=*/true); } -// Gets a name for the kernel caller func, calculated from the first template argument. +// Gets a name for the kernel caller func, calculated from the first template +// argument. static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, MangleContext &MC) { QualType KernelNameType = calculateKernelNameType(S.getASTContext(), KernelCallerFunc); if (S.getLangOpts().SYCLUnnamedLambda) - return PredefinedExpr::ComputeName( - S.getASTContext(), PredefinedExpr::UniqueStableNameType, KernelNameType); + return PredefinedExpr::ComputeName(S.getASTContext(), + PredefinedExpr::UniqueStableNameType, + KernelNameType); SmallString<256> Result; llvm::raw_svector_ostream Out(Result); From 2373337afa1a261ad22231b1a0ab9b2d0ece47a6 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 9 Apr 2020 18:53:10 -0700 Subject: [PATCH 03/38] Massive set of additions/changes. Implemneted the declaration generator. Still need IntHeader and kernel body generation Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 491 ++++++++++++++++++++---------------- 1 file changed, 279 insertions(+), 212 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f5550c5e83df7..e8d5a49f7cce4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1012,33 +1012,8 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { // in the following function we extract types of kernel object fields and add it // to the array with kernel parameters descriptors. // Returns true if all arguments are successfully built. -static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, +/*static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, SmallVectorImpl &ParamDescs) { - auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) { - // Create a parameter descriptor and append it to the result - ParamDescs.push_back(makeParamDesc(Fld, ArgType)); - }; - - // Creates a parameter descriptor for SYCL special object - SYCL accessor or - // sampler. - // All special SYCL objects must have __init method. We extract types for - // kernel parameters from __init method parameters. We will use __init method - // and kernel parameters which we build here to initialize special objects in - // the kernel body. - auto createSpecialSYCLObjParamDesc = [&](const FieldDecl *Fld, - const QualType &ArgTy) { - const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); - assert(RecordDecl && "Special SYCL object must be of a record type"); - - CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); - assert(InitMethod && "The accessor/sampler must have the __init method"); - unsigned NumParams = InitMethod->getNumParams(); - for (size_t I = 0; I < NumParams; ++I) { - ParmVarDecl *PD = InitMethod->getParamDecl(I); - CreateAndAddPrmDsc(Fld, PD->getType().getCanonicalType()); - } - }; - // Create parameter descriptor for accessor in case when it's wrapped with // some class. // TODO: Do we need support case when sampler is wrapped with some class or @@ -1079,65 +1054,16 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, QualType ArgTy = Fld->getType(); if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) { createSpecialSYCLObjParamDesc(Fld, ArgTy); - } else if (Util::isSyclSpecConstantType(ArgTy)) { - // Specialization constants are not added as arguments. } else if (ArgTy->isStructureOrClassType()) { - if (Context.getLangOpts().SYCLStdLayoutKernelParams) { - if (!ArgTy->isStandardLayoutType()) { - Context.getDiagnostics().Report(Fld->getLocation(), - diag::err_sycl_non_std_layout_type) - << ArgTy; - AllArgsAreValid = false; - continue; - } - } - - CXXRecordDecl *RD = - cast(ArgTy->getAs()->getDecl()); - if (!RD->hasTrivialCopyConstructor()) { - Context.getDiagnostics().Report( - Fld->getLocation(), - diag::err_sycl_non_trivially_copy_ctor_dtor_type) - << 0 << ArgTy; - AllArgsAreValid = false; - continue; - } - if (!RD->hasTrivialDestructor()) { - Context.getDiagnostics().Report( - Fld->getLocation(), - diag::err_sycl_non_trivially_copy_ctor_dtor_type) - << 1 << ArgTy; - AllArgsAreValid = false; - continue; - } - CreateAndAddPrmDsc(Fld, ArgTy); // Create descriptors for each accessor field in the class or struct createParamDescForWrappedAccessors(Fld, ArgTy); - } else if (ArgTy->isReferenceType()) { - Context.getDiagnostics().Report( - Fld->getLocation(), diag::err_bad_kernel_param_type) << ArgTy; - AllArgsAreValid = false; - } else if (ArgTy->isPointerType()) { - // Pointer Arguments need to be in the global address space - QualType PointeeTy = ArgTy->getPointeeType(); - Qualifiers Quals = PointeeTy.getQualifiers(); - Quals.setAddressSpace(LangAS::opencl_global); - PointeeTy = - Context.getQualifiedType(PointeeTy.getUnqualifiedType(), Quals); - QualType ModTy = Context.getPointerType(PointeeTy); - - CreateAndAddPrmDsc(Fld, ModTy); - } else if (ArgTy->isScalarType()) { - CreateAndAddPrmDsc(Fld, ArgTy); - } else { - llvm_unreachable("Unsupported kernel parameter type"); - } } return AllArgsAreValid; } +*/ /// Adds necessary data describing given kernel to the integration header. /// \param H the integration header object @@ -1256,47 +1182,6 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, } } -static FunctionDecl * -CreateOpenCLKernelDeclaration(ASTContext &Context, StringRef Name, - ArrayRef ParamDescs) { - - DeclContext *DC = Context.getTranslationUnitDecl(); - QualType RetTy = Context.VoidTy; - SmallVector ArgTys; - - // Extract argument types from the descriptor array: - std::transform( - ParamDescs.begin(), ParamDescs.end(), std::back_inserter(ArgTys), - [](const ParamDesc &PD) -> QualType { return std::get<0>(PD); }); - FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); - QualType FuncTy = Context.getFunctionType(RetTy, ArgTys, Info); - DeclarationName DN = DeclarationName(&Context.Idents.get(Name)); - - FunctionDecl *OpenCLKernel = FunctionDecl::Create( - Context, DC, SourceLocation(), SourceLocation(), DN, FuncTy, - Context.getTrivialTypeSourceInfo(RetTy), SC_None); - - llvm::SmallVector Params; - int i = 0; - for (const auto &PD : ParamDescs) { - auto P = ParmVarDecl::Create(Context, OpenCLKernel, SourceLocation(), - SourceLocation(), std::get<1>(PD), - std::get<0>(PD), std::get<2>(PD), SC_None, 0); - P->setScopeInfo(0, i++); - P->setIsUsed(); - Params.push_back(P); - } - OpenCLKernel->setParams(Params); - - OpenCLKernel->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); - OpenCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); - OpenCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context)); - - // Add kernel to translation unit to see it in AST-dump - DC->addDecl(OpenCLKernel); - return OpenCLKernel; -} - // The first template argument to the kernel function is used to identify the // kernel itself. static QualType calculateKernelNameType(ASTContext &Ctx, @@ -1331,6 +1216,39 @@ static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, // anonymous namespace so these don't get linkage. namespace { +template +static void VisitKernelFields(RecordDecl::field_range Fields, + Handlers &... handlers) { + +// Implements the 'for-each-visitor' pattern. +#define KF_FOR_EACH(FUNC) \ + (void)std::initializer_list { (handlers.FUNC(Field, ArgTy), 0)... } + + for (const auto &Field : Fields) { + QualType ArgTy = Field->getType(); + + if (Util::isSyclAccessorType(ArgTy)) + KF_FOR_EACH(handleSyclAccessorType); + else if (Util::isSyclSamplerType(ArgTy)) + KF_FOR_EACH(handleSyclSamplerType); + else if (Util::isSyclSpecConstantType(ArgTy)) + KF_FOR_EACH(handleSyclSpecConstantType); + else if (ArgTy->isStructureOrClassType()) + KF_FOR_EACH(handleStructType); + else if (ArgTy->isReferenceType()) + KF_FOR_EACH(handleReferenceType); + else if (ArgTy->isPointerType()) + KF_FOR_EACH(handlePointerType); + else if (ArgTy->isArrayType()) + KF_FOR_EACH(handleArrayType); + else if (ArgTy->isScalarType()) + KF_FOR_EACH(handleScalarType); + else + KF_FOR_EACH(handleOtherType); + } +#undef KF_FOR_EACH +} + // A base type that the SYCL OpenCL Kernel construction task uses to implement // individual tasks. template @@ -1340,120 +1258,263 @@ class SyclKernelFieldHandler { SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} public: - // TODO: Should the default behavior DO anything? Also, do these need any - // more params? - void handleSyclAccessorType(const FieldDecl *, QualType){} - void handleSyclSamplerType(const FieldDecl *, QualType){} - void handleSyclSpecConstantType(const FieldDecl *, QualType){} - void handleStructType(const FieldDecl *, QualType){} - void handleReferenceType(const FieldDecl *, QualType){} - void handlePointerType(const FieldDecl *, QualType){} - void handleArrayType(const FieldDecl *, QualType){} - void handleScalarType(const FieldDecl *, QualType){} + // Mark these virutal so that we can use override in the implementer classes, + // despite virtual dispatch never being used. + virtual void handleSyclAccessorType(const FieldDecl *, QualType){} + virtual void handleSyclSamplerType(const FieldDecl *, QualType){} + virtual void handleSyclSpecConstantType(const FieldDecl *, QualType){} + virtual void handleStructType(const FieldDecl *, QualType){} + virtual void handleReferenceType(const FieldDecl *, QualType){} + virtual void handlePointerType(const FieldDecl *, QualType){} + virtual void handleArrayType(const FieldDecl *, QualType){} + virtual void handleScalarType(const FieldDecl *, QualType){} + // Most handlers shouldn't be handling this, just the field checker. + virtual void handleOtherType(const FieldDecl *, QualType){} }; // A type to check the valididty of all of the argument types. -class SYCLKernelFieldChecker : public SyclKernelFieldHandler { - bool AllArgsValid = true; +class SyclKernelFieldChecker + : public SyclKernelFieldHandler { + bool IsInvalid = false; DiagnosticsEngine &Diag; public: - SYCLKernelFieldChecker(Sema &S) + SyclKernelFieldChecker(Sema &S) : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} - bool isValid() { return AllArgsValid; } + bool isValid() { return !IsInvalid; } - void handleSyclReferenceType(const FieldDecl *FD, QualType ArgTy) { - AllArgsValid = false; - Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) << ArgTy; + void handleReferenceType(const FieldDecl *FD, QualType ArgTy) final { + IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) + << ArgTy; } - void handleSyclStructType(const FieldDecl *FD, QualType ArgTy) { + void handleStructType(const FieldDecl *FD, QualType ArgTy) final { if (SemaRef.getASTContext().getLangOpts().SYCLStdLayoutKernelParams && - !ArgTy->isStandardLayoutType()) { - AllArgsValid = false; - Diag.Report(FD->getLocation(), diag::err_sycl_non_std_layout_type) + !ArgTy->isStandardLayoutType()) + IsInvalid = + Diag.Report(FD->getLocation(), diag::err_sycl_non_std_layout_type) << ArgTy; - } else { + else { CXXRecordDecl *RD = ArgTy->getAsCXXRecordDecl(); - if (!RD->hasTrivialCopyConstructor()) { - AllArgsValid = false; - Diag.Report(FD->getLocation(), - diag::err_sycl_non_trivially_copy_ctor_dtor_type) + if (!RD->hasTrivialCopyConstructor()) + + IsInvalid = + Diag.Report(FD->getLocation(), + diag::err_sycl_non_trivially_copy_ctor_dtor_type) << 0 << ArgTy; - } else if (!RD->hasTrivialDestructor()) { - AllArgsValid = false; - Diag.Report(FD->getLocation(), - diag::err_sycl_non_trivially_copy_ctor_dtor_type) + else if (!RD->hasTrivialDestructor()) + IsInvalid = + Diag.Report(FD->getLocation(), + diag::err_sycl_non_trivially_copy_ctor_dtor_type) << 1 << ArgTy; - } } } + + // We should be able to ahndle this, so we made it part of the visitor, but + // this is 'to be implemented'. + void handleArrayType(const FieldDecl *FD, QualType ArgTy) final { + IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) + << ArgTy; + } + + void handleOtherType(const FieldDecl *FD, QualType ArgTy) final { + IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) + << ArgTy; + } }; -// A type to Create and own the FunctionDecl for the kernel. -class SYCLKernelHeader : public SyclKernelFieldHandler { - SYCLKernelFieldChecker &ArgChecker; - FunctionDecl *KernelObj = nullptr; +// A type that handles the accessor recursion and acts as a base for +// SyclKernelDeclCreator. It doesn't 'own' anything other than the KernelObj +// pointer and functionality required to add a param. +class SyclKernelDeclBase + : public SyclKernelFieldHandler { +protected: + FunctionDecl *KernelObj; + llvm::SmallVectorImpl &ArgTys; + llvm::SmallVectorImpl &Params; + + void addParam(const FieldDecl *FD, QualType ArgTy) { + ASTContext &Ctx = SemaRef.getASTContext(); + // TODO: should we split this function up? These ops NEED to happen in + // lockstep, so leaning toward leaving htis as just a somewhat long function + // :/ + + // Create a new ParmVarDecl based on the new info. + ParamDesc newParamDesc = makeParamDesc(FD, ArgTy); + auto *NewParam = ParmVarDecl::Create( + SemaRef.getASTContext(), KernelObj, SourceLocation(), SourceLocation(), + std::get<1>(newParamDesc), std::get<0>(newParamDesc), + std::get<2>(newParamDesc), SC_None, /*DefArg*/ nullptr); + + NewParam->setScopeInfo(0, Params.size()); + NewParam->setIsUsed(); + + Params.push_back(NewParam); + + // Create the new type. + ArgTys.push_back(ArgTy); + FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); + + // Set type, note until the destructor of the owner of KernelObj is called + // (to set the parameters), we cannot access the parameters from KernelObj + // without memory problems. + KernelObj->setType(FuncType); + } + + // All special SYCL objects must have __init method. We extract types for + // kernel parameters from __init method parameters. We will use __init method + // and kernel parameters which we build here to initialize special objects in + // the kernel body. + void handleSpecialType(const FieldDecl *FD, QualType ArgTy) { + const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); + assert(RecordDecl && "Special SYCL object must be of a record type"); + + CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); + assert(InitMethod && "The accessor/sampler must have the __init method"); + + for (const ParmVarDecl *Param : InitMethod->parameters()) + addParam(FD, Param->getType().getCanonicalType()); + } + public: - // TODO: More params necessary for the kernel obj header. - SYCLKernelHeader(Sema &S, SYCLKernelFieldChecker &ArgChecker) - : SyclKernelFieldHandler(S), ArgChecker(ArgChecker) {} - ~SYCLKernelHeader() { - // TODO: Should 'body' do this? Does it need to do stuff in its destructor? - if (KernelObj && ArgChecker.isValid()) - SemaRef.addSyclDeviceDecl(KernelObj); + SyclKernelDeclBase(Sema &S, FunctionDecl *KernelObj, + llvm::SmallVectorImpl &ArgTys, + llvm::SmallVectorImpl &Params) + : SyclKernelFieldHandler(S), KernelObj(KernelObj), ArgTys(ArgTys), + Params(Params) {} + + void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { + handleSpecialType(FD, ArgTy); + } + + void handleStructType(const FieldDecl *FD, QualType ArgTy) override { + addParam(FD, ArgTy); + + // Create descriptors for each accessor field in the class or struct + const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); + VisitKernelFields(Wrapper->fields(), *this); } }; -class SYCLKernelBody : public SyclKernelFieldHandler { - SYCLKernelHeader &Header; +// A type to Create and own the FunctionDecl for the kernel. +class SyclKernelDeclCreator + : public SyclKernelDeclBase { + // TODO: rather than this, should we consider a 'commit' function that + // finalizes under success only? + SyclKernelFieldChecker &ArgChecker; + Sema::ContextRAII FuncContext; + // Yes, the list of Parameters contains this info, but we use it often enough + // we shouldn't be recreating it constantly. QualTypes are cheap anyway. + llvm::SmallVector ArgTys; + llvm::SmallVector Params; + + static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD, StringRef Name) { + // Set implict attributes. + FD->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); + FD->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); + FD->addAttr(ArtificialAttr::CreateImplicit(Context)); + } + + static FunctionDecl *initKernelObj(ASTContext &Ctx, StringRef Name, + SourceLocation Loc, bool IsInline) { + // Create this with no prototype, and we can fix this up after we've seen + // all the params. + FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, {}, Info); + + FunctionDecl *FD = FunctionDecl::Create( + Ctx, Ctx.getTranslationUnitDecl(), Loc, Loc, &Ctx.Idents.get(Name), + FuncType, Ctx.getTrivialTypeSourceInfo(Ctx.VoidTy), SC_None); + FD->setImplicitlyInline(IsInline); + setKernelImplicitAttrs(Ctx, FD, Name); + + // Add kernel to translation unit to see it in AST-dump. + Ctx.getTranslationUnitDecl()->addDecl(FD); + return FD; + } public: - SYCLKernelBody(Sema &S, SYCLKernelHeader &H) - : SyclKernelFieldHandler(S), Header(H) {} -}; -} + SyclKernelDeclCreator(Sema &S, SyclKernelFieldChecker &ArgChecker, + StringRef Name, SourceLocation Loc, bool IsInline) + : SyclKernelDeclBase( + S, initKernelObj(S.getASTContext(), Name, Loc, IsInline), ArgTys, + Params), + ArgChecker(ArgChecker), FuncContext(SemaRef, KernelObj) {} -template -void VisitKernelFields(RecordDecl::field_range Fields, Handlers& ... handlers) { + ~SyclKernelDeclCreator() { + KernelObj->setParams(Params); - // Implements the 'for-each-visitor' pattern. -#define KF_FOR_EACH(FUNC) \ - (void)std::initializer_list{(handlers.FUNC(Field, ArgTy), 0)...} + if (ArgChecker.isValid()) + SemaRef.addSyclDeviceDecl(KernelObj); + } - for (const auto &Field : Fields){ - QualType ArgTy = Field->getType(); + void handleSyclSamplerType(const FieldDecl *FD, QualType ArgTy) final { + handleSpecialType(FD, ArgTy); + } - if (Util::isSyclAccessorType(ArgTy)) - KF_FOR_EACH(handleSyclAccessorType); - else if (Util::isSyclSamplerType(ArgTy)) - KF_FOR_EACH(handleSyclSamplerType); - else if (Util::isSyclSpecConstantType(ArgTy)) - KF_FOR_EACH(handleSyclSpecConstantType); - else if (ArgTy->isStructureOrClassType()) - KF_FOR_EACH(handleStructType); - else if (ArgTy->isReferenceType()) - KF_FOR_EACH(handleReferenceType); - else if (ArgTy->isPointerType()) - KF_FOR_EACH(handlePointerType); - else if (ArgTy->isArrayType()) - KF_FOR_EACH(handleArrayType); - else if (ArgTy->isScalarType()) - KF_FOR_EACH(handleScalarType); - else - llvm_unreachable("Unsupported kernel parameter type"); + void handlePointerType(const FieldDecl *FD, QualType ArgTy) final { + // TODO: Can we document what the heck this is doing?! + QualType PointeeTy = ArgTy->getPointeeType(); + Qualifiers Quals = PointeeTy.getQualifiers(); + Quals.setAddressSpace(LangAS::opencl_global); + PointeeTy = SemaRef.getASTContext().getQualifiedType( + PointeeTy.getUnqualifiedType(), Quals); + QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy); + addParam(FD, ModTy); } -#undef KF_FOR_EACH + void handleScalarType(const FieldDecl *FD, QualType ArgTy) final { + addParam(FD, ArgTy); + } + + // This is implemented here because this is the only case where the recurse + // object is required. The base type is pretty cheap, so we might opt + // to just always create it (the way this one is implemented) and just put + // this implementation in the base. + void handleStructType(const FieldDecl *FD, QualType ArgTy) final { + addParam(FD, ArgTy); + + // Create descriptors for each accessor field in the class or struct + const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); + SyclKernelDeclBase Recurse(SemaRef, KernelObj, ArgTys, Params); + VisitKernelFields(Wrapper->fields(), Recurse); + } + + void setBody(CompoundStmt *KB) { + KernelObj->setBody(KB); + } +}; + +class SyclKernelBodyCreator + : public SyclKernelFieldHandler { + SyclKernelDeclCreator &DeclCreator; + // TODO: When/Where does this get created? + CompoundStmt *KernelBody = nullptr; + +public: + SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC) + : SyclKernelFieldHandler(S), DeclCreator(DC) { + } + ~SyclKernelBodyCreator() { + DeclCreator.setBody(KernelBody); + } +}; + +class SyclKernelIntHeaderCreator + : public SyclKernelFieldHandler { + public: + SyclKernelIntHeaderCreator(Sema &S) : SyclKernelFieldHandler(S) {} +}; } // Generates the OpenCL kernel using KernelCallerFunc (kernel caller -// function) defined is SYCL headers. +// function) defined is Sycl headers. // Generated OpenCL kernel contains the body of the kernel caller function, // receives OpenCL like parameters and additionally does some manipulation to // initialize captured lambda/functor fields with these parameters. -// SYCL runtime marks kernel caller function with sycl_kernel attribute. +// Sycl runtime marks kernel caller function with sycl_kernel attribute. // To be able to generate OpenCL kernel from KernelCallerFunc we put -// the following requirements to the function which SYCL runtime can mark with +// the following requirements to the function which Sycl runtime can mark with // sycl_kernel attribute: // - Must be template function with at least two template parameters. // First parameter must represent "unique kernel name" @@ -1475,41 +1536,47 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, assert(KernelLambda && "invalid kernel caller"); std::string KernelName = constructKernelName(*this, KernelCallerFunc, MC); - SYCLKernelFieldChecker checker(*this); - SYCLKernelHeader header(*this, checker); - SYCLKernelBody body(*this, header); + SyclKernelFieldChecker checker(*this); + SyclKernelDeclCreator kernel_decl(*this, checker, KernelName, + KernelLambda->getLocation(), + KernelCallerFunc->isInlined()); + SyclKernelBodyCreator kernel_body(*this, kernel_decl); + SyclKernelIntHeaderCreator int_header(*this); - VisitKernelFields(KernelLambda->fields(), checker, header, body); + ConstructingOpenCLKernel = true; + VisitKernelFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, + int_header); + ConstructingOpenCLKernel = false; /* // Build list of kernel arguments - llvm::SmallVector ParamDescs; - if (!buildArgTys(getASTContext(), LE, ParamDescs)) - return; + //llvm::SmallVector ParamDescs; + //if (!buildArgTys(getASTContext(), LE, ParamDescs)) + // return; // TODO Maybe don't emit integration header inside the Sema? - populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); + ***populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); - FunctionDecl *OpenCLKernel = - CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); + //FunctionDecl *OpenCLKernel = + // CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); - ContextRAII FuncContext(*this, OpenCLKernel); + //ContextRAII FuncContext(*this, OpenCLKernel); // Let's copy source location of a functor/lambda to emit nicer diagnostics - OpenCLKernel->setLocation(LE->getLocation()); + //OpenCLKernel->setLocation(LE->getLocation()); // If the source function is implicitly inline, the kernel should be marked // such as well. This allows the kernel to be ODR'd if there are multiple uses // in different translation units. - OpenCLKernel->setImplicitlyInline(KernelCallerFunc->isInlined()); + //OpenCLKernel->setImplicitlyInline(KernelCallerFunc->isInlined()); - ConstructingOpenCLKernel = true; - CompoundStmt *OpenCLKernelBody = + //ConstructingOpenCLKernel = true; + ****CompoundStmt *OpenCLKernelBody = CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); - ConstructingOpenCLKernel = false; - OpenCLKernel->setBody(OpenCLKernelBody); - addSyclDeviceDecl(OpenCLKernel); + //ConstructingOpenCLKernel = false; + //OpenCLKernel->setBody(OpenCLKernelBody); + //addSyclDeviceDecl(OpenCLKernel); */ } From 4b5abb1da82c79e3121ad8b03c3d4b15875a6d0c Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 9 Apr 2020 18:57:23 -0700 Subject: [PATCH 04/38] clang-format changes Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index e8d5a49f7cce4..a092bbe40ba32 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1260,16 +1260,16 @@ class SyclKernelFieldHandler { public: // Mark these virutal so that we can use override in the implementer classes, // despite virtual dispatch never being used. - virtual void handleSyclAccessorType(const FieldDecl *, QualType){} - virtual void handleSyclSamplerType(const FieldDecl *, QualType){} - virtual void handleSyclSpecConstantType(const FieldDecl *, QualType){} - virtual void handleStructType(const FieldDecl *, QualType){} - virtual void handleReferenceType(const FieldDecl *, QualType){} - virtual void handlePointerType(const FieldDecl *, QualType){} - virtual void handleArrayType(const FieldDecl *, QualType){} - virtual void handleScalarType(const FieldDecl *, QualType){} + virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} + virtual void handleSyclSamplerType(const FieldDecl *, QualType) {} + virtual void handleSyclSpecConstantType(const FieldDecl *, QualType) {} + virtual void handleStructType(const FieldDecl *, QualType) {} + virtual void handleReferenceType(const FieldDecl *, QualType) {} + virtual void handlePointerType(const FieldDecl *, QualType) {} + virtual void handleArrayType(const FieldDecl *, QualType) {} + virtual void handleScalarType(const FieldDecl *, QualType) {} // Most handlers shouldn't be handling this, just the field checker. - virtual void handleOtherType(const FieldDecl *, QualType){} + virtual void handleOtherType(const FieldDecl *, QualType) {} }; // A type to check the valididty of all of the argument types. From 3ef047327358bfa3e5dbc41bb090dedb70cf4ca3 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 9 Apr 2020 19:02:19 -0700 Subject: [PATCH 05/38] more clang-format changes Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 36 ++++++++++++++---------------------- 1 file changed, 14 insertions(+), 22 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a092bbe40ba32..9741fb7468e15 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1251,8 +1251,7 @@ static void VisitKernelFields(RecordDecl::field_range Fields, // A base type that the SYCL OpenCL Kernel construction task uses to implement // individual tasks. -template -class SyclKernelFieldHandler { +template class SyclKernelFieldHandler { protected: Sema &SemaRef; SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} @@ -1277,6 +1276,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { bool IsInvalid = false; DiagnosticsEngine &Diag; + public: SyclKernelFieldChecker(Sema &S) : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} @@ -1324,12 +1324,11 @@ class SyclKernelFieldChecker // A type that handles the accessor recursion and acts as a base for // SyclKernelDeclCreator. It doesn't 'own' anything other than the KernelObj // pointer and functionality required to add a param. -class SyclKernelDeclBase - : public SyclKernelFieldHandler { +class SyclKernelDeclBase : public SyclKernelFieldHandler { protected: FunctionDecl *KernelObj; llvm::SmallVectorImpl &ArgTys; - llvm::SmallVectorImpl &Params; + llvm::SmallVectorImpl &Params; void addParam(const FieldDecl *FD, QualType ArgTy) { ASTContext &Ctx = SemaRef.getASTContext(); @@ -1375,7 +1374,6 @@ class SyclKernelDeclBase addParam(FD, Param->getType().getCanonicalType()); } - public: SyclKernelDeclBase(Sema &S, FunctionDecl *KernelObj, llvm::SmallVectorImpl &ArgTys, @@ -1397,8 +1395,7 @@ class SyclKernelDeclBase }; // A type to Create and own the FunctionDecl for the kernel. -class SyclKernelDeclCreator - : public SyclKernelDeclBase { +class SyclKernelDeclCreator : public SyclKernelDeclBase { // TODO: rather than this, should we consider a 'commit' function that // finalizes under success only? SyclKernelFieldChecker &ArgChecker; @@ -1406,9 +1403,10 @@ class SyclKernelDeclCreator // Yes, the list of Parameters contains this info, but we use it often enough // we shouldn't be recreating it constantly. QualTypes are cheap anyway. llvm::SmallVector ArgTys; - llvm::SmallVector Params; + llvm::SmallVector Params; - static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD, StringRef Name) { + static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD, + StringRef Name) { // Set implict attributes. FD->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); FD->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); @@ -1479,9 +1477,7 @@ class SyclKernelDeclCreator VisitKernelFields(Wrapper->fields(), Recurse); } - void setBody(CompoundStmt *KB) { - KernelObj->setBody(KB); - } + void setBody(CompoundStmt *KB) { KernelObj->setBody(KB); } }; class SyclKernelBodyCreator @@ -1492,20 +1488,16 @@ class SyclKernelBodyCreator public: SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC) - : SyclKernelFieldHandler(S), DeclCreator(DC) { - } - ~SyclKernelBodyCreator() { - DeclCreator.setBody(KernelBody); - } + : SyclKernelFieldHandler(S), DeclCreator(DC) {} + ~SyclKernelBodyCreator() { DeclCreator.setBody(KernelBody); } }; class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { - public: - SyclKernelIntHeaderCreator(Sema &S) : SyclKernelFieldHandler(S) {} +public: + SyclKernelIntHeaderCreator(Sema &S) : SyclKernelFieldHandler(S) {} }; -} - +} // namespace // Generates the OpenCL kernel using KernelCallerFunc (kernel caller // function) defined is Sycl headers. From 07bb9beb514f9243c095c75fb4ce55a344e50ec1 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 06:37:40 -0700 Subject: [PATCH 06/38] A few spelling changes, move the kernel type changing Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 56 +++++++++++++++++-------------------- 1 file changed, 26 insertions(+), 30 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9741fb7468e15..8381eab8a43a4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1216,32 +1216,33 @@ static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, // anonymous namespace so these don't get linkage. namespace { +// A visitor function that dispatches to functions as defined in +// SyclKernelFieldHandler for the purposes of kernel generation. template -static void VisitKernelFields(RecordDecl::field_range Fields, +static void VisitRecordFields(RecordDecl::field_range Fields, Handlers &... handlers) { - // Implements the 'for-each-visitor' pattern. #define KF_FOR_EACH(FUNC) \ - (void)std::initializer_list { (handlers.FUNC(Field, ArgTy), 0)... } + (void)std::initializer_list { (handlers.FUNC(Field, FieldTy), 0)... } for (const auto &Field : Fields) { - QualType ArgTy = Field->getType(); + QualType FieldTy = Field->getType(); - if (Util::isSyclAccessorType(ArgTy)) + if (Util::isSyclAccessorType(FieldTy)) KF_FOR_EACH(handleSyclAccessorType); - else if (Util::isSyclSamplerType(ArgTy)) + else if (Util::isSyclSamplerType(FieldTy)) KF_FOR_EACH(handleSyclSamplerType); - else if (Util::isSyclSpecConstantType(ArgTy)) + else if (Util::isSyclSpecConstantType(FieldTy)) KF_FOR_EACH(handleSyclSpecConstantType); - else if (ArgTy->isStructureOrClassType()) + else if (FieldTy->isStructureOrClassType()) KF_FOR_EACH(handleStructType); - else if (ArgTy->isReferenceType()) + else if (FieldTy->isReferenceType()) KF_FOR_EACH(handleReferenceType); - else if (ArgTy->isPointerType()) + else if (FieldTy->isPointerType()) KF_FOR_EACH(handlePointerType); - else if (ArgTy->isArrayType()) + else if (FieldTy->isArrayType()) KF_FOR_EACH(handleArrayType); - else if (ArgTy->isScalarType()) + else if (FieldTy->isScalarType()) KF_FOR_EACH(handleScalarType); else KF_FOR_EACH(handleOtherType); @@ -1331,7 +1332,6 @@ class SyclKernelDeclBase : public SyclKernelFieldHandler { llvm::SmallVectorImpl &Params; void addParam(const FieldDecl *FD, QualType ArgTy) { - ASTContext &Ctx = SemaRef.getASTContext(); // TODO: should we split this function up? These ops NEED to happen in // lockstep, so leaning toward leaving htis as just a somewhat long function // :/ @@ -1348,15 +1348,7 @@ class SyclKernelDeclBase : public SyclKernelFieldHandler { Params.push_back(NewParam); - // Create the new type. ArgTys.push_back(ArgTy); - FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); - QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); - - // Set type, note until the destructor of the owner of KernelObj is called - // (to set the parameters), we cannot access the parameters from KernelObj - // without memory problems. - KernelObj->setType(FuncType); } // All special SYCL objects must have __init method. We extract types for @@ -1390,7 +1382,7 @@ class SyclKernelDeclBase : public SyclKernelFieldHandler { // Create descriptors for each accessor field in the class or struct const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); - VisitKernelFields(Wrapper->fields(), *this); + VisitRecordFields(Wrapper->fields(), *this); } }; @@ -1413,8 +1405,8 @@ class SyclKernelDeclCreator : public SyclKernelDeclBase { FD->addAttr(ArtificialAttr::CreateImplicit(Context)); } - static FunctionDecl *initKernelObj(ASTContext &Ctx, StringRef Name, - SourceLocation Loc, bool IsInline) { + static FunctionDecl *createKernelDecl(ASTContext &Ctx, StringRef Name, + SourceLocation Loc, bool IsInline) { // Create this with no prototype, and we can fix this up after we've seen // all the params. FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); @@ -1435,11 +1427,15 @@ class SyclKernelDeclCreator : public SyclKernelDeclBase { SyclKernelDeclCreator(Sema &S, SyclKernelFieldChecker &ArgChecker, StringRef Name, SourceLocation Loc, bool IsInline) : SyclKernelDeclBase( - S, initKernelObj(S.getASTContext(), Name, Loc, IsInline), ArgTys, + S, createKernelDecl(S.getASTContext(), Name, Loc, IsInline), ArgTys, Params), ArgChecker(ArgChecker), FuncContext(SemaRef, KernelObj) {} ~SyclKernelDeclCreator() { + ASTContext &Ctx = SemaRef.getASTContext(); + FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); + KernelObj->setType(FuncType); KernelObj->setParams(Params); if (ArgChecker.isValid()) @@ -1474,7 +1470,7 @@ class SyclKernelDeclCreator : public SyclKernelDeclBase { // Create descriptors for each accessor field in the class or struct const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); SyclKernelDeclBase Recurse(SemaRef, KernelObj, ArgTys, Params); - VisitKernelFields(Wrapper->fields(), Recurse); + VisitRecordFields(Wrapper->fields(), Recurse); } void setBody(CompoundStmt *KB) { KernelObj->setBody(KB); } @@ -1500,13 +1496,13 @@ class SyclKernelIntHeaderCreator } // namespace // Generates the OpenCL kernel using KernelCallerFunc (kernel caller -// function) defined is Sycl headers. +// function) defined is SYCL headers. // Generated OpenCL kernel contains the body of the kernel caller function, // receives OpenCL like parameters and additionally does some manipulation to // initialize captured lambda/functor fields with these parameters. -// Sycl runtime marks kernel caller function with sycl_kernel attribute. +// SYCL runtime marks kernel caller function with sycl_kernel attribute. // To be able to generate OpenCL kernel from KernelCallerFunc we put -// the following requirements to the function which Sycl runtime can mark with +// the following requirements to the function which SYCL runtime can mark with // sycl_kernel attribute: // - Must be template function with at least two template parameters. // First parameter must represent "unique kernel name" @@ -1536,7 +1532,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, SyclKernelIntHeaderCreator int_header(*this); ConstructingOpenCLKernel = true; - VisitKernelFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, + VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, int_header); ConstructingOpenCLKernel = false; From ec1e68d54d82f2c2845d1fe667be854234fc6a9b Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 07:31:58 -0700 Subject: [PATCH 07/38] remove TODO comment, the function has now returned to a managable size :) --- clang/lib/Sema/SemaSYCL.cpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8381eab8a43a4..61464d292965e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1332,10 +1332,6 @@ class SyclKernelDeclBase : public SyclKernelFieldHandler { llvm::SmallVectorImpl &Params; void addParam(const FieldDecl *FD, QualType ArgTy) { - // TODO: should we split this function up? These ops NEED to happen in - // lockstep, so leaning toward leaving htis as just a somewhat long function - // :/ - // Create a new ParmVarDecl based on the new info. ParamDesc newParamDesc = makeParamDesc(FD, ArgTy); auto *NewParam = ParmVarDecl::Create( From ec93877a191d51cf468ff2d0aa1b95bfb57eb461 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 10:56:09 -0700 Subject: [PATCH 08/38] WIP with integration header Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 309 +++++++++++++++++------------------- 1 file changed, 146 insertions(+), 163 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 61464d292965e..c40b2ed413db7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1007,64 +1007,6 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue()); } -// Creates list of kernel parameters descriptors using KernelObj (kernel object) -// Fields of kernel object must be initialized with SYCL kernel arguments so -// in the following function we extract types of kernel object fields and add it -// to the array with kernel parameters descriptors. -// Returns true if all arguments are successfully built. -/*static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj, - SmallVectorImpl &ParamDescs) { - // Create parameter descriptor for accessor in case when it's wrapped with - // some class. - // TODO: Do we need support case when sampler is wrapped with some class or - // struct? - std::function - createParamDescForWrappedAccessors = - [&](const FieldDecl *Fld, const QualType &ArgTy) { - const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); - for (const auto *WrapperFld : Wrapper->fields()) { - QualType FldType = WrapperFld->getType(); - if (FldType->isStructureOrClassType()) { - if (Util::isSyclAccessorType(FldType)) { - // Accessor field is found - create descriptor. - createSpecialSYCLObjParamDesc(WrapperFld, FldType); - } else if (Util::isSyclSpecConstantType(FldType)) { - // Don't try recursive search below. - } else { - // Field is some class or struct - recursively check for - // accessor fields. - createParamDescForWrappedAccessors(WrapperFld, FldType); - } - } - } - }; - - bool AllArgsAreValid = true; - // Run through kernel object fields and create corresponding kernel - // parameters descriptors. There are a several possible cases: - // - Kernel object field is a SYCL special object (SYCL accessor or SYCL - // sampler). These objects has a special initialization scheme - using - // __init method. - // - Kernel object field has a scalar type. In this case we should add - // kernel parameter with the same type. - // - Kernel object field has a structure or class type. Same handling as a - // scalar but we should check if this structure/class contains accessors - // and add parameter decriptor for them properly. - for (const auto *Fld : KernelObj->fields()) { - QualType ArgTy = Fld->getType(); - if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) { - createSpecialSYCLObjParamDesc(Fld, ArgTy); - } else if (ArgTy->isStructureOrClassType()) { - CreateAndAddPrmDsc(Fld, ArgTy); - - // Create descriptors for each accessor field in the class or struct - createParamDescForWrappedAccessors(Fld, ArgTy); - } - - return AllArgsAreValid; -} -*/ - /// Adds necessary data describing given kernel to the integration header. /// \param H the integration header object /// \param Name kernel name @@ -1072,29 +1014,7 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { /// of single_task, parallel_for, etc) /// \param KernelObjTy kernel object type static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, - QualType NameType, CXXRecordDecl *KernelObjTy) { - - ASTContext &Ctx = KernelObjTy->getASTContext(); - const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(KernelObjTy); - const std::string StableName = PredefinedExpr::ComputeName( - Ctx, PredefinedExpr::UniqueStableNameExpr, NameType); - H.startKernel(Name, NameType, StableName, KernelObjTy->getLocation()); - - auto populateHeaderForAccessor = [&](const QualType &ArgTy, uint64_t Offset) { - // The parameter is a SYCL accessor object. - // The Info field of the parameter descriptor for accessor contains - // two template parameters packed into an integer field: - // - target (e.g. global_buffer, constant_buffer, local); - // - dimension of the accessor. - const auto *AccTy = ArgTy->getAsCXXRecordDecl(); - assert(AccTy && "accessor must be of a record type"); - const auto *AccTmplTy = cast(AccTy); - int Dims = static_cast( - AccTmplTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); - int Info = getAccessTarget(AccTmplTy) | (Dims << 11); - H.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, Offset); - }; - + QualType NameType, CXXRecordDecl *KernelObjTy) {/* std::function populateHeaderForWrappedAccessors = [&](const QualType &ArgTy, uint64_t Offset) { @@ -1180,7 +1100,7 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, llvm_unreachable("unsupported kernel parameter type"); } } -} +*/} // The first template argument to the kernel function is used to identify the // kernel itself. @@ -1198,15 +1118,14 @@ static QualType calculateKernelNameType(ASTContext &Ctx, // Gets a name for the kernel caller func, calculated from the first template // argument. static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, - MangleContext &MC) { + MangleContext &MC, bool StableName) { QualType KernelNameType = calculateKernelNameType(S.getASTContext(), KernelCallerFunc); - if (S.getLangOpts().SYCLUnnamedLambda) + if (StableName) return PredefinedExpr::ComputeName(S.getASTContext(), PredefinedExpr::UniqueStableNameType, KernelNameType); - SmallString<256> Result; llvm::raw_svector_ostream Out(Result); @@ -1216,15 +1135,32 @@ static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, // anonymous namespace so these don't get linkage. namespace { +// Implements the 'for-each-visitor' pattern. +#define KF_FOR_EACH(FUNC) \ + (void)std::initializer_list { (handlers.FUNC(Field, FieldTy), 0)... } + +template +static void VisitAccessorWrapperFields(RecordDecl::field_range Fields, + Handlers &... handlers) { + // TODO: Does this need to handle other types to support other things? I + // don't think so, but we'll see. Also want to see if any consumers need to + // handle these 'sub' structs. If so, we likely need to split the + // 'handleStructType' function into two. Do we need to do the same with + // sampler or spec constant? + for (const auto &Field : Fields) { + QualType FieldTy = Field->getType(); + if (Util::isSyclAccessorType(FieldTy)) + KF_FOR_EACH(handleSyclAccessorType); + else if (FieldTy->isStructureOrClassType()) + VisitAccessorWrapperFields(FieldTy->getAsRecordDecl()->fields(), + handlers...); + } +} // A visitor function that dispatches to functions as defined in // SyclKernelFieldHandler for the purposes of kernel generation. template static void VisitRecordFields(RecordDecl::field_range Fields, Handlers &... handlers) { -// Implements the 'for-each-visitor' pattern. -#define KF_FOR_EACH(FUNC) \ - (void)std::initializer_list { (handlers.FUNC(Field, FieldTy), 0)... } - for (const auto &Field : Fields) { QualType FieldTy = Field->getType(); @@ -1234,9 +1170,11 @@ static void VisitRecordFields(RecordDecl::field_range Fields, KF_FOR_EACH(handleSyclSamplerType); else if (Util::isSyclSpecConstantType(FieldTy)) KF_FOR_EACH(handleSyclSpecConstantType); - else if (FieldTy->isStructureOrClassType()) + else if (FieldTy->isStructureOrClassType()) { KF_FOR_EACH(handleStructType); - else if (FieldTy->isReferenceType()) + VisitAccessorWrapperFields(FieldTy->getAsRecordDecl()->fields(), + handlers...); + } else if (FieldTy->isReferenceType()) KF_FOR_EACH(handleReferenceType); else if (FieldTy->isPointerType()) KF_FOR_EACH(handlePointerType); @@ -1247,8 +1185,8 @@ static void VisitRecordFields(RecordDecl::field_range Fields, else KF_FOR_EACH(handleOtherType); } -#undef KF_FOR_EACH } +#undef KF_FOR_EACH // A base type that the SYCL OpenCL Kernel construction task uses to implement // individual tasks. @@ -1322,14 +1260,13 @@ class SyclKernelFieldChecker } }; -// A type that handles the accessor recursion and acts as a base for -// SyclKernelDeclCreator. It doesn't 'own' anything other than the KernelObj -// pointer and functionality required to add a param. -class SyclKernelDeclBase : public SyclKernelFieldHandler { -protected: +// A type to Create and own the FunctionDecl for the kernel. +class SyclKernelDeclCreator + : public SyclKernelFieldHandler { FunctionDecl *KernelObj; - llvm::SmallVectorImpl &ArgTys; - llvm::SmallVectorImpl &Params; + llvm::SmallVector Params; + SyclKernelFieldChecker &ArgChecker; + Sema::ContextRAII FuncContext; void addParam(const FieldDecl *FD, QualType ArgTy) { // Create a new ParmVarDecl based on the new info. @@ -1343,8 +1280,6 @@ class SyclKernelDeclBase : public SyclKernelFieldHandler { NewParam->setIsUsed(); Params.push_back(NewParam); - - ArgTys.push_back(ArgTy); } // All special SYCL objects must have __init method. We extract types for @@ -1353,8 +1288,7 @@ class SyclKernelDeclBase : public SyclKernelFieldHandler { // the kernel body. void handleSpecialType(const FieldDecl *FD, QualType ArgTy) { const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); - assert(RecordDecl && "Special SYCL object must be of a record type"); - + assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method"); @@ -1362,36 +1296,6 @@ class SyclKernelDeclBase : public SyclKernelFieldHandler { addParam(FD, Param->getType().getCanonicalType()); } -public: - SyclKernelDeclBase(Sema &S, FunctionDecl *KernelObj, - llvm::SmallVectorImpl &ArgTys, - llvm::SmallVectorImpl &Params) - : SyclKernelFieldHandler(S), KernelObj(KernelObj), ArgTys(ArgTys), - Params(Params) {} - - void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { - handleSpecialType(FD, ArgTy); - } - - void handleStructType(const FieldDecl *FD, QualType ArgTy) override { - addParam(FD, ArgTy); - - // Create descriptors for each accessor field in the class or struct - const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); - VisitRecordFields(Wrapper->fields(), *this); - } -}; - -// A type to Create and own the FunctionDecl for the kernel. -class SyclKernelDeclCreator : public SyclKernelDeclBase { - // TODO: rather than this, should we consider a 'commit' function that - // finalizes under success only? - SyclKernelFieldChecker &ArgChecker; - Sema::ContextRAII FuncContext; - // Yes, the list of Parameters contains this info, but we use it often enough - // we shouldn't be recreating it constantly. QualTypes are cheap anyway. - llvm::SmallVector ArgTys; - llvm::SmallVector Params; static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD, StringRef Name) { @@ -1422,14 +1326,19 @@ class SyclKernelDeclCreator : public SyclKernelDeclBase { public: SyclKernelDeclCreator(Sema &S, SyclKernelFieldChecker &ArgChecker, StringRef Name, SourceLocation Loc, bool IsInline) - : SyclKernelDeclBase( - S, createKernelDecl(S.getASTContext(), Name, Loc, IsInline), ArgTys, - Params), + : SyclKernelFieldHandler(S), + KernelObj(createKernelDecl(S.getASTContext(), Name, Loc, IsInline)), ArgChecker(ArgChecker), FuncContext(SemaRef, KernelObj) {} ~SyclKernelDeclCreator() { ASTContext &Ctx = SemaRef.getASTContext(); FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); + + SmallVector ArgTys; + std::transform(std::begin(Params), std::end(Params), + std::back_inserter(ArgTys), + [](const ParmVarDecl *PVD) { return PVD->getType(); }); + QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); KernelObj->setType(FuncType); KernelObj->setParams(Params); @@ -1438,6 +1347,10 @@ class SyclKernelDeclCreator : public SyclKernelDeclBase { SemaRef.addSyclDeviceDecl(KernelObj); } + void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { + handleSpecialType(FD, ArgTy); + } + void handleSyclSamplerType(const FieldDecl *FD, QualType ArgTy) final { handleSpecialType(FD, ArgTy); } @@ -1452,6 +1365,7 @@ class SyclKernelDeclCreator : public SyclKernelDeclBase { QualType ModTy = SemaRef.getASTContext().getPointerType(PointeeTy); addParam(FD, ModTy); } + void handleScalarType(const FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy); } @@ -1462,11 +1376,6 @@ class SyclKernelDeclCreator : public SyclKernelDeclBase { // this implementation in the base. void handleStructType(const FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy); - - // Create descriptors for each accessor field in the class or struct - const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); - SyclKernelDeclBase Recurse(SemaRef, KernelObj, ArgTys, Params); - VisitRecordFields(Wrapper->fields(), Recurse); } void setBody(CompoundStmt *KB) { KernelObj->setBody(KB); } @@ -1486,8 +1395,90 @@ class SyclKernelBodyCreator class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { + SYCLIntegrationHeader &Header; + // Required for calculating the field offsets. + const ASTRecordLayout &Layout; + + uint64_t getFieldOffset(const FieldDecl *FD) const { + // TODO: FIX THIS FOR THE RECURSE CASE! + return Layout.getFieldOffset(FD->getFieldIndex()) / 8; + } + + void addParam(const FieldDecl *FD, QualType ArgTy) { + // TODO: fieldOffset is WRONG if this is in a wrapper! + uint64_t Size = + SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); + Header.addParamDesc(SYCLIntegrationHeader::kind_std_layout, + static_cast(Size), + static_cast(getFieldOffset(FD))); + } + public: - SyclKernelIntHeaderCreator(Sema &S) : SyclKernelFieldHandler(S) {} + SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, + const ASTRecordLayout &LambdaLayout, + SourceLocation KernelLoc, QualType NameType, + StringRef Name, StringRef StableName) + : SyclKernelFieldHandler(S), Header(H), Layout(LambdaLayout) { + Header.startKernel(Name, NameType, StableName, KernelLoc); + } + + void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { + // TODO: offset stuff is wrong again in the recursion case!? + const auto *AccTy = + cast(ArgTy->getAsRecordDecl()); + // TODO: Is this the right assert here? or is it exactly 2? + assert(AccTy->getTemplateArgs().size() >= 2 && + "Incorrect template args for Accessor Type"); + int Dims = static_cast( + AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); + int Info = getAccessTarget(AccTy) | (Dims << 11); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, + getFieldOffset(FD)); + } + + void handleSyclSamplerType(const FieldDecl *FD, QualType ArgTy) final { + const auto *SamplerTy = ArgTy->getAsCXXRecordDecl(); + assert(SamplerTy && "Sampler type must be a C++ record type"); + CXXMethodDecl *InitMethod = getMethodByName(SamplerTy, InitMethodName); + assert(InitMethod && "sampler must have __init method"); + + // sampler __init method has only one argument + const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0); + assert(SamplerArg && "sampler __init method must have sampler parameter"); + + addParam(FD, SamplerArg->getType()); + } + + void handleSyclSpecConstantType(const FieldDecl *FD, QualType ArgTy) final { + const TemplateArgumentList &TemplateArgs = + cast(ArgTy->getAsRecordDecl()) + ->getTemplateInstantiationArgs(); + // TODO: Is this the right assert here? or is it exactly 2? + assert(TemplateArgs.size() >= 2 && + "Incorrect template args for Accessor Type"); + // Get specialization constant ID type, which is the second template + // argument. + QualType SpecConstIDTy = + TypeName::getFullyQualifiedType(TemplateArgs.get(1).getAsType(), + SemaRef.getASTContext(), true) + .getCanonicalType(); + const std::string SpecConstName = PredefinedExpr::ComputeName( + SemaRef.getASTContext(), PredefinedExpr::UniqueStableNameType, + SpecConstIDTy); + Header.addSpecConstant(SpecConstName, SpecConstIDTy); + } + + void handlePointerType(const FieldDecl *FD, QualType ArgTy) final { + addParam(FD, ArgTy); + } + void handleStructType(const FieldDecl *FD, QualType ArgTy) final { + addParam(FD, ArgTy); + } + void handleScalarType(const FieldDecl *FD, QualType ArgTy) final { + addParam(FD, ArgTy); + } + + }; } // namespace @@ -1518,14 +1509,25 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, // The first argument to the KernelCallerFunc is the lambda object. CXXRecordDecl *KernelLambda = getKernelObjectType(KernelCallerFunc); assert(KernelLambda && "invalid kernel caller"); - std::string KernelName = constructKernelName(*this, KernelCallerFunc, MC); + + // Calculate both names, since Integration headers need both. + std::string CalculatedName = + constructKernelName(*this, KernelCallerFunc, MC, /*StableName*/ false); + std::string StableName = + constructKernelName(*this, KernelCallerFunc, MC, /*StableName*/ true); + StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName + : CalculatedName); SyclKernelFieldChecker checker(*this); SyclKernelDeclCreator kernel_decl(*this, checker, KernelName, KernelLambda->getLocation(), KernelCallerFunc->isInlined()); SyclKernelBodyCreator kernel_body(*this, kernel_decl); - SyclKernelIntHeaderCreator int_header(*this); + SyclKernelIntHeaderCreator int_header( + *this, getSyclIntegrationHeader(), + Context.getASTRecordLayout(KernelLambda), KernelLambda->getLocation(), + calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, + StableName); ConstructingOpenCLKernel = true; VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, @@ -1533,28 +1535,9 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, ConstructingOpenCLKernel = false; /* - // Build list of kernel arguments - //llvm::SmallVector ParamDescs; - //if (!buildArgTys(getASTContext(), LE, ParamDescs)) - // return; - // TODO Maybe don't emit integration header inside the Sema? ***populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); - //FunctionDecl *OpenCLKernel = - // CreateOpenCLKernelDeclaration(getASTContext(), Name, ParamDescs); - - //ContextRAII FuncContext(*this, OpenCLKernel); - - // Let's copy source location of a functor/lambda to emit nicer diagnostics - //OpenCLKernel->setLocation(LE->getLocation()); - - // If the source function is implicitly inline, the kernel should be marked - // such as well. This allows the kernel to be ODR'd if there are multiple - uses - // in different translation units. - //OpenCLKernel->setImplicitlyInline(KernelCallerFunc->isInlined()); - //ConstructingOpenCLKernel = true; ****CompoundStmt *OpenCLKernelBody = CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); From e52f39c73d1415c7d1a6486bfafe3cfe80f32b8f Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 12:34:28 -0700 Subject: [PATCH 09/38] Fix int header 'kind' calculation. Also, make sure we visit bases. Still need to fix the integration header offset calculation/keeping track Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 74 ++++++++++++++++++++++++++----------- 1 file changed, 52 insertions(+), 22 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c40b2ed413db7..7c7ebe22fc0a0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1151,9 +1151,17 @@ static void VisitAccessorWrapperFields(RecordDecl::field_range Fields, QualType FieldTy = Field->getType(); if (Util::isSyclAccessorType(FieldTy)) KF_FOR_EACH(handleSyclAccessorType); - else if (FieldTy->isStructureOrClassType()) + else if (FieldTy->isStructureOrClassType()) { + CXXRecordDecl *RD = FieldTy->getAsRecordDecl(); + RD->forAllBases( + [](CXXRecordDecl *Base) { + VisitAccessorWrapperFields(Base->fields(), handlers...); + return true; + }, + /*AllowShortCircuit*/ false); VisitAccessorWrapperFields(FieldTy->getAsRecordDecl()->fields(), handlers...); + } } } // A visitor function that dispatches to functions as defined in @@ -1172,6 +1180,19 @@ static void VisitRecordFields(RecordDecl::field_range Fields, KF_FOR_EACH(handleSyclSpecConstantType); else if (FieldTy->isStructureOrClassType()) { KF_FOR_EACH(handleStructType); + + CXXRecordDecl *RD = FieldTy->getAsRecordDecl(); + // Go through the fields of bases as well, the previous implementation + // missed these, so I presume this is going to be fixing a bug. This goes + // through all the bases in a non-guaranteed way, though it skips VBases + // which are otherwise not allowed anyway. + RD->forAllBases( + [](CXXRecordDecl *Base) { + VisitAccessorWrapperFields(Base->fields(), handlers...); + return true; + }, + /*AllowShortCircuit*/ false); + VisitAccessorWrapperFields(FieldTy->getAsRecordDecl()->fields(), handlers...); } else if (FieldTy->isReferenceType()) @@ -1190,12 +1211,13 @@ static void VisitRecordFields(RecordDecl::field_range Fields, // A base type that the SYCL OpenCL Kernel construction task uses to implement // individual tasks. -template class SyclKernelFieldHandler { +template class SyclKernelFieldVisitor { protected: Sema &SemaRef; SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} public: + // Mark these virutal so that we can use override in the implementer classes, // despite virtual dispatch never being used. virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} @@ -1396,38 +1418,47 @@ class SyclKernelBodyCreator class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { SYCLIntegrationHeader &Header; - // Required for calculating the field offsets. - const ASTRecordLayout &Layout; + const CXXRecordDecl *KernelLambda; + + // Keeping track of offsets as we go along is a little awkward, so see if just + // calculating each time is worth doing. Presumably, if the structure depth doesn't + // get insane, we shouldn't have a problem. + uint64_t getFieldOffsetHelper(const CXXRecordDecl *RD, const FieldDecl *FD) { + // TODO! + return 0; + } uint64_t getFieldOffset(const FieldDecl *FD) const { - // TODO: FIX THIS FOR THE RECURSE CASE! - return Layout.getFieldOffset(FD->getFieldIndex()) / 8; + // TODO: Figure out a better way to do this, having to recalculate this + // constantly is going to be expensive. + // TODO: Figure out how to calc lower down the structs. + // uint64_t CurOffset = SemaRef.getASTContext().getFieldOffset(FD) / 8; + return 0; } - void addParam(const FieldDecl *FD, QualType ArgTy) { - // TODO: fieldOffset is WRONG if this is in a wrapper! + void addParam(const FieldDecl *FD, QualType ArgTy, + SYCLIntegrationHeader::kernel_param_kind_t Kind) { uint64_t Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); - Header.addParamDesc(SYCLIntegrationHeader::kind_std_layout, - static_cast(Size), + Header.addParamDesc(Kind, static_cast(Size), static_cast(getFieldOffset(FD))); } public: SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, - const ASTRecordLayout &LambdaLayout, + const CXXRecordDecl *KernelLambda, SourceLocation KernelLoc, QualType NameType, StringRef Name, StringRef StableName) - : SyclKernelFieldHandler(S), Header(H), Layout(LambdaLayout) { - Header.startKernel(Name, NameType, StableName, KernelLoc); + : SyclKernelFieldHandler(S), Header(H), + KernelLambda(KernelLambda) { + Header.startKernel(Name, NameType, StableName, KernelLambda->GetLocation()); } void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { // TODO: offset stuff is wrong again in the recursion case!? const auto *AccTy = cast(ArgTy->getAsRecordDecl()); - // TODO: Is this the right assert here? or is it exactly 2? - assert(AccTy->getTemplateArgs().size() >= 2 && + assert(AccTy->getTemplateArgs().size() == 2 && "Incorrect template args for Accessor Type"); int Dims = static_cast( AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); @@ -1446,15 +1477,14 @@ class SyclKernelIntHeaderCreator const ParmVarDecl *SamplerArg = InitMethod->getParamDecl(0); assert(SamplerArg && "sampler __init method must have sampler parameter"); - addParam(FD, SamplerArg->getType()); + addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler); } void handleSyclSpecConstantType(const FieldDecl *FD, QualType ArgTy) final { const TemplateArgumentList &TemplateArgs = cast(ArgTy->getAsRecordDecl()) ->getTemplateInstantiationArgs(); - // TODO: Is this the right assert here? or is it exactly 2? - assert(TemplateArgs.size() >= 2 && + assert(TemplateArgs.size() == 2 && "Incorrect template args for Accessor Type"); // Get specialization constant ID type, which is the second template // argument. @@ -1469,13 +1499,13 @@ class SyclKernelIntHeaderCreator } void handlePointerType(const FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy); + addParam(FD, ArgTy, SYCLIntegrationHeader::kind_pointer); } void handleStructType(const FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy); + addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); } void handleScalarType(const FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy); + addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); } @@ -1525,7 +1555,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, SyclKernelBodyCreator kernel_body(*this, kernel_decl); SyclKernelIntHeaderCreator int_header( *this, getSyclIntegrationHeader(), - Context.getASTRecordLayout(KernelLambda), KernelLambda->getLocation(), + KernelLambda, Context.getASTRecordLayout(KernelLambda), KernelLambda->getLocation(), calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, StableName); From bf828fc962ad37399b994ae38e29504f52290b8b Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 14:07:53 -0700 Subject: [PATCH 10/38] small refactoring for base classes --- clang/lib/Sema/SemaSYCL.cpp | 267 +++++++++++++++--------------------- 1 file changed, 110 insertions(+), 157 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7c7ebe22fc0a0..12ecdf2f2e3f0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1001,107 +1001,22 @@ static ParamDesc makeParamDesc(const FieldDecl *Src, QualType Ty) { Ctx.getTrivialTypeSourceInfo(Ty)); } +static ParamDesc makeParamDesc(ASTContext &Ctx, const CXXBaseSpecifier &Src, + QualType Ty) { + // TODO: There is no name for the base available, but duplicate names are + // seemingly already possible, so we'll give them all the same name for now. + // This only happens with the accessor types. + std::string Name = "_arg__base"; + return std::make_tuple(Ty, &Ctx.Idents.get(Name), + Ctx.getTrivialTypeSourceInfo(Ty)); +} + /// \return the target of given SYCL accessor type static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { return static_cast( AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue()); } -/// Adds necessary data describing given kernel to the integration header. -/// \param H the integration header object -/// \param Name kernel name -/// \param NameType type representing kernel name (first template argument -/// of single_task, parallel_for, etc) -/// \param KernelObjTy kernel object type -static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, - QualType NameType, CXXRecordDecl *KernelObjTy) {/* - std::function - populateHeaderForWrappedAccessors = [&](const QualType &ArgTy, - uint64_t Offset) { - const auto *Wrapper = ArgTy->getAsCXXRecordDecl(); - for (const auto *WrapperFld : Wrapper->fields()) { - QualType FldType = WrapperFld->getType(); - if (FldType->isStructureOrClassType()) { - ASTContext &WrapperCtx = Wrapper->getASTContext(); - const ASTRecordLayout &WrapperLayout = - WrapperCtx.getASTRecordLayout(Wrapper); - // Get offset (in bytes) of the field in wrapper class or struct - uint64_t OffsetInWrapper = - WrapperLayout.getFieldOffset(WrapperFld->getFieldIndex()) / 8; - if (Util::isSyclAccessorType(FldType)) { - // This is an accesor - populate the header appropriately - populateHeaderForAccessor(FldType, Offset + OffsetInWrapper); - } else { - // This is an other class or struct - recursively search for an - // accessor field - populateHeaderForWrappedAccessors(FldType, - Offset + OffsetInWrapper); - } - } - } - }; - - for (const auto Fld : KernelObjTy->fields()) { - QualType ActualArgType; - QualType ArgTy = Fld->getType(); - - // Get offset in bytes - uint64_t Offset = Layout.getFieldOffset(Fld->getFieldIndex()) / 8; - - if (Util::isSyclAccessorType(ArgTy)) { - populateHeaderForAccessor(ArgTy, Offset); - } else if (Util::isSyclSamplerType(ArgTy)) { - // The parameter is a SYCL sampler object - const auto *SamplerTy = ArgTy->getAsCXXRecordDecl(); - assert(SamplerTy && "sampler must be of a record type"); - - CXXMethodDecl *InitMethod = getMethodByName(SamplerTy, InitMethodName); - assert(InitMethod && "sampler must have __init method"); - - // sampler __init method has only one argument - auto *FuncDecl = cast(InitMethod); - ParmVarDecl *SamplerArg = FuncDecl->getParamDecl(0); - assert(SamplerArg && "sampler __init method must have sampler parameter"); - uint64_t Sz = Ctx.getTypeSizeInChars(SamplerArg->getType()).getQuantity(); - H.addParamDesc(SYCLIntegrationHeader::kind_sampler, - static_cast(Sz), static_cast(Offset)); - } else if (ArgTy->isPointerType()) { - uint64_t Sz = Ctx.getTypeSizeInChars(Fld->getType()).getQuantity(); - H.addParamDesc(SYCLIntegrationHeader::kind_pointer, - static_cast(Sz), static_cast(Offset)); - } else if (Util::isSyclSpecConstantType(ArgTy)) { - // Add specialization constant ID to the header. - auto *TmplSpec = - cast(ArgTy->getAsCXXRecordDecl()); - const TemplateArgumentList *TemplateArgs = - &TmplSpec->getTemplateInstantiationArgs(); - // Get specialization constant ID type, which is the second template - // argument. - QualType SpecConstIDTy = TypeName::getFullyQualifiedType( - TemplateArgs->get(1).getAsType(), Ctx, true) - .getCanonicalType(); - const std::string SpecConstName = PredefinedExpr::ComputeName( - Ctx, PredefinedExpr::UniqueStableNameExpr, SpecConstIDTy); - H.addSpecConstant(SpecConstName, SpecConstIDTy); - // Spec constant lambda capture does not become a kernel argument. - } else if (ArgTy->isStructureOrClassType() || ArgTy->isScalarType()) { - // the parameter is an object of standard layout type or scalar; - // the check for standard layout is done elsewhere - uint64_t Sz = Ctx.getTypeSizeInChars(Fld->getType()).getQuantity(); - H.addParamDesc(SYCLIntegrationHeader::kind_std_layout, - static_cast(Sz), static_cast(Offset)); - - // check for accessor fields in structure or class and populate the - // integration header appropriately - if (ArgTy->isStructureOrClassType()) { - populateHeaderForWrappedAccessors(ArgTy, Offset); - } - } else { - llvm_unreachable("unsupported kernel parameter type"); - } - } -*/} - // The first template argument to the kernel function is used to identify the // kernel itself. static QualType calculateKernelNameType(ASTContext &Ctx, @@ -1136,39 +1051,46 @@ static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, // anonymous namespace so these don't get linkage. namespace { // Implements the 'for-each-visitor' pattern. -#define KF_FOR_EACH(FUNC) \ - (void)std::initializer_list { (handlers.FUNC(Field, FieldTy), 0)... } - template -static void VisitAccessorWrapperFields(RecordDecl::field_range Fields, - Handlers &... handlers) { - // TODO: Does this need to handle other types to support other things? I - // don't think so, but we'll see. Also want to see if any consumers need to - // handle these 'sub' structs. If so, we likely need to split the - // 'handleStructType' function into two. Do we need to do the same with - // sampler or spec constant? - for (const auto &Field : Fields) { - QualType FieldTy = Field->getType(); - if (Util::isSyclAccessorType(FieldTy)) - KF_FOR_EACH(handleSyclAccessorType); - else if (FieldTy->isStructureOrClassType()) { - CXXRecordDecl *RD = FieldTy->getAsRecordDecl(); - RD->forAllBases( - [](CXXRecordDecl *Base) { - VisitAccessorWrapperFields(Base->fields(), handlers...); - return true; - }, - /*AllowShortCircuit*/ false); - VisitAccessorWrapperFields(FieldTy->getAsRecordDecl()->fields(), - handlers...); +static void VisitAccessorWrapper(CXXRecordDecl *Wrapper, + Handlers &... handlers); + +QualType getItemType(const FieldDecl *FD) { + return FD->getType(); +} +QualType getItemType(const CXXBaseSpecifier &BS) { + return BS.getType(); +} + +template +static void VisitAccessorWrapperHelper(RangeTy Range, Handlers &... handlers) { + for (const auto &Item : Range) { + QualType ItemTy = getItemType(Item); + if (Util::isSyclAccessorType(ItemTy)) + (void)std::initializer_list{ + (handlers.handleAccessorType(Item, ItemTy), 0)...}; + else if (ItemTy->isStructureOrClassType()) { + VisitAccessorWrapper(ItemTy->getAsCXXRecordDecl()); } } } + +template +static void VisitAccessorWrapper(CXXRecordDecl *Wrapper, + Handlers &... handlers) { + + VisitAccessorWrapperHelper(Wrapper->bases(), handlers...); + VisitAccessorWrapperHelper(Wrapper->fields(), handlers...); +} + // A visitor function that dispatches to functions as defined in // SyclKernelFieldHandler for the purposes of kernel generation. template static void VisitRecordFields(RecordDecl::field_range Fields, Handlers &... handlers) { +#define KF_FOR_EACH(FUNC) \ + (void)std::initializer_list { (handlers.FUNC(Field, FieldTy), 0)... } + for (const auto &Field : Fields) { QualType FieldTy = Field->getType(); @@ -1180,21 +1102,8 @@ static void VisitRecordFields(RecordDecl::field_range Fields, KF_FOR_EACH(handleSyclSpecConstantType); else if (FieldTy->isStructureOrClassType()) { KF_FOR_EACH(handleStructType); - - CXXRecordDecl *RD = FieldTy->getAsRecordDecl(); - // Go through the fields of bases as well, the previous implementation - // missed these, so I presume this is going to be fixing a bug. This goes - // through all the bases in a non-guaranteed way, though it skips VBases - // which are otherwise not allowed anyway. - RD->forAllBases( - [](CXXRecordDecl *Base) { - VisitAccessorWrapperFields(Base->fields(), handlers...); - return true; - }, - /*AllowShortCircuit*/ false); - - VisitAccessorWrapperFields(FieldTy->getAsRecordDecl()->fields(), - handlers...); + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); + VisitAccessorWrapper(RD); } else if (FieldTy->isReferenceType()) KF_FOR_EACH(handleReferenceType); else if (FieldTy->isPointerType()) @@ -1206,20 +1115,26 @@ static void VisitRecordFields(RecordDecl::field_range Fields, else KF_FOR_EACH(handleOtherType); } -} #undef KF_FOR_EACH +} // A base type that the SYCL OpenCL Kernel construction task uses to implement // individual tasks. -template class SyclKernelFieldVisitor { +template class SyclKernelFieldHandler { protected: Sema &SemaRef; SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} public: - // Mark these virutal so that we can use override in the implementer classes, // despite virtual dispatch never being used. + + //// TODO: Can these return 'bool' and we can short-circuit the handling? That + // way the field checker cna return true/false based on whether the rest + // should be still working. + + // Accessor can be a base class or a field decl, so both must be handled. + virtual void handleSyclAccessorType(const CXXBaseSpecifier &, QualType) {} virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} virtual void handleSyclSamplerType(const FieldDecl *, QualType) {} virtual void handleSyclSpecConstantType(const FieldDecl *, QualType) {} @@ -1291,8 +1206,17 @@ class SyclKernelDeclCreator Sema::ContextRAII FuncContext; void addParam(const FieldDecl *FD, QualType ArgTy) { - // Create a new ParmVarDecl based on the new info. ParamDesc newParamDesc = makeParamDesc(FD, ArgTy); + addParam(newParamDesc, ArgTy); + } + + void addParam(const CXXBaseSpecifier &BS, QualType ArgTy) { + ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), BS, ArgTy); + addParam(newParamDesc, ArgTy); + } + + void addParam(ParamDesc newParamDesc, QualType ArgTy) { + // Create a new ParmVarDecl based on the new info. auto *NewParam = ParmVarDecl::Create( SemaRef.getASTContext(), KernelObj, SourceLocation(), SourceLocation(), std::get<1>(newParamDesc), std::get<0>(newParamDesc), @@ -1369,6 +1293,17 @@ class SyclKernelDeclCreator SemaRef.addSyclDeviceDecl(KernelObj); } + void handleSyclAccessorType(const CXXBaseSpecifier &BS, + QualType ArgTy) final { + const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); + assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); + CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); + assert(InitMethod && "The accessor/sampler must have the __init method"); + + for (const ParmVarDecl *Param : InitMethod->parameters()) + addParam(BS, Param->getType().getCanonicalType()); + } + void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { handleSpecialType(FD, ArgTy); } @@ -1419,52 +1354,71 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { SYCLIntegrationHeader &Header; const CXXRecordDecl *KernelLambda; + int64_t CurOffset = 0; // Keeping track of offsets as we go along is a little awkward, so see if just // calculating each time is worth doing. Presumably, if the structure depth doesn't - // get insane, we shouldn't have a problem. - uint64_t getFieldOffsetHelper(const CXXRecordDecl *RD, const FieldDecl *FD) { + // get insane, we shouldn't have a problem with run time. + uint64_t getFieldOffset(const CXXRecordDecl *RD, const FieldDecl *FD) { + // TODO! + return 0; + } + uint64_t getBaseOffset(const CXXRecordDecl *RD, const FieldDecl *FD) { // TODO! return 0; } - uint64_t getFieldOffset(const FieldDecl *FD) const { - // TODO: Figure out a better way to do this, having to recalculate this - // constantly is going to be expensive. - // TODO: Figure out how to calc lower down the structs. - // uint64_t CurOffset = SemaRef.getASTContext().getFieldOffset(FD) / 8; + uint64_t getOffset(const CXXRecordDecl *RD) const { + // TODO: Figure this out! Offset of a base class. return 0; } + uint64_t getOffset(const FieldDecl *FD) const { + // TODO: Figure out how to calc lower down the structs, currently only gives + // the 'base' value. + return SemaRef.getASTContext().getFieldOffset(FD) / 8; + } void addParam(const FieldDecl *FD, QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind) { uint64_t Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); Header.addParamDesc(Kind, static_cast(Size), - static_cast(getFieldOffset(FD))); + static_cast(getOffset(FD))); } public: SyclKernelIntHeaderCreator(Sema &S, SYCLIntegrationHeader &H, const CXXRecordDecl *KernelLambda, - SourceLocation KernelLoc, QualType NameType, - StringRef Name, StringRef StableName) - : SyclKernelFieldHandler(S), Header(H), - KernelLambda(KernelLambda) { - Header.startKernel(Name, NameType, StableName, KernelLambda->GetLocation()); + QualType NameType, StringRef Name, + StringRef StableName) + : SyclKernelFieldHandler(S), Header(H), KernelLambda(KernelLambda) { + Header.startKernel(Name, NameType, StableName, KernelLambda->getLocation()); + } + + void handleSyclAccessorType(const CXXBaseSpecifier &BC, + QualType ArgTy) final { + const auto *AccTy = + cast(ArgTy->getAsRecordDecl()); + assert(AccTy->getTemplateArgs().size() >= 2 && + "Incorrect template args for Accessor Type"); + int Dims = static_cast( + AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); + int Info = getAccessTarget(AccTy) | (Dims << 11); + Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, + // TODO: is this the right way? + getOffset(BC.getType()->getAsCXXRecordDecl())); } void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { - // TODO: offset stuff is wrong again in the recursion case!? const auto *AccTy = cast(ArgTy->getAsRecordDecl()); - assert(AccTy->getTemplateArgs().size() == 2 && + assert(AccTy->getTemplateArgs().size() >= 2 && "Incorrect template args for Accessor Type"); int Dims = static_cast( AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); int Info = getAccessTarget(AccTy) | (Dims << 11); Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, - getFieldOffset(FD)); + getOffset(FD)); } void handleSyclSamplerType(const FieldDecl *FD, QualType ArgTy) final { @@ -1555,9 +1509,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, SyclKernelBodyCreator kernel_body(*this, kernel_decl); SyclKernelIntHeaderCreator int_header( *this, getSyclIntegrationHeader(), - KernelLambda, Context.getASTRecordLayout(KernelLambda), KernelLambda->getLocation(), - calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, - StableName); + KernelLambda, calculateKernelNameType(Context, KernelCallerFunc), + CalculatedName, StableName); ConstructingOpenCLKernel = true; VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, From 5baa5ad654c80ee8d2e5c794e00cc859a2148e5d Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 14:49:01 -0700 Subject: [PATCH 11/38] First run at getting the offset correct... not sure if this works --- clang/lib/Sema/SemaSYCL.cpp | 82 +++++++++++++++++++++++++------------ 1 file changed, 55 insertions(+), 27 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 12ecdf2f2e3f0..111c33f92ef3d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1050,10 +1050,6 @@ static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, // anonymous namespace so these don't get linkage. namespace { -// Implements the 'for-each-visitor' pattern. -template -static void VisitAccessorWrapper(CXXRecordDecl *Wrapper, - Handlers &... handlers); QualType getItemType(const FieldDecl *FD) { return FD->getType(); @@ -1062,25 +1058,39 @@ QualType getItemType(const CXXBaseSpecifier &BS) { return BS.getType(); } +// Implements the 'for-each-visitor' pattern. +template +static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent, + CXXRecordDecl *Wrapper, + Handlers &... handlers); + template -static void VisitAccessorWrapperHelper(RangeTy Range, Handlers &... handlers) { +static void VisitAccessorWrapperHelper(CXXRecordDecl *Owner, RangeTy Range, + Handlers &... handlers) { for (const auto &Item : Range) { QualType ItemTy = getItemType(Item); if (Util::isSyclAccessorType(ItemTy)) (void)std::initializer_list{ - (handlers.handleAccessorType(Item, ItemTy), 0)...}; + (handlers.handleSyclAccessorType(Item, ItemTy), 0)...}; else if (ItemTy->isStructureOrClassType()) { - VisitAccessorWrapper(ItemTy->getAsCXXRecordDecl()); + VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(), + handlers...); } } } -template -static void VisitAccessorWrapper(CXXRecordDecl *Wrapper, - Handlers &... handlers) { - - VisitAccessorWrapperHelper(Wrapper->bases(), handlers...); - VisitAccessorWrapperHelper(Wrapper->fields(), handlers...); +// poorly named Parent is the 'how we got here', basically just enough info for +// the offset adjustment to know what to do about the enter-struct info. +template +static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent, + CXXRecordDecl *Wrapper, + Handlers &... handlers) { + (void)std::initializer_list{ + (handlers.enterStruct(Owner, Parent), 0)...}; + VisitAccessorWrapperHelper(Wrapper, Wrapper->bases(), handlers...); + VisitAccessorWrapperHelper(Wrapper, Wrapper->fields(), handlers...); + (void)std::initializer_list{ + (handlers.leaveStruct(Owner, Parent), 0)...}; } // A visitor function that dispatches to functions as defined in @@ -1103,7 +1113,7 @@ static void VisitRecordFields(RecordDecl::field_range Fields, else if (FieldTy->isStructureOrClassType()) { KF_FOR_EACH(handleStructType); CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); - VisitAccessorWrapper(RD); + VisitAccessorWrapper(nullptr, Field, RD, handlers...); } else if (FieldTy->isReferenceType()) KF_FOR_EACH(handleReferenceType); else if (FieldTy->isPointerType()) @@ -1145,6 +1155,17 @@ template class SyclKernelFieldHandler { virtual void handleScalarType(const FieldDecl *, QualType) {} // Most handlers shouldn't be handling this, just the field checker. virtual void handleOtherType(const FieldDecl *, QualType) {} + + // The following are only used for keeping track of where we are in the base + // class/field graph. Int Headers use this to calculate offset, most others + // don't have a need for these. + + virtual void enterStruct(const CXXRecordDecl *, const FieldDecl *) {} + virtual void leaveStruct(const CXXRecordDecl *, const FieldDecl *) {} + virtual void enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} + virtual void leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} + // virtual void enterStruct(const FieldDecl *, CXXRecordDecl *Struct); + // virtual void leaveStruct(const FieldDecl *, CXXRecordDecl *Struct); }; // A type to check the valididty of all of the argument types. @@ -1356,18 +1377,6 @@ class SyclKernelIntHeaderCreator const CXXRecordDecl *KernelLambda; int64_t CurOffset = 0; - // Keeping track of offsets as we go along is a little awkward, so see if just - // calculating each time is worth doing. Presumably, if the structure depth doesn't - // get insane, we shouldn't have a problem with run time. - uint64_t getFieldOffset(const CXXRecordDecl *RD, const FieldDecl *FD) { - // TODO! - return 0; - } - uint64_t getBaseOffset(const CXXRecordDecl *RD, const FieldDecl *FD) { - // TODO! - return 0; - } - uint64_t getOffset(const CXXRecordDecl *RD) const { // TODO: Figure this out! Offset of a base class. return 0; @@ -1375,7 +1384,7 @@ class SyclKernelIntHeaderCreator uint64_t getOffset(const FieldDecl *FD) const { // TODO: Figure out how to calc lower down the structs, currently only gives // the 'base' value. - return SemaRef.getASTContext().getFieldOffset(FD) / 8; + return CurOffset + SemaRef.getASTContext().getFieldOffset(FD) / 8; } void addParam(const FieldDecl *FD, QualType ArgTy, @@ -1462,7 +1471,26 @@ class SyclKernelIntHeaderCreator addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); } + // Keep track of the current struct offset. + void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; + } + + void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + CurOffset -= SemaRef.getASTContext().getFieldOffset(FD) / 8; + } + void enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + const ASTRecordLayout &Layout = + SemaRef.getASTContext().getASTRecordLayout(RD); + CurOffset += Layout.getBaseClassOffset(BS.getType()); + } + + void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + const ASTRecordLayout &Layout = + SemaRef.getASTContext().getASTRecordLayout(RD); + CurOffset -= Layout.getBaseClassOffset(BS.getType()); + } }; } // namespace From c551aad0a0935333488abd78f69f54774d0d8026 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 14:56:51 -0700 Subject: [PATCH 12/38] Got the offset logic compiling/looking right. Also did a clang format. Quitting for the weekend, pushing so that others can have a look at the progress/help how they can --- clang/lib/Sema/SemaSYCL.cpp | 23 +++++++++-------------- 1 file changed, 9 insertions(+), 14 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 111c33f92ef3d..a318d879d504c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1051,12 +1051,8 @@ static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, // anonymous namespace so these don't get linkage. namespace { -QualType getItemType(const FieldDecl *FD) { - return FD->getType(); -} -QualType getItemType(const CXXBaseSpecifier &BS) { - return BS.getType(); -} +QualType getItemType(const FieldDecl *FD) { return FD->getType(); } +QualType getItemType(const CXXBaseSpecifier &BS) { return BS.getType(); } // Implements the 'for-each-visitor' pattern. template @@ -1085,12 +1081,10 @@ template static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent, CXXRecordDecl *Wrapper, Handlers &... handlers) { - (void)std::initializer_list{ - (handlers.enterStruct(Owner, Parent), 0)...}; + (void)std::initializer_list{(handlers.enterStruct(Owner, Parent), 0)...}; VisitAccessorWrapperHelper(Wrapper, Wrapper->bases(), handlers...); VisitAccessorWrapperHelper(Wrapper, Wrapper->fields(), handlers...); - (void)std::initializer_list{ - (handlers.leaveStruct(Owner, Parent), 0)...}; + (void)std::initializer_list{(handlers.leaveStruct(Owner, Parent), 0)...}; } // A visitor function that dispatches to functions as defined in @@ -1263,7 +1257,6 @@ class SyclKernelDeclCreator addParam(FD, Param->getType().getCanonicalType()); } - static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD, StringRef Name) { // Set implict attributes. @@ -1483,14 +1476,16 @@ class SyclKernelIntHeaderCreator void enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { const ASTRecordLayout &Layout = SemaRef.getASTContext().getASTRecordLayout(RD); - CurOffset += Layout.getBaseClassOffset(BS.getType()); + CurOffset += Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl()) + .getQuantity(); } void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { const ASTRecordLayout &Layout = SemaRef.getASTContext().getASTRecordLayout(RD); - CurOffset -= Layout.getBaseClassOffset(BS.getType()); - } + CurOffset -= Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl()) + .getQuantity(); + } }; } // namespace From 6a5609c6a6af2005c789741cdb8daa463389c751 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Fri, 10 Apr 2020 15:07:32 -0700 Subject: [PATCH 13/38] Format the createOpenCLKernel func, remove int header from comment --- clang/lib/Sema/SemaSYCL.cpp | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a318d879d504c..97fccbbb4d4bc 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1531,9 +1531,9 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, KernelCallerFunc->isInlined()); SyclKernelBodyCreator kernel_body(*this, kernel_decl); SyclKernelIntHeaderCreator int_header( - *this, getSyclIntegrationHeader(), - KernelLambda, calculateKernelNameType(Context, KernelCallerFunc), - CalculatedName, StableName); + *this, getSyclIntegrationHeader(), KernelLambda, + calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, + StableName); ConstructingOpenCLKernel = true; VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, @@ -1541,9 +1541,6 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, ConstructingOpenCLKernel = false; /* - // TODO Maybe don't emit integration header inside the Sema? - ***populateIntHeader(getSyclIntegrationHeader(), Name, KernelNameType, LE); - //ConstructingOpenCLKernel = true; ****CompoundStmt *OpenCLKernelBody = CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); From 542a078c1b5ad5886b808f78854464db9ef800aa Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 13 Apr 2020 05:58:39 -0700 Subject: [PATCH 14/38] Rename kernelobj to kerneldecl Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 97fccbbb4d4bc..94bdbdc1fef73 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1215,7 +1215,7 @@ class SyclKernelFieldChecker // A type to Create and own the FunctionDecl for the kernel. class SyclKernelDeclCreator : public SyclKernelFieldHandler { - FunctionDecl *KernelObj; + FunctionDecl *KernelDecl; llvm::SmallVector Params; SyclKernelFieldChecker &ArgChecker; Sema::ContextRAII FuncContext; @@ -1233,7 +1233,7 @@ class SyclKernelDeclCreator void addParam(ParamDesc newParamDesc, QualType ArgTy) { // Create a new ParmVarDecl based on the new info. auto *NewParam = ParmVarDecl::Create( - SemaRef.getASTContext(), KernelObj, SourceLocation(), SourceLocation(), + SemaRef.getASTContext(), KernelDecl, SourceLocation(), SourceLocation(), std::get<1>(newParamDesc), std::get<0>(newParamDesc), std::get<2>(newParamDesc), SC_None, /*DefArg*/ nullptr); @@ -1287,8 +1287,8 @@ class SyclKernelDeclCreator SyclKernelDeclCreator(Sema &S, SyclKernelFieldChecker &ArgChecker, StringRef Name, SourceLocation Loc, bool IsInline) : SyclKernelFieldHandler(S), - KernelObj(createKernelDecl(S.getASTContext(), Name, Loc, IsInline)), - ArgChecker(ArgChecker), FuncContext(SemaRef, KernelObj) {} + KernelDecl(createKernelDecl(S.getASTContext(), Name, Loc, IsInline)), + ArgChecker(ArgChecker), FuncContext(SemaRef, KernelDecl) {} ~SyclKernelDeclCreator() { ASTContext &Ctx = SemaRef.getASTContext(); @@ -1300,11 +1300,11 @@ class SyclKernelDeclCreator [](const ParmVarDecl *PVD) { return PVD->getType(); }); QualType FuncType = Ctx.getFunctionType(Ctx.VoidTy, ArgTys, Info); - KernelObj->setType(FuncType); - KernelObj->setParams(Params); + KernelDecl->setType(FuncType); + KernelDecl->setParams(Params); if (ArgChecker.isValid()) - SemaRef.addSyclDeviceDecl(KernelObj); + SemaRef.addSyclDeviceDecl(KernelDecl); } void handleSyclAccessorType(const CXXBaseSpecifier &BS, @@ -1349,7 +1349,7 @@ class SyclKernelDeclCreator addParam(FD, ArgTy); } - void setBody(CompoundStmt *KB) { KernelObj->setBody(KB); } + void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } }; class SyclKernelBodyCreator From 7d3bb74e0eac4ba38137743ab483a6e27117dd23 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 13 Apr 2020 08:28:30 -0700 Subject: [PATCH 15/38] Framwork for body, added the stream type as a special type Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 73 ++++++++++++++++++++++++++++++++++--- 1 file changed, 67 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 94bdbdc1fef73..8a5631bf69f3c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -969,6 +969,7 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S, // need to replace all refs to this kernel oject with refs to our clone // declared inside kernel body. Stmt *FunctionBody = KernelCallerFunc->getBody(); + ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); // DeclRefExpr with valid source location but with decl which is not marked @@ -1068,6 +1069,9 @@ static void VisitAccessorWrapperHelper(CXXRecordDecl *Owner, RangeTy Range, if (Util::isSyclAccessorType(ItemTy)) (void)std::initializer_list{ (handlers.handleSyclAccessorType(Item, ItemTy), 0)...}; + else if (Util::isSyclStreamType(ItemTy)) + (void)std::initializer_list{ + (handlers.handleSyclStreamType(Item, ItemTy), 0)...}; else if (ItemTy->isStructureOrClassType()) { VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(), handlers...); @@ -1104,6 +1108,8 @@ static void VisitRecordFields(RecordDecl::field_range Fields, KF_FOR_EACH(handleSyclSamplerType); else if (Util::isSyclSpecConstantType(FieldTy)) KF_FOR_EACH(handleSyclSpecConstantType); + else if (Util::isSyclStreamType(FieldTy)) + KF_FOR_EACH(handleSyclStreamType); else if (FieldTy->isStructureOrClassType()) { KF_FOR_EACH(handleStructType); CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); @@ -1142,6 +1148,7 @@ template class SyclKernelFieldHandler { virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} virtual void handleSyclSamplerType(const FieldDecl *, QualType) {} virtual void handleSyclSpecConstantType(const FieldDecl *, QualType) {} + virtual void handleSyclStreamType(const FieldDecl *, QualType) {} virtual void handleStructType(const FieldDecl *, QualType) {} virtual void handleReferenceType(const FieldDecl *, QualType) {} virtual void handlePointerType(const FieldDecl *, QualType) {} @@ -1350,18 +1357,71 @@ class SyclKernelDeclCreator } void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } + + FunctionDecl *getKernelDecl() { return KernelDecl; } }; class SyclKernelBodyCreator : public SyclKernelFieldHandler { SyclKernelDeclCreator &DeclCreator; - // TODO: When/Where does this get created? - CompoundStmt *KernelBody = nullptr; + llvm::SmallVector BodyStmts; + llvm::SmallVector FinalizeStmts; + llvm::SmallVector InitExprs; + + // Using the statements/init expressions that we've created, this generates + // the kernel body compound stmt. CompoundStmt needs to know its number of + // statements in advance to allocate it, so we cannot do this as we go along. + CompountStmt *createKernelBody() { + // TODO: Can we hold off on creating KernelObjClone to here? + + Expr *ILE = new (SemaRef.getASTContext()) + InitListExpr(S.Context, SourceLocation(), InitExprs, SourceLocation()); + // TODO!!! ILE->setType(QualType(LC->getTypeForDecl(), 0)); + // KernelObjectClone->setInit(ILE); + + // TODO: More kernel object init with KernelBodyTransform. + + BodyStmts.insert(std::end(BodyStmts), std::begin(FinalizeStmts), + std::begin(FinalizeStmts)); + return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {}); + } + + // TODO: not sure what this does yet, name is a placeholder for future use. + void doSomethingForParallelForWorkGroup() { + } public: - SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC) - : SyclKernelFieldHandler(S), DeclCreator(DC) {} - ~SyclKernelBodyCreator() { DeclCreator.setBody(KernelBody); } + SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, KernelInvocationKind K) + : SyclKernelFieldHandler(S), DeclCreator(DC) { + // TODO: Something special with the lambda when InvokeParallelForWorkGroup. + if (K == InvokeParallelForWorkGroup) + do somethingForparalellForWorkGroup(); + } + ~SyclKernelBodyCreator() { + CompoundStmt *KernelBody = createKernelBody(); + DeclCreator.setBody(KernelBody); + } + + void handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { + // TODO: Creates init sequence and inits special sycl obj + } + + void handleSyclSamplerType(FieldDecl *FD, QualType Ty) final { + // TODO: Creates init sequence and inits special sycl obj + } + + void handleSyclStreamType(FieldDecl *FD, QualType Ty) final { + // TODO: Creates init/finalize sequence and inits special sycl obj + } + + void handleStructType(FieldDecl *FD, QualType Ty) final { + // TODO: a bunch of work doing inits, note this has a little more than + // scalar. + } + void handleScalarType(FieldDecl *FD, QualType Ty) final { + // TODO: a bunch of work doing inits. + } + }; class SyclKernelIntHeaderCreator @@ -1529,7 +1589,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, SyclKernelDeclCreator kernel_decl(*this, checker, KernelName, KernelLambda->getLocation(), KernelCallerFunc->isInlined()); - SyclKernelBodyCreator kernel_body(*this, kernel_decl); + SyclKernelBodyCreator kernel_body(*this, kernel_decl, + getKernelInvocationKind(KernelCallerFunc)); SyclKernelIntHeaderCreator int_header( *this, getSyclIntegrationHeader(), KernelLambda, calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, From 9f582cdfcc8b62beb25bf9eddb41d4f7c2a3b434 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 13 Apr 2020 08:36:09 -0700 Subject: [PATCH 16/38] Fix all the build errors from the last patch Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 39 ++++++++++++++++++++++++------------- 1 file changed, 25 insertions(+), 14 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8a5631bf69f3c..d8a04d6d85d9e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1148,6 +1148,7 @@ template class SyclKernelFieldHandler { virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} virtual void handleSyclSamplerType(const FieldDecl *, QualType) {} virtual void handleSyclSpecConstantType(const FieldDecl *, QualType) {} + virtual void handleSyclStreamType(const CXXBaseSpecifier &, QualType) {} virtual void handleSyclStreamType(const FieldDecl *, QualType) {} virtual void handleStructType(const FieldDecl *, QualType) {} virtual void handleReferenceType(const FieldDecl *, QualType) {} @@ -1371,11 +1372,11 @@ class SyclKernelBodyCreator // Using the statements/init expressions that we've created, this generates // the kernel body compound stmt. CompoundStmt needs to know its number of // statements in advance to allocate it, so we cannot do this as we go along. - CompountStmt *createKernelBody() { + CompoundStmt *createKernelBody() { // TODO: Can we hold off on creating KernelObjClone to here? - Expr *ILE = new (SemaRef.getASTContext()) - InitListExpr(S.Context, SourceLocation(), InitExprs, SourceLocation()); + Expr *ILE = new (SemaRef.getASTContext()) InitListExpr( + SemaRef.getASTContext(), SourceLocation(), InitExprs, SourceLocation()); // TODO!!! ILE->setType(QualType(LC->getTypeForDecl(), 0)); // KernelObjectClone->setInit(ILE); @@ -1391,37 +1392,47 @@ class SyclKernelBodyCreator } public: - SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, KernelInvocationKind K) + SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, + KernelInvocationKind K) : SyclKernelFieldHandler(S), DeclCreator(DC) { - // TODO: Something special with the lambda when InvokeParallelForWorkGroup. - if (K == InvokeParallelForWorkGroup) - do somethingForparalellForWorkGroup(); - } + // TODO: Something special with the lambda when InvokeParallelForWorkGroup. + if (K == InvokeParallelForWorkGroup) + doSomethingForParallelForWorkGroup(); + } ~SyclKernelBodyCreator() { CompoundStmt *KernelBody = createKernelBody(); DeclCreator.setBody(KernelBody); } - void handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { + void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) final { + // TODO: Creates init sequence and inits special sycl obj + } + + void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { // TODO: Creates init sequence and inits special sycl obj } - void handleSyclSamplerType(FieldDecl *FD, QualType Ty) final { + + void handleSyclSamplerType(const FieldDecl *FD, QualType Ty) final { // TODO: Creates init sequence and inits special sycl obj } - void handleSyclStreamType(FieldDecl *FD, QualType Ty) final { + void handleSyclStreamType(const FieldDecl *FD, QualType Ty) final { + // TODO: Creates init/finalize sequence and inits special sycl obj + } + + void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) final { // TODO: Creates init/finalize sequence and inits special sycl obj } - void handleStructType(FieldDecl *FD, QualType Ty) final { + + void handleStructType(const FieldDecl *FD, QualType Ty) final { // TODO: a bunch of work doing inits, note this has a little more than // scalar. } - void handleScalarType(FieldDecl *FD, QualType Ty) final { + void handleScalarType(const FieldDecl *FD, QualType Ty) final { // TODO: a bunch of work doing inits. } - }; class SyclKernelIntHeaderCreator From 6b987c36440982b6d0bf90059d8fbbff0f84737b Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 13 Apr 2020 11:11:41 -0700 Subject: [PATCH 17/38] Add getParamVarDeclsForCurrentField, which provides a way for the body creator to get the list of parameters for the currently handled field Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d8a04d6d85d9e..30ea342073209 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1227,6 +1227,9 @@ class SyclKernelDeclCreator llvm::SmallVector Params; SyclKernelFieldChecker &ArgChecker; Sema::ContextRAII FuncContext; + // Holds the last handled field's first parameter. This doesn't store an + // iterator as push_back invalidates iterators. + size_t LastParamIndex = 0; void addParam(const FieldDecl *FD, QualType ArgTy) { ParamDesc newParamDesc = makeParamDesc(FD, ArgTy); @@ -1248,6 +1251,7 @@ class SyclKernelDeclCreator NewParam->setScopeInfo(0, Params.size()); NewParam->setIsUsed(); + LastParamIndex = Params.size(); Params.push_back(NewParam); } @@ -1261,8 +1265,12 @@ class SyclKernelDeclCreator CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method"); + // Don't do -1 here because we count on this to be the first parameter added + // (if any). + size_t ParamIndex = Params.size(); for (const ParmVarDecl *Param : InitMethod->parameters()) addParam(FD, Param->getType().getCanonicalType()); + LastParamIndex = ParamIndex; } static void setKernelImplicitAttrs(ASTContext &Context, FunctionDecl *FD, @@ -1322,8 +1330,12 @@ class SyclKernelDeclCreator CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method"); + // Don't do -1 here because we count on this to be the first parameter added + // (if any). + size_t ParamIndex = Params.size(); for (const ParmVarDecl *Param : InitMethod->parameters()) addParam(BS, Param->getType().getCanonicalType()); + LastParamIndex = ParamIndex; } void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { @@ -1360,6 +1372,11 @@ class SyclKernelDeclCreator void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } FunctionDecl *getKernelDecl() { return KernelDecl; } + + llvm::ArrayRef getParamVarDeclsForCurrentField() { + return ArrayRef(std::begin(Params) + LastParamIndex, + std::end(Params)); + } }; class SyclKernelBodyCreator From 726212bba4723a54b1b1df84e105315ad74c09f8 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 14 Apr 2020 13:49:46 +0300 Subject: [PATCH 18/38] Implement body generation for accessor, struct and scalar Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 137 +++++++++++++++++++++++++++++++++--- 1 file changed, 128 insertions(+), 9 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 30ea342073209..48f00aa8a8d0c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1385,17 +1385,22 @@ class SyclKernelBodyCreator llvm::SmallVector BodyStmts; llvm::SmallVector FinalizeStmts; llvm::SmallVector InitExprs; + VarDecl *KernelObjClone; + InitializedEntity VarEntity; + DeclRefExpr *KernelObjCloneRef; + CXXRecordDecl *KernelObj; + Expr *CurBase; + llvm::SmallVector Bases; // Using the statements/init expressions that we've created, this generates // the kernel body compound stmt. CompoundStmt needs to know its number of // statements in advance to allocate it, so we cannot do this as we go along. CompoundStmt *createKernelBody() { - // TODO: Can we hold off on creating KernelObjClone to here? Expr *ILE = new (SemaRef.getASTContext()) InitListExpr( SemaRef.getASTContext(), SourceLocation(), InitExprs, SourceLocation()); - // TODO!!! ILE->setType(QualType(LC->getTypeForDecl(), 0)); - // KernelObjectClone->setInit(ILE); + ILE->setType(QualType(KernelObj->getTypeForDecl(), 0)); + KernelObjClone->setInit(ILE); // TODO: More kernel object init with KernelBodyTransform. @@ -1408,13 +1413,80 @@ class SyclKernelBodyCreator void doSomethingForParallelForWorkGroup() { } + MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { + DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); + MemberExpr *Result = SemaRef.BuildMemberExpr( + Base, /*IsArrow */ false, SourceLocation(), NestedNameSpecifierLoc(), + SourceLocation(), Member, MemberDAP, + /*HadMultipleCandidates*/ false, + DeclarationNameInfo(Member->getDeclName(), SourceLocation()), + Member->getType(), VK_LValue, OK_Ordinary); + return Result; + } + + void createSpecialMethodCall(const CXXRecordDecl *SpecialClass, Expr *Base, + const std::string &MethodName, FieldDecl *Field) { + CXXMethodDecl *Method = getMethodByName(SpecialClass, MethodName); + assert(Method && + "The accessor/sampler/stream must have the __init method. Stream" + " must also have __finalize method"); + unsigned NumParams = Method->getNumParams(); + llvm::SmallVector ParamDREs(NumParams); + llvm::ArrayRef KernelParameters = + DeclCreator.getParamVarDeclsForCurrentField(); + for (size_t I = 0; I < NumParams; ++I) { + QualType ParamType = KernelParameters[I]->getOriginalType(); + ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType, + VK_LValue, SourceLocation()); + } + + MemberExpr *SpecialObjME = BuildMemberExpr(Base, Field); + + MemberExpr *MethodME = BuildMemberExpr(SpecialObjME, Method); + + QualType ResultTy = Method->getReturnType(); + ExprValueKind VK = Expr::getValueKindForType(ResultTy); + ResultTy = ResultTy.getNonLValueExprType(SemaRef.Context); + llvm::SmallVector ParamStmts; + const auto *Proto = cast(Method->getType()); + SemaRef.GatherArgumentsForCall(SourceLocation(), Method, Proto, 0, + ParamDREs, ParamStmts); + // [kernel_obj or wrapper object].accessor.__init(_ValueType*, + // range, range, id) + CXXMemberCallExpr *Call = + CXXMemberCallExpr::Create(SemaRef.Context, MethodME, ParamStmts, + ResultTy, VK, SourceLocation()); + BodyStmts.push_back(Call); + + } + public: SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, - KernelInvocationKind K) - : SyclKernelFieldHandler(S), DeclCreator(DC) { + KernelInvocationKind K, CXXRecordDecl *KernelObj) + : SyclKernelFieldHandler(S), DeclCreator(DC), + VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), + KernelObj(KernelObj) { // TODO: Something special with the lambda when InvokeParallelForWorkGroup. if (K == InvokeParallelForWorkGroup) doSomethingForParallelForWorkGroup(); + + // TODO get rid of kernel obj clone + TypeSourceInfo *TSInfo = + KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr; + KernelObjClone = VarDecl::Create( + SemaRef.Context, DeclCreator.getKernelDecl(), SourceLocation(), + SourceLocation(), KernelObj->getIdentifier(), + QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None); + Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), + SourceLocation(), SourceLocation()); + BodyStmts.push_back(DS); + KernelObjCloneRef = DeclRefExpr::Create( + S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, + false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), + VK_LValue); + VarEntity = InitializedEntity::InitializeVariable(KernelObjClone); + CurBase = KernelObjCloneRef; + Bases.push_back(CurBase); } ~SyclKernelBodyCreator() { CompoundStmt *KernelBody = createKernelBody(); @@ -1423,6 +1495,19 @@ class SyclKernelBodyCreator void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) final { // TODO: Creates init sequence and inits special sycl obj + const auto *AccDecl = Ty->getAsCXXRecordDecl(); + // TODO : we don't need all this init stuff if remove kernel obj clone + InitializedEntity Entity = InitializedEntity::InitializeMember( + const_cast(FD), &VarEntity); + // Initialize with the default constructor. + InitializationKind InitKind = + InitializationKind::CreateDefault(SourceLocation()); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); + ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); + InitExprs.push_back(MemberInit.get()); + // TODO don't do const-cast + createSpecialMethodCall(AccDecl, CurBase, InitMethodName, + const_cast(FD)); } void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { @@ -1442,13 +1527,46 @@ class SyclKernelBodyCreator // TODO: Creates init/finalize sequence and inits special sycl obj } + void CreateExprForStructOrScalar(FieldDecl *FD) { + ParmVarDecl *KernelParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity); + QualType ParamType = KernelParameter->getOriginalType(); + Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, + SourceLocation()); + InitializationKind InitKind = + InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE); + + ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE); + InitExprs.push_back(MemberInit.get()); + } void handleStructType(const FieldDecl *FD, QualType Ty) final { - // TODO: a bunch of work doing inits, note this has a little more than - // scalar. + CreateExprForStructOrScalar(const_cast(FD)); } + void handleScalarType(const FieldDecl *FD, QualType Ty) final { - // TODO: a bunch of work doing inits. + CreateExprForStructOrScalar(const_cast(FD)); + } + + void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + CurBase = BuildMemberExpr(CurBase, const_cast(FD)); + Bases.push_back(CurBase); + } + + void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + CurBase = Bases.back(); + Bases.pop_back(); + } + + void enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + // TODO : do something here? + } + + void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + // TODO : do something here? } }; @@ -1618,7 +1736,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, KernelLambda->getLocation(), KernelCallerFunc->isInlined()); SyclKernelBodyCreator kernel_body(*this, kernel_decl, - getKernelInvocationKind(KernelCallerFunc)); + getKernelInvocationKind(KernelCallerFunc), + KernelLambda); SyclKernelIntHeaderCreator int_header( *this, getSyclIntegrationHeader(), KernelLambda, calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, From 3b2e900e89cfa0fc5d499e0529bd419d82acc6e6 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 14 Apr 2020 17:44:03 +0300 Subject: [PATCH 19/38] Apply clang format and a couple of comments Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 107 +++++++++++++++++------------------- 1 file changed, 50 insertions(+), 57 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 48f00aa8a8d0c..bfeaae80aad36 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1389,7 +1389,6 @@ class SyclKernelBodyCreator InitializedEntity VarEntity; DeclRefExpr *KernelObjCloneRef; CXXRecordDecl *KernelObj; - Expr *CurBase; llvm::SmallVector Bases; // Using the statements/init expressions that we've created, this generates @@ -1410,8 +1409,7 @@ class SyclKernelBodyCreator } // TODO: not sure what this does yet, name is a placeholder for future use. - void doSomethingForParallelForWorkGroup() { - } + void doSomethingForParallelForWorkGroup() {} MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); @@ -1424,40 +1422,55 @@ class SyclKernelBodyCreator return Result; } - void createSpecialMethodCall(const CXXRecordDecl *SpecialClass, Expr *Base, - const std::string &MethodName, FieldDecl *Field) { - CXXMethodDecl *Method = getMethodByName(SpecialClass, MethodName); - assert(Method && - "The accessor/sampler/stream must have the __init method. Stream" - " must also have __finalize method"); - unsigned NumParams = Method->getNumParams(); - llvm::SmallVector ParamDREs(NumParams); - llvm::ArrayRef KernelParameters = - DeclCreator.getParamVarDeclsForCurrentField(); - for (size_t I = 0; I < NumParams; ++I) { - QualType ParamType = KernelParameters[I]->getOriginalType(); - ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType, - VK_LValue, SourceLocation()); - } + void CreateExprForStructOrScalar(FieldDecl *FD) { + ParmVarDecl *KernelParameter = + DeclCreator.getParamVarDeclsForCurrentField()[0]; + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity); + QualType ParamType = KernelParameter->getOriginalType(); + Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, + SourceLocation()); + InitializationKind InitKind = + InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE); - MemberExpr *SpecialObjME = BuildMemberExpr(Base, Field); + ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE); + InitExprs.push_back(MemberInit.get()); + } - MemberExpr *MethodME = BuildMemberExpr(SpecialObjME, Method); + void createSpecialMethodCall(const CXXRecordDecl *SpecialClass, Expr *Base, + const std::string &MethodName, + FieldDecl *Field) { + CXXMethodDecl *Method = getMethodByName(SpecialClass, MethodName); + assert(Method && + "The accessor/sampler/stream must have the __init method. Stream" + " must also have __finalize method"); + unsigned NumParams = Method->getNumParams(); + llvm::SmallVector ParamDREs(NumParams); + llvm::ArrayRef KernelParameters = + DeclCreator.getParamVarDeclsForCurrentField(); + for (size_t I = 0; I < NumParams; ++I) { + QualType ParamType = KernelParameters[I]->getOriginalType(); + ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType, + VK_LValue, SourceLocation()); + } - QualType ResultTy = Method->getReturnType(); - ExprValueKind VK = Expr::getValueKindForType(ResultTy); - ResultTy = ResultTy.getNonLValueExprType(SemaRef.Context); - llvm::SmallVector ParamStmts; - const auto *Proto = cast(Method->getType()); - SemaRef.GatherArgumentsForCall(SourceLocation(), Method, Proto, 0, - ParamDREs, ParamStmts); - // [kernel_obj or wrapper object].accessor.__init(_ValueType*, - // range, range, id) - CXXMemberCallExpr *Call = - CXXMemberCallExpr::Create(SemaRef.Context, MethodME, ParamStmts, - ResultTy, VK, SourceLocation()); - BodyStmts.push_back(Call); + MemberExpr *SpecialObjME = BuildMemberExpr(Base, Field); + MemberExpr *MethodME = BuildMemberExpr(SpecialObjME, Method); + + QualType ResultTy = Method->getReturnType(); + ExprValueKind VK = Expr::getValueKindForType(ResultTy); + ResultTy = ResultTy.getNonLValueExprType(SemaRef.Context); + llvm::SmallVector ParamStmts; + const auto *Proto = cast(Method->getType()); + SemaRef.GatherArgumentsForCall(SourceLocation(), Method, Proto, 0, + ParamDREs, ParamStmts); + // [kernel_obj or wrapper object].accessor.__init(_ValueType*, + // range, range, id) + CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( + SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, SourceLocation()); + BodyStmts.push_back(Call); } public: @@ -1468,7 +1481,7 @@ class SyclKernelBodyCreator KernelObj(KernelObj) { // TODO: Something special with the lambda when InvokeParallelForWorkGroup. if (K == InvokeParallelForWorkGroup) - doSomethingForParallelForWorkGroup(); + doSomethingForParallelForWorkGroup(); // TODO get rid of kernel obj clone TypeSourceInfo *TSInfo = @@ -1485,8 +1498,7 @@ class SyclKernelBodyCreator false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), VK_LValue); VarEntity = InitializedEntity::InitializeVariable(KernelObjClone); - CurBase = KernelObjCloneRef; - Bases.push_back(CurBase); + Bases.push_back(KernelObjCloneRef); } ~SyclKernelBodyCreator() { CompoundStmt *KernelBody = createKernelBody(); @@ -1506,7 +1518,7 @@ class SyclKernelBodyCreator ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); InitExprs.push_back(MemberInit.get()); // TODO don't do const-cast - createSpecialMethodCall(AccDecl, CurBase, InitMethodName, + createSpecialMethodCall(AccDecl, Bases.back(), InitMethodName, const_cast(FD)); } @@ -1514,7 +1526,6 @@ class SyclKernelBodyCreator // TODO: Creates init sequence and inits special sycl obj } - void handleSyclSamplerType(const FieldDecl *FD, QualType Ty) final { // TODO: Creates init sequence and inits special sycl obj } @@ -1527,22 +1538,6 @@ class SyclKernelBodyCreator // TODO: Creates init/finalize sequence and inits special sycl obj } - void CreateExprForStructOrScalar(FieldDecl *FD) { - ParmVarDecl *KernelParameter = - DeclCreator.getParamVarDeclsForCurrentField()[0]; - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity); - QualType ParamType = KernelParameter->getOriginalType(); - Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, - SourceLocation()); - InitializationKind InitKind = - InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE); - - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, DRE); - InitExprs.push_back(MemberInit.get()); - } - void handleStructType(const FieldDecl *FD, QualType Ty) final { CreateExprForStructOrScalar(const_cast(FD)); } @@ -1552,12 +1547,10 @@ class SyclKernelBodyCreator } void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { - CurBase = BuildMemberExpr(CurBase, const_cast(FD)); - Bases.push_back(CurBase); + Bases.push_back(BuildMemberExpr(Bases.back(), const_cast(FD))); } void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { - CurBase = Bases.back(); Bases.pop_back(); } From d642fb97f5340c11cc7b33995deed2248f8b056f Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 15 Apr 2020 13:07:11 +0300 Subject: [PATCH 20/38] Make LIT tests pass Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 54 ++++++++++++++++++++++++------------- 1 file changed, 36 insertions(+), 18 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index bfeaae80aad36..8334f2930d7b4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1387,9 +1387,9 @@ class SyclKernelBodyCreator llvm::SmallVector InitExprs; VarDecl *KernelObjClone; InitializedEntity VarEntity; - DeclRefExpr *KernelObjCloneRef; CXXRecordDecl *KernelObj; llvm::SmallVector Bases; + FunctionDecl *KernelCallerFunc; // Using the statements/init expressions that we've created, this generates // the kernel body compound stmt. CompoundStmt needs to know its number of @@ -1400,8 +1400,22 @@ class SyclKernelBodyCreator SemaRef.getASTContext(), SourceLocation(), InitExprs, SourceLocation()); ILE->setType(QualType(KernelObj->getTypeForDecl(), 0)); KernelObjClone->setInit(ILE); + Stmt *FunctionBody = KernelCallerFunc->getBody(); - // TODO: More kernel object init with KernelBodyTransform. + ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); + + // DeclRefExpr with valid source location but with decl which is not marked + // as used is invalid. + KernelObjClone->setIsUsed(); + std::pair MappingPair; + MappingPair.first = KernelObjParam; + MappingPair.second = KernelObjClone; + + // Function scope might be empty, so we do push + SemaRef.PushFunctionScope(); + KernelBodyTransform KBT(MappingPair, SemaRef); + Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); + BodyStmts.push_back(NewBody); BodyStmts.insert(std::end(BodyStmts), std::begin(FinalizeStmts), std::begin(FinalizeStmts)); @@ -1422,7 +1436,7 @@ class SyclKernelBodyCreator return Result; } - void CreateExprForStructOrScalar(FieldDecl *FD) { + void createExprForStructOrScalar(FieldDecl *FD) { ParmVarDecl *KernelParameter = DeclCreator.getParamVarDeclsForCurrentField()[0]; InitializedEntity Entity = @@ -1475,10 +1489,11 @@ class SyclKernelBodyCreator public: SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, - KernelInvocationKind K, CXXRecordDecl *KernelObj) + KernelInvocationKind K, CXXRecordDecl *KernelObj, + FunctionDecl *KernelCallerFunc) : SyclKernelFieldHandler(S), DeclCreator(DC), VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), - KernelObj(KernelObj) { + KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { // TODO: Something special with the lambda when InvokeParallelForWorkGroup. if (K == InvokeParallelForWorkGroup) doSomethingForParallelForWorkGroup(); @@ -1493,30 +1508,33 @@ class SyclKernelBodyCreator Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), SourceLocation(), SourceLocation()); BodyStmts.push_back(DS); - KernelObjCloneRef = DeclRefExpr::Create( + DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create( S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), VK_LValue); VarEntity = InitializedEntity::InitializeVariable(KernelObjClone); Bases.push_back(KernelObjCloneRef); } + ~SyclKernelBodyCreator() { CompoundStmt *KernelBody = createKernelBody(); DeclCreator.setBody(KernelBody); } void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) final { - // TODO: Creates init sequence and inits special sycl obj const auto *AccDecl = Ty->getAsCXXRecordDecl(); // TODO : we don't need all this init stuff if remove kernel obj clone - InitializedEntity Entity = InitializedEntity::InitializeMember( - const_cast(FD), &VarEntity); - // Initialize with the default constructor. - InitializationKind InitKind = - InitializationKind::CreateDefault(SourceLocation()); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); - InitExprs.push_back(MemberInit.get()); + // Perform initialization only if it is field of kernel object + if (Bases.size() == 1) { + InitializedEntity Entity = InitializedEntity::InitializeMember( + const_cast(FD), &VarEntity); + // Initialize with the default constructor. + InitializationKind InitKind = + InitializationKind::CreateDefault(SourceLocation()); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); + ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); + InitExprs.push_back(MemberInit.get()); + } // TODO don't do const-cast createSpecialMethodCall(AccDecl, Bases.back(), InitMethodName, const_cast(FD)); @@ -1539,11 +1557,11 @@ class SyclKernelBodyCreator } void handleStructType(const FieldDecl *FD, QualType Ty) final { - CreateExprForStructOrScalar(const_cast(FD)); + createExprForStructOrScalar(const_cast(FD)); } void handleScalarType(const FieldDecl *FD, QualType Ty) final { - CreateExprForStructOrScalar(const_cast(FD)); + createExprForStructOrScalar(const_cast(FD)); } void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { @@ -1730,7 +1748,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, KernelCallerFunc->isInlined()); SyclKernelBodyCreator kernel_body(*this, kernel_decl, getKernelInvocationKind(KernelCallerFunc), - KernelLambda); + KernelLambda, KernelCallerFunc); SyclKernelIntHeaderCreator int_header( *this, getSyclIntegrationHeader(), KernelLambda, calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, From c4da2ce912698854165f1cd928a3adc6a98770df Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 15 Apr 2020 14:56:04 +0300 Subject: [PATCH 21/38] Remove const casts Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 100 +++++++++++++++++++++++++++--------- 1 file changed, 76 insertions(+), 24 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8334f2930d7b4..00cb15410d414 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1144,25 +1144,78 @@ template class SyclKernelFieldHandler { // should be still working. // Accessor can be a base class or a field decl, so both must be handled. - virtual void handleSyclAccessorType(const CXXBaseSpecifier &, QualType) {} - virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} - virtual void handleSyclSamplerType(const FieldDecl *, QualType) {} - virtual void handleSyclSpecConstantType(const FieldDecl *, QualType) {} - virtual void handleSyclStreamType(const CXXBaseSpecifier &, QualType) {} - virtual void handleSyclStreamType(const FieldDecl *, QualType) {} - virtual void handleStructType(const FieldDecl *, QualType) {} - virtual void handleReferenceType(const FieldDecl *, QualType) {} - virtual void handlePointerType(const FieldDecl *, QualType) {} - virtual void handleArrayType(const FieldDecl *, QualType) {} - virtual void handleScalarType(const FieldDecl *, QualType) {} + virtual void handleSyclAccessorType(CXXBaseSpecifier &, QualType) {} + virtual void handleSyclAccessorType(FieldDecl *, QualType) {} + virtual void handleSyclSamplerType(FieldDecl *, QualType) {} + virtual void handleSyclSpecConstantType(FieldDecl *, QualType) {} + virtual void handleSyclStreamType(CXXBaseSpecifier &, QualType) {} + virtual void handleSyclStreamType(FieldDecl *, QualType) {} + virtual void handleStructType(FieldDecl *, QualType) {} + virtual void handleReferenceType(FieldDecl *, QualType) {} + virtual void handlePointerType(FieldDecl *, QualType) {} + virtual void handleArrayType(FieldDecl *, QualType) {} + virtual void handleScalarType(FieldDecl *, QualType) {} // Most handlers shouldn't be handling this, just the field checker. - virtual void handleOtherType(const FieldDecl *, QualType) {} + virtual void handleOtherType(FieldDecl *, QualType) {} + + // TODO: fix warnings from derived classes + virtual void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) { + static_cast(this)->handleSyclAccessorType( + const_cast(BS), Ty); + } + virtual void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleSyclAccessorType( + const_cast(FD), Ty); + } + virtual void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) { + static_cast(this)->handleSyclStreamType( + const_cast(BS), Ty); + } + virtual void handleSyclStreamType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleSyclStreamType( + const_cast(FD), Ty); + } + virtual void handleSyclSamplerType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleSyclSamplerType( + const_cast(FD), Ty); + } + virtual void handleSyclSpecConstantType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleSyclSpecConstantType( + const_cast(FD), Ty); + } + virtual void handleStructType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleStructType( + const_cast(FD), Ty); + } + virtual void handleReferenceType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleReferenceType( + const_cast(FD), Ty); + } + virtual void handlePointerType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handlePointerType( + const_cast(FD), Ty); + } + virtual void handleArrayType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleArrayType( + const_cast(FD), Ty); + } + virtual void handleScalarType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleScalarType( + const_cast(FD), Ty); + } + virtual void handleOtherType(const FieldDecl *FD, QualType Ty) { + static_cast(this)->handleOtherType( + const_cast(FD), Ty); + } // The following are only used for keeping track of where we are in the base // class/field graph. Int Headers use this to calculate offset, most others // don't have a need for these. + virtual void enterStruct(const CXXRecordDecl *, FieldDecl *) {} - virtual void enterStruct(const CXXRecordDecl *, const FieldDecl *) {} + virtual void enterStruct(const CXXRecordDecl *RD, const FieldDecl *FD) { + static_cast(this)->enterStruct(RD, const_cast(FD)); + } virtual void leaveStruct(const CXXRecordDecl *, const FieldDecl *) {} virtual void enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} virtual void leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} @@ -1521,13 +1574,13 @@ class SyclKernelBodyCreator DeclCreator.setBody(KernelBody); } - void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) final { + void handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { const auto *AccDecl = Ty->getAsCXXRecordDecl(); // TODO : we don't need all this init stuff if remove kernel obj clone // Perform initialization only if it is field of kernel object if (Bases.size() == 1) { - InitializedEntity Entity = InitializedEntity::InitializeMember( - const_cast(FD), &VarEntity); + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity); // Initialize with the default constructor. InitializationKind InitKind = InitializationKind::CreateDefault(SourceLocation()); @@ -1536,8 +1589,7 @@ class SyclKernelBodyCreator InitExprs.push_back(MemberInit.get()); } // TODO don't do const-cast - createSpecialMethodCall(AccDecl, Bases.back(), InitMethodName, - const_cast(FD)); + createSpecialMethodCall(AccDecl, Bases.back(), InitMethodName, FD); } void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { @@ -1556,16 +1608,16 @@ class SyclKernelBodyCreator // TODO: Creates init/finalize sequence and inits special sycl obj } - void handleStructType(const FieldDecl *FD, QualType Ty) final { - createExprForStructOrScalar(const_cast(FD)); + void handleStructType(FieldDecl *FD, QualType Ty) final { + createExprForStructOrScalar(FD); } - void handleScalarType(const FieldDecl *FD, QualType Ty) final { - createExprForStructOrScalar(const_cast(FD)); + void handleScalarType(FieldDecl *FD, QualType Ty) final { + createExprForStructOrScalar(FD); } - void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { - Bases.push_back(BuildMemberExpr(Bases.back(), const_cast(FD))); + void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { + Bases.push_back(BuildMemberExpr(Bases.back(), FD)); } void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { From d37f1df2d0949a9132ab875437bc59d42294af6c Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 15 Apr 2020 16:25:39 +0300 Subject: [PATCH 22/38] Revert "Remove const casts" This reverts commit c4da2ce912698854165f1cd928a3adc6a98770df. --- clang/lib/Sema/SemaSYCL.cpp | 100 +++++++++--------------------------- 1 file changed, 24 insertions(+), 76 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 00cb15410d414..8334f2930d7b4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1144,78 +1144,25 @@ template class SyclKernelFieldHandler { // should be still working. // Accessor can be a base class or a field decl, so both must be handled. - virtual void handleSyclAccessorType(CXXBaseSpecifier &, QualType) {} - virtual void handleSyclAccessorType(FieldDecl *, QualType) {} - virtual void handleSyclSamplerType(FieldDecl *, QualType) {} - virtual void handleSyclSpecConstantType(FieldDecl *, QualType) {} - virtual void handleSyclStreamType(CXXBaseSpecifier &, QualType) {} - virtual void handleSyclStreamType(FieldDecl *, QualType) {} - virtual void handleStructType(FieldDecl *, QualType) {} - virtual void handleReferenceType(FieldDecl *, QualType) {} - virtual void handlePointerType(FieldDecl *, QualType) {} - virtual void handleArrayType(FieldDecl *, QualType) {} - virtual void handleScalarType(FieldDecl *, QualType) {} + virtual void handleSyclAccessorType(const CXXBaseSpecifier &, QualType) {} + virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} + virtual void handleSyclSamplerType(const FieldDecl *, QualType) {} + virtual void handleSyclSpecConstantType(const FieldDecl *, QualType) {} + virtual void handleSyclStreamType(const CXXBaseSpecifier &, QualType) {} + virtual void handleSyclStreamType(const FieldDecl *, QualType) {} + virtual void handleStructType(const FieldDecl *, QualType) {} + virtual void handleReferenceType(const FieldDecl *, QualType) {} + virtual void handlePointerType(const FieldDecl *, QualType) {} + virtual void handleArrayType(const FieldDecl *, QualType) {} + virtual void handleScalarType(const FieldDecl *, QualType) {} // Most handlers shouldn't be handling this, just the field checker. - virtual void handleOtherType(FieldDecl *, QualType) {} - - // TODO: fix warnings from derived classes - virtual void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) { - static_cast(this)->handleSyclAccessorType( - const_cast(BS), Ty); - } - virtual void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleSyclAccessorType( - const_cast(FD), Ty); - } - virtual void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) { - static_cast(this)->handleSyclStreamType( - const_cast(BS), Ty); - } - virtual void handleSyclStreamType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleSyclStreamType( - const_cast(FD), Ty); - } - virtual void handleSyclSamplerType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleSyclSamplerType( - const_cast(FD), Ty); - } - virtual void handleSyclSpecConstantType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleSyclSpecConstantType( - const_cast(FD), Ty); - } - virtual void handleStructType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleStructType( - const_cast(FD), Ty); - } - virtual void handleReferenceType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleReferenceType( - const_cast(FD), Ty); - } - virtual void handlePointerType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handlePointerType( - const_cast(FD), Ty); - } - virtual void handleArrayType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleArrayType( - const_cast(FD), Ty); - } - virtual void handleScalarType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleScalarType( - const_cast(FD), Ty); - } - virtual void handleOtherType(const FieldDecl *FD, QualType Ty) { - static_cast(this)->handleOtherType( - const_cast(FD), Ty); - } + virtual void handleOtherType(const FieldDecl *, QualType) {} // The following are only used for keeping track of where we are in the base // class/field graph. Int Headers use this to calculate offset, most others // don't have a need for these. - virtual void enterStruct(const CXXRecordDecl *, FieldDecl *) {} - virtual void enterStruct(const CXXRecordDecl *RD, const FieldDecl *FD) { - static_cast(this)->enterStruct(RD, const_cast(FD)); - } + virtual void enterStruct(const CXXRecordDecl *, const FieldDecl *) {} virtual void leaveStruct(const CXXRecordDecl *, const FieldDecl *) {} virtual void enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} virtual void leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} @@ -1574,13 +1521,13 @@ class SyclKernelBodyCreator DeclCreator.setBody(KernelBody); } - void handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { + void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) final { const auto *AccDecl = Ty->getAsCXXRecordDecl(); // TODO : we don't need all this init stuff if remove kernel obj clone // Perform initialization only if it is field of kernel object if (Bases.size() == 1) { - InitializedEntity Entity = - InitializedEntity::InitializeMember(FD, &VarEntity); + InitializedEntity Entity = InitializedEntity::InitializeMember( + const_cast(FD), &VarEntity); // Initialize with the default constructor. InitializationKind InitKind = InitializationKind::CreateDefault(SourceLocation()); @@ -1589,7 +1536,8 @@ class SyclKernelBodyCreator InitExprs.push_back(MemberInit.get()); } // TODO don't do const-cast - createSpecialMethodCall(AccDecl, Bases.back(), InitMethodName, FD); + createSpecialMethodCall(AccDecl, Bases.back(), InitMethodName, + const_cast(FD)); } void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { @@ -1608,16 +1556,16 @@ class SyclKernelBodyCreator // TODO: Creates init/finalize sequence and inits special sycl obj } - void handleStructType(FieldDecl *FD, QualType Ty) final { - createExprForStructOrScalar(FD); + void handleStructType(const FieldDecl *FD, QualType Ty) final { + createExprForStructOrScalar(const_cast(FD)); } - void handleScalarType(FieldDecl *FD, QualType Ty) final { - createExprForStructOrScalar(FD); + void handleScalarType(const FieldDecl *FD, QualType Ty) final { + createExprForStructOrScalar(const_cast(FD)); } - void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { - Bases.push_back(BuildMemberExpr(Bases.back(), FD)); + void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + Bases.push_back(BuildMemberExpr(Bases.back(), const_cast(FD))); } void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { From 7e3ff18a485729ded13f7a4c3ca9fb2f028b34b3 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Thu, 16 Apr 2020 12:46:16 +0300 Subject: [PATCH 23/38] Add support for sampler, start struggling with stream Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 106 ++++++++++++++++++++++++------------ 1 file changed, 71 insertions(+), 35 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8334f2930d7b4..a4fd85e374bb2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1069,12 +1069,12 @@ static void VisitAccessorWrapperHelper(CXXRecordDecl *Owner, RangeTy Range, if (Util::isSyclAccessorType(ItemTy)) (void)std::initializer_list{ (handlers.handleSyclAccessorType(Item, ItemTy), 0)...}; - else if (Util::isSyclStreamType(ItemTy)) - (void)std::initializer_list{ - (handlers.handleSyclStreamType(Item, ItemTy), 0)...}; else if (ItemTy->isStructureOrClassType()) { VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(), handlers...); + if (Util::isSyclStreamType(ItemTy)) + (void)std::initializer_list{ + (handlers.handleSyclStreamType(Item, ItemTy), 0)...}; } } } @@ -1108,9 +1108,12 @@ static void VisitRecordFields(RecordDecl::field_range Fields, KF_FOR_EACH(handleSyclSamplerType); else if (Util::isSyclSpecConstantType(FieldTy)) KF_FOR_EACH(handleSyclSpecConstantType); - else if (Util::isSyclStreamType(FieldTy)) + else if (Util::isSyclStreamType(FieldTy)) { + // Stream actually wraps accessors, so do recursion KF_FOR_EACH(handleSyclStreamType); - else if (FieldTy->isStructureOrClassType()) { + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); + VisitAccessorWrapper(nullptr, Field, RD, handlers...); + } else if (FieldTy->isStructureOrClassType()) { KF_FOR_EACH(handleStructType); CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); VisitAccessorWrapper(nullptr, Field, RD, handlers...); @@ -1369,6 +1372,14 @@ class SyclKernelDeclCreator addParam(FD, ArgTy); } + void handleSyclStreamType(const FieldDecl *FD, QualType ArgTy) final { + addParam(FD, ArgTy); + } + + void handleSyclStreamType(const CXXBaseSpecifier &, QualType ArgTy) final { + // TODO do something + } + void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } FunctionDecl *getKernelDecl() { return KernelDecl; } @@ -1417,8 +1428,8 @@ class SyclKernelBodyCreator Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); BodyStmts.push_back(NewBody); - BodyStmts.insert(std::end(BodyStmts), std::begin(FinalizeStmts), - std::begin(FinalizeStmts)); + BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(), + FinalizeStmts.end()); return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {}); } @@ -1470,7 +1481,6 @@ class SyclKernelBodyCreator } MemberExpr *SpecialObjME = BuildMemberExpr(Base, Field); - MemberExpr *MethodME = BuildMemberExpr(SpecialObjME, Method); QualType ResultTy = Method->getReturnType(); @@ -1484,7 +1494,39 @@ class SyclKernelBodyCreator // range, range, id) CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, SourceLocation()); - BodyStmts.push_back(Call); + if (MethodName == FinalizeMethodName) + FinalizeStmts.push_back(Call); + else + BodyStmts.push_back(Call); + } + + // TODO get rid of kernel obj clone + static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC, + CXXRecordDecl *KernelObj) { + TypeSourceInfo *TSInfo = + KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr; + VarDecl *VD = VarDecl::Create( + Ctx, DC, SourceLocation(), SourceLocation(), KernelObj->getIdentifier(), + QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None); + + return VD; + } + + void handleSpecialType(FieldDecl *FD, QualType Ty) { + const auto *RecordDecl = Ty->getAsCXXRecordDecl(); + // TODO : we don't need all this init stuff if remove kernel obj clone + // Perform initialization only if it is field of kernel object + if (Bases.size() == 1) { + InitializedEntity Entity = + InitializedEntity::InitializeMember(FD, &VarEntity); + // Initialize with the default constructor. + InitializationKind InitKind = + InitializationKind::CreateDefault(SourceLocation()); + InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); + ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); + InitExprs.push_back(MemberInit.get()); + } + createSpecialMethodCall(RecordDecl, Bases.back(), InitMethodName, FD); } public: @@ -1492,19 +1534,14 @@ class SyclKernelBodyCreator KernelInvocationKind K, CXXRecordDecl *KernelObj, FunctionDecl *KernelCallerFunc) : SyclKernelFieldHandler(S), DeclCreator(DC), + KernelObjClone(createKernelObjClone(S.getASTContext(), + DC.getKernelDecl(), KernelObj)), VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { // TODO: Something special with the lambda when InvokeParallelForWorkGroup. if (K == InvokeParallelForWorkGroup) doSomethingForParallelForWorkGroup(); - // TODO get rid of kernel obj clone - TypeSourceInfo *TSInfo = - KernelObj->isLambda() ? KernelObj->getLambdaTypeInfo() : nullptr; - KernelObjClone = VarDecl::Create( - SemaRef.Context, DeclCreator.getKernelDecl(), SourceLocation(), - SourceLocation(), KernelObj->getIdentifier(), - QualType(KernelObj->getTypeForDecl(), 0), TSInfo, SC_None); Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), SourceLocation(), SourceLocation()); BodyStmts.push_back(DS); @@ -1512,7 +1549,6 @@ class SyclKernelBodyCreator S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), VK_LValue); - VarEntity = InitializedEntity::InitializeVariable(KernelObjClone); Bases.push_back(KernelObjCloneRef); } @@ -1522,22 +1558,7 @@ class SyclKernelBodyCreator } void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) final { - const auto *AccDecl = Ty->getAsCXXRecordDecl(); - // TODO : we don't need all this init stuff if remove kernel obj clone - // Perform initialization only if it is field of kernel object - if (Bases.size() == 1) { - InitializedEntity Entity = InitializedEntity::InitializeMember( - const_cast(FD), &VarEntity); - // Initialize with the default constructor. - InitializationKind InitKind = - InitializationKind::CreateDefault(SourceLocation()); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, None); - ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); - InitExprs.push_back(MemberInit.get()); - } - // TODO don't do const-cast - createSpecialMethodCall(AccDecl, Bases.back(), InitMethodName, - const_cast(FD)); + handleSpecialType(const_cast(FD), Ty); } void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { @@ -1545,11 +1566,19 @@ class SyclKernelBodyCreator } void handleSyclSamplerType(const FieldDecl *FD, QualType Ty) final { - // TODO: Creates init sequence and inits special sycl obj + handleSpecialType(const_cast(FD), Ty); } void handleSyclStreamType(const FieldDecl *FD, QualType Ty) final { - // TODO: Creates init/finalize sequence and inits special sycl obj + const auto *StreamDecl = Ty->getAsCXXRecordDecl(); + // Hmm okay, we created "stream" kernel argument and did copy initialization + // but why? + // We also don't have SEMA tests for stream, only pretty small CodeGen test + createExprForStructOrScalar(const_cast(FD)); + createSpecialMethodCall(StreamDecl, Bases.back(), InitMethodName, + const_cast(FD)); + createSpecialMethodCall(StreamDecl, Bases.back(), FinalizeMethodName, + const_cast(FD)); } void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) final { @@ -1681,6 +1710,13 @@ class SyclKernelIntHeaderCreator addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); } + void handleSyclStreamType(const FieldDecl *FD, QualType ArgTy) final { + addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); + } + void handleSyclStreamType(const CXXBaseSpecifier &BC, QualType ArgTy) final { + // TODO implement + } + // Keep track of the current struct offset. void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; From f77fc36a76002e96ff79d679a3641ebdde2e2f3b Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Thu, 16 Apr 2020 14:18:34 +0300 Subject: [PATCH 24/38] Initialize stream after internal accessors Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a4fd85e374bb2..b741bd04fb1a4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1110,9 +1110,9 @@ static void VisitRecordFields(RecordDecl::field_range Fields, KF_FOR_EACH(handleSyclSpecConstantType); else if (Util::isSyclStreamType(FieldTy)) { // Stream actually wraps accessors, so do recursion - KF_FOR_EACH(handleSyclStreamType); CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); VisitAccessorWrapper(nullptr, Field, RD, handlers...); + KF_FOR_EACH(handleSyclStreamType); } else if (FieldTy->isStructureOrClassType()) { KF_FOR_EACH(handleStructType); CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); From a8982bc7a5d5ae577df159303700c9850892673a Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Thu, 16 Apr 2020 14:38:37 +0300 Subject: [PATCH 25/38] Handle pointers Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b741bd04fb1a4..d99a1136cc441 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1455,6 +1455,12 @@ class SyclKernelBodyCreator QualType ParamType = KernelParameter->getOriginalType(); Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, SourceLocation()); + if (FD->getType()->isPointerType() && + FD->getType()->getPointeeType().getAddressSpace() != + ParamType->getPointeeType().getAddressSpace()) + DRE = ImplicitCastExpr::Create(SemaRef.Context, FD->getType(), + CK_AddressSpaceConversion, DRE, nullptr, + VK_RValue); InitializationKind InitKind = InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); InitializationSequence InitSeq(SemaRef, Entity, InitKind, DRE); @@ -1585,6 +1591,10 @@ class SyclKernelBodyCreator // TODO: Creates init/finalize sequence and inits special sycl obj } + void handlePointerType(const FieldDecl *FD, QualType ArgTy) final { + createExprForStructOrScalar(const_cast(FD)); + } + void handleStructType(const FieldDecl *FD, QualType Ty) final { createExprForStructOrScalar(const_cast(FD)); } From e564a88fa587358e9a51704ee2b5936cb4bfd99c Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Thu, 16 Apr 2020 15:05:56 +0300 Subject: [PATCH 26/38] Fix hp Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 69 ++++++++++++++++++------------------- 1 file changed, 33 insertions(+), 36 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d99a1136cc441..cbf35fa690c18 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -723,34 +723,6 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S, CXXRecordDecl *LC = getKernelObjectType(KernelCallerFunc); assert(LC && "Kernel object must be available"); - if (getKernelInvocationKind(KernelCallerFunc) == InvokeParallelForWorkGroup) { - CXXRecordDecl *LambdaObjTy = - KernelCallerFunc->getParamDecl(0)->getType()->getAsCXXRecordDecl(); - assert(LambdaObjTy && - "unexpected kernel_parallel_for_work_group parameter type"); - FindPFWGLambdaFnVisitor V(LambdaObjTy); - V.TraverseStmt(KernelCallerFunc->getBody()); - CXXMethodDecl *WGLambdaFn = V.getLambdaFn(); - assert(WGLambdaFn && "PFWG lambda not found"); - // Mark the function that it "works" in a work group scope: - // NOTE: In case of parallel_for_work_item the marker call itself is marked - // with work item scope attribute, here the '()' operator of the - // object passed as parameter is marked. This is an optimization - - // there are a lot of locals created at parallel_for_work_group scope - // before calling the lambda - it is more efficient to have all of - // them in the private address space rather then sharing via the local - // AS. See parallel_for_work_group implementation in the SYCL headers. - if (!WGLambdaFn->hasAttr()) { - WGLambdaFn->addAttr(SYCLScopeAttr::CreateImplicit( - S.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); - // Search and mark parallel_for_work_item calls: - MarkWIScopeFnVisitor MarkWIScope(S.getASTContext()); - MarkWIScope.TraverseDecl(WGLambdaFn); - // Now mark local variables declared in the PFWG lambda with work group - // scope attribute - addScopeAttrToLocalVars(*WGLambdaFn); - } - } TypeSourceInfo *TSInfo = LC->isLambda() ? LC->getLambdaTypeInfo() : nullptr; @@ -1433,8 +1405,35 @@ class SyclKernelBodyCreator return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, {}, {}); } - // TODO: not sure what this does yet, name is a placeholder for future use. - void doSomethingForParallelForWorkGroup() {} + void markParallelWorkItemCalls() { + if (getKernelInvocationKind(KernelCallerFunc) == + InvokeParallelForWorkGroup) { + FindPFWGLambdaFnVisitor V(KernelObj); + V.TraverseStmt(KernelCallerFunc->getBody()); + CXXMethodDecl *WGLambdaFn = V.getLambdaFn(); + assert(WGLambdaFn && "PFWG lambda not found"); + // Mark the function that it "works" in a work group scope: + // NOTE: In case of parallel_for_work_item the marker call itself is + // marked + // with work item scope attribute, here the '()' operator of the + // object passed as parameter is marked. This is an optimization - + // there are a lot of locals created at parallel_for_work_group + // scope before calling the lambda - it is more efficient to have + // all of them in the private address space rather then sharing via + // the local AS. See parallel_for_work_group implementation in the + // SYCL headers. + if (!WGLambdaFn->hasAttr()) { + WGLambdaFn->addAttr(SYCLScopeAttr::CreateImplicit( + SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); + // Search and mark parallel_for_work_item calls: + MarkWIScopeFnVisitor MarkWIScope(SemaRef.getASTContext()); + MarkWIScope.TraverseDecl(WGLambdaFn); + // Now mark local variables declared in the PFWG lambda with work group + // scope attribute + addScopeAttrToLocalVars(*WGLambdaFn); + } + } + } MemberExpr *BuildMemberExpr(Expr *Base, ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); @@ -1537,7 +1536,7 @@ class SyclKernelBodyCreator public: SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, - KernelInvocationKind K, CXXRecordDecl *KernelObj, + CXXRecordDecl *KernelObj, FunctionDecl *KernelCallerFunc) : SyclKernelFieldHandler(S), DeclCreator(DC), KernelObjClone(createKernelObjClone(S.getASTContext(), @@ -1545,8 +1544,7 @@ class SyclKernelBodyCreator VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { // TODO: Something special with the lambda when InvokeParallelForWorkGroup. - if (K == InvokeParallelForWorkGroup) - doSomethingForParallelForWorkGroup(); + markParallelWorkItemCalls(); Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), SourceLocation(), SourceLocation()); @@ -1792,9 +1790,8 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, SyclKernelDeclCreator kernel_decl(*this, checker, KernelName, KernelLambda->getLocation(), KernelCallerFunc->isInlined()); - SyclKernelBodyCreator kernel_body(*this, kernel_decl, - getKernelInvocationKind(KernelCallerFunc), - KernelLambda, KernelCallerFunc); + SyclKernelBodyCreator kernel_body(*this, kernel_decl, KernelLambda, + KernelCallerFunc); SyclKernelIntHeaderCreator int_header( *this, getSyclIntegrationHeader(), KernelLambda, calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, From 39f128b6e3e7001b64e8cd1f014a02f4c584d304 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Thu, 16 Apr 2020 11:22:15 -0700 Subject: [PATCH 27/38] Fix integration header problem with unnamed lambdas Integration header generation was being given the wrong name for the kernel because Erich misunderstood how the integration header generation used names. --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index cbf35fa690c18..75f602f4c5eb8 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1794,7 +1794,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, KernelCallerFunc); SyclKernelIntHeaderCreator int_header( *this, getSyclIntegrationHeader(), KernelLambda, - calculateKernelNameType(Context, KernelCallerFunc), CalculatedName, + calculateKernelNameType(Context, KernelCallerFunc), KernelName, StableName); ConstructingOpenCLKernel = true; From 3c7fed952345369254aaf4cc15dd2b351d2e3c4a Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Fri, 17 Apr 2020 08:51:50 +0300 Subject: [PATCH 28/38] Remove CreateOpenCLKernelBody Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 262 ------------------------------------ 1 file changed, 262 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 75f602f4c5eb8..b5b53088e1f0f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -710,259 +710,6 @@ static CXXRecordDecl *getKernelObjectType(FunctionDecl *Caller) { return (*Caller->param_begin())->getType()->getAsCXXRecordDecl(); } -// Creates body for new OpenCL kernel. This body contains initialization of SYCL -// kernel object fields with kernel parameters and a little bit transformed body -// of the kernel caller function. -static CompoundStmt *CreateOpenCLKernelBody(Sema &S, - FunctionDecl *KernelCallerFunc, - DeclContext *KernelDecl) { - using BodyStmtsT = llvm::SmallVector; - - BodyStmtsT BodyStmts; - BodyStmtsT FinalizeStmts; - CXXRecordDecl *LC = getKernelObjectType(KernelCallerFunc); - assert(LC && "Kernel object must be available"); - - - TypeSourceInfo *TSInfo = LC->isLambda() ? LC->getLambdaTypeInfo() : nullptr; - - // Create a local kernel object (lambda or functor) assembled from the - // incoming formal parameters - auto KernelObjClone = VarDecl::Create( - S.Context, KernelDecl, SourceLocation(), SourceLocation(), - LC->getIdentifier(), QualType(LC->getTypeForDecl(), 0), TSInfo, SC_None); - Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), - SourceLocation(), SourceLocation()); - BodyStmts.push_back(DS); - auto KernelObjCloneRef = - DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), SourceLocation(), - KernelObjClone, false, DeclarationNameInfo(), - QualType(LC->getTypeForDecl(), 0), VK_LValue); - - auto KernelFuncDecl = cast(KernelDecl); - auto KernelFuncParam = - KernelFuncDecl->param_begin(); // Iterator to ParamVarDecl (VarDecl) - if (KernelFuncParam) { - llvm::SmallVector InitExprs; - InitializedEntity VarEntity = - InitializedEntity::InitializeVariable(KernelObjClone); - for (auto Field : LC->fields()) { - // Creates Expression for special SYCL object: accessor or sampler. - // All special SYCL objects must have __init method, here we use it to - // initialize them. We create call of __init method and pass built kernel - // arguments as parameters to the __init method. - auto getExprForSpecialSYCLObj = [&](const QualType ¶mTy, - FieldDecl *Field, - const CXXRecordDecl *CRD, - Expr *Base, - const std::string &MethodName, - BodyStmtsT &Statements) { - CXXMethodDecl *Method = getMethodByName(CRD, MethodName); - assert(Method && - "The accessor/sampler/stream must have the __init method. Stream" - " must also have __finalize method"); - unsigned NumParams = Method->getNumParams(); - llvm::SmallVector ParamDREs(NumParams); - auto KFP = KernelFuncParam; - for (size_t I = 0; I < NumParams; ++KFP, ++I) { - QualType ParamType = (*KFP)->getOriginalType(); - ParamDREs[I] = DeclRefExpr::Create( - S.Context, NestedNameSpecifierLoc(), SourceLocation(), *KFP, - false, DeclarationNameInfo(), ParamType, VK_LValue); - } - - if (NumParams) - std::advance(KernelFuncParam, NumParams - 1); - - DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); - // [kernel_obj or wrapper object].special_obj - auto SpecialObjME = MemberExpr::Create( - S.Context, Base, false, SourceLocation(), NestedNameSpecifierLoc(), - SourceLocation(), Field, FieldDAP, - DeclarationNameInfo(Field->getDeclName(), SourceLocation()), - nullptr, Field->getType(), VK_LValue, OK_Ordinary, NOUR_None); - - // [kernel_obj or wrapper object].special_obj.__init - DeclAccessPair MethodDAP = DeclAccessPair::make(Method, AS_none); - auto ME = MemberExpr::Create( - S.Context, SpecialObjME, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), Method, MethodDAP, - DeclarationNameInfo(Method->getDeclName(), SourceLocation()), - nullptr, Method->getType(), VK_LValue, OK_Ordinary, NOUR_None); - - // Not referenced -> not emitted - S.MarkFunctionReferenced(SourceLocation(), Method, true); - - QualType ResultTy = Method->getReturnType(); - ExprValueKind VK = Expr::getValueKindForType(ResultTy); - ResultTy = ResultTy.getNonLValueExprType(S.Context); - - llvm::SmallVector ParamStmts; - const auto *Proto = cast(Method->getType()); - S.GatherArgumentsForCall(SourceLocation(), Method, Proto, 0, - ParamDREs, ParamStmts); - // [kernel_obj or wrapper object].accessor.__init(_ValueType*, - // range, range, id) - CXXMemberCallExpr *Call = CXXMemberCallExpr::Create( - S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation()); - Statements.push_back(Call); - }; - - // Recursively search for accessor fields to initialize them with kernel - // parameters - std::function - getExprForWrappedAccessorInit = - [&](const CXXRecordDecl *CRD, Expr *Base) { - for (auto *WrapperFld : CRD->fields()) { - QualType FldType = WrapperFld->getType(); - CXXRecordDecl *WrapperFldCRD = FldType->getAsCXXRecordDecl(); - if (FldType->isStructureOrClassType()) { - if (Util::isSyclAccessorType(FldType)) { - // Accessor field found - create expr to initialize this - // accessor object. Need to start from the next target - // function parameter, since current one is the wrapper - // object or parameter of the previous processed accessor - // object. - KernelFuncParam++; - getExprForSpecialSYCLObj(FldType, WrapperFld, - WrapperFldCRD, Base, - InitMethodName, BodyStmts); - } else if (Util::isSyclSpecConstantType(FldType)) { - // Specialization constants are "invisible" to the - // kernel argument creation and device-side SYCL object - // materialization infrastructure in this source. - // It is OK not to really materialize them on the kernel - // side, because their only use can be via - // 'spec_const_obj.get()' method, which is translated to - // an intrinsic and 'this' is really never used. - } else { - // Field is a structure or class so change the wrapper - // object and recursively search for accessor field. - DeclAccessPair WrapperFieldDAP = - DeclAccessPair::make(WrapperFld, AS_none); - auto NewBase = MemberExpr::Create( - S.Context, Base, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), - WrapperFld, WrapperFieldDAP, - DeclarationNameInfo(WrapperFld->getDeclName(), - SourceLocation()), - nullptr, WrapperFld->getType(), VK_LValue, - OK_Ordinary, NOUR_None); - getExprForWrappedAccessorInit(WrapperFldCRD, NewBase); - } - } - } - }; - - // Run through kernel object fields and add initialization for them using - // built kernel parameters. There are a several possible cases: - // - Kernel object field is a SYCL special object (SYCL accessor or SYCL - // sampler). These objects has a special initialization scheme - using - // __init method. - // - Kernel object field has a scalar type. In this case we should add - // simple initialization. - // - Kernel object field has a structure or class type. Same handling as - // a scalar but we should check if this structure/class contains - // accessors and add initialization for them properly. - QualType FieldType = Field->getType(); - CXXRecordDecl *CRD = FieldType->getAsCXXRecordDecl(); - InitializedEntity Entity = - InitializedEntity::InitializeMember(Field, &VarEntity); - if (Util::isSyclAccessorType(FieldType) || - Util::isSyclSamplerType(FieldType)) { - // Initialize with the default constructor. - InitializationKind InitKind = - InitializationKind::CreateDefault(SourceLocation()); - InitializationSequence InitSeq(S, Entity, InitKind, None); - ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, None); - InitExprs.push_back(MemberInit.get()); - getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef, - InitMethodName, BodyStmts); - } else if (Util::isSyclSpecConstantType(FieldType)) { - // Just skip specialization constants - not part of signature. - } else if (CRD || FieldType->isScalarType()) { - // If field has built-in or a structure/class type just initialize - // this field with corresponding kernel argument using copy - // initialization. - QualType ParamType = (*KernelFuncParam)->getOriginalType(); - Expr *DRE = - DeclRefExpr::Create(S.Context, NestedNameSpecifierLoc(), - SourceLocation(), *KernelFuncParam, false, - DeclarationNameInfo(), ParamType, VK_LValue); - - if (FieldType->isPointerType() && - FieldType->getPointeeType().getAddressSpace() != - ParamType->getPointeeType().getAddressSpace()) - DRE = ImplicitCastExpr::Create(S.Context, FieldType, - CK_AddressSpaceConversion, DRE, - nullptr, VK_RValue); - InitializationKind InitKind = - InitializationKind::CreateCopy(SourceLocation(), SourceLocation()); - InitializationSequence InitSeq(S, Entity, InitKind, DRE); - - ExprResult MemberInit = InitSeq.Perform(S, Entity, InitKind, DRE); - InitExprs.push_back(MemberInit.get()); - - if (CRD) { - // If a structure/class type has accessor fields then we need to - // initialize these accessors in proper way by calling __init method - // of the accessor and passing corresponding kernel parameters. - DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none); - auto Lhs = MemberExpr::Create( - S.Context, KernelObjCloneRef, false, SourceLocation(), - NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP, - DeclarationNameInfo(Field->getDeclName(), SourceLocation()), - nullptr, Field->getType(), VK_LValue, OK_Ordinary, NOUR_None); - getExprForWrappedAccessorInit(CRD, Lhs); - if (Util::isSyclStreamType(FieldType)) { - // Generate call to the __init method of the stream class after - // initializing accessors wrapped by this stream object - getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef, - InitMethodName, BodyStmts); - - // Generate call to the __finalize method of stream class. - // Will put it later to the end of function body. - getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef, - FinalizeMethodName, FinalizeStmts); - } - } - } else { - llvm_unreachable("Unsupported field type"); - } - KernelFuncParam++; - } - Expr *ILE = new (S.Context) - InitListExpr(S.Context, SourceLocation(), InitExprs, SourceLocation()); - ILE->setType(QualType(LC->getTypeForDecl(), 0)); - KernelObjClone->setInit(ILE); - } - - // In the kernel caller function kernel object is a function parameter, so we - // need to replace all refs to this kernel oject with refs to our clone - // declared inside kernel body. - Stmt *FunctionBody = KernelCallerFunc->getBody(); - - ParmVarDecl *KernelObjParam = *(KernelCallerFunc->param_begin()); - - // DeclRefExpr with valid source location but with decl which is not marked - // as used is invalid. - KernelObjClone->setIsUsed(); - std::pair MappingPair; - MappingPair.first = KernelObjParam; - MappingPair.second = KernelObjClone; - - // Function scope might be empty, so we do push - S.PushFunctionScope(); - KernelBodyTransform KBT(MappingPair, S); - Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); - BodyStmts.push_back(NewBody); - - BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(), FinalizeStmts.end()); - - return CompoundStmt::Create(S.Context, BodyStmts, SourceLocation(), - SourceLocation()); -} - /// Creates a kernel parameter descriptor /// \param Src field declaration to construct name from /// \param Ty the desired parameter type @@ -1801,15 +1548,6 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body, int_header); ConstructingOpenCLKernel = false; - - /* - //ConstructingOpenCLKernel = true; - ****CompoundStmt *OpenCLKernelBody = - CreateOpenCLKernelBody(*this, KernelCallerFunc, OpenCLKernel); - //ConstructingOpenCLKernel = false; - //OpenCLKernel->setBody(OpenCLKernelBody); - //addSyclDeviceDecl(OpenCLKernel); - */ } void Sema::MarkDevice(void) { From 948f1baf37e839d0ccdd68deff5df679304cba42 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Fri, 17 Apr 2020 09:39:31 +0300 Subject: [PATCH 29/38] Cleanups and renamings Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 114 ++++++++++++++++++------------------ 1 file changed, 57 insertions(+), 57 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b5b53088e1f0f..e2a4fff2c9970 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -867,16 +867,16 @@ template class SyclKernelFieldHandler { // Accessor can be a base class or a field decl, so both must be handled. virtual void handleSyclAccessorType(const CXXBaseSpecifier &, QualType) {} - virtual void handleSyclAccessorType(const FieldDecl *, QualType) {} - virtual void handleSyclSamplerType(const FieldDecl *, QualType) {} - virtual void handleSyclSpecConstantType(const FieldDecl *, QualType) {} + virtual void handleSyclAccessorType(FieldDecl *, QualType) {} + virtual void handleSyclSamplerType(FieldDecl *, QualType) {} + virtual void handleSyclSpecConstantType(FieldDecl *, QualType) {} virtual void handleSyclStreamType(const CXXBaseSpecifier &, QualType) {} - virtual void handleSyclStreamType(const FieldDecl *, QualType) {} - virtual void handleStructType(const FieldDecl *, QualType) {} - virtual void handleReferenceType(const FieldDecl *, QualType) {} - virtual void handlePointerType(const FieldDecl *, QualType) {} - virtual void handleArrayType(const FieldDecl *, QualType) {} - virtual void handleScalarType(const FieldDecl *, QualType) {} + virtual void handleSyclStreamType(FieldDecl *, QualType) {} + virtual void handleStructType(FieldDecl *, QualType) {} + virtual void handleReferenceType(FieldDecl *, QualType) {} + virtual void handlePointerType(FieldDecl *, QualType) {} + virtual void handleArrayType(FieldDecl *, QualType) {} + virtual void handleScalarType(FieldDecl *, QualType) {} // Most handlers shouldn't be handling this, just the field checker. virtual void handleOtherType(const FieldDecl *, QualType) {} @@ -884,7 +884,7 @@ template class SyclKernelFieldHandler { // class/field graph. Int Headers use this to calculate offset, most others // don't have a need for these. - virtual void enterStruct(const CXXRecordDecl *, const FieldDecl *) {} + virtual void enterStruct(const CXXRecordDecl *, FieldDecl *) {} virtual void leaveStruct(const CXXRecordDecl *, const FieldDecl *) {} virtual void enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} virtual void leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} @@ -903,11 +903,11 @@ class SyclKernelFieldChecker : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} bool isValid() { return !IsInvalid; } - void handleReferenceType(const FieldDecl *FD, QualType ArgTy) final { + void handleReferenceType(FieldDecl *FD, QualType ArgTy) final { IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) << ArgTy; } - void handleStructType(const FieldDecl *FD, QualType ArgTy) final { + void handleStructType(FieldDecl *FD, QualType ArgTy) final { if (SemaRef.getASTContext().getLangOpts().SYCLStdLayoutKernelParams && !ArgTy->isStandardLayoutType()) IsInvalid = @@ -929,9 +929,9 @@ class SyclKernelFieldChecker } } - // We should be able to ahndle this, so we made it part of the visitor, but + // We should be able to handle this, so we made it part of the visitor, but // this is 'to be implemented'. - void handleArrayType(const FieldDecl *FD, QualType ArgTy) final { + void handleArrayType(FieldDecl *FD, QualType ArgTy) final { IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) << ArgTy; } @@ -981,7 +981,7 @@ class SyclKernelDeclCreator // kernel parameters from __init method parameters. We will use __init method // and kernel parameters which we build here to initialize special objects in // the kernel body. - void handleSpecialType(const FieldDecl *FD, QualType ArgTy) { + void handleSpecialType(FieldDecl *FD, QualType ArgTy) { const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); @@ -1060,15 +1060,15 @@ class SyclKernelDeclCreator LastParamIndex = ParamIndex; } - void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { + void handleSyclAccessorType(FieldDecl *FD, QualType ArgTy) final { handleSpecialType(FD, ArgTy); } - void handleSyclSamplerType(const FieldDecl *FD, QualType ArgTy) final { + void handleSyclSamplerType(FieldDecl *FD, QualType ArgTy) final { handleSpecialType(FD, ArgTy); } - void handlePointerType(const FieldDecl *FD, QualType ArgTy) final { + void handlePointerType(FieldDecl *FD, QualType ArgTy) final { // TODO: Can we document what the heck this is doing?! QualType PointeeTy = ArgTy->getPointeeType(); Qualifiers Quals = PointeeTy.getQualifiers(); @@ -1079,7 +1079,7 @@ class SyclKernelDeclCreator addParam(FD, ModTy); } - void handleScalarType(const FieldDecl *FD, QualType ArgTy) final { + void handleScalarType(FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy); } @@ -1087,16 +1087,16 @@ class SyclKernelDeclCreator // object is required. The base type is pretty cheap, so we might opt // to just always create it (the way this one is implemented) and just put // this implementation in the base. - void handleStructType(const FieldDecl *FD, QualType ArgTy) final { + void handleStructType(FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy); } - void handleSyclStreamType(const FieldDecl *FD, QualType ArgTy) final { + void handleSyclStreamType(FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy); } void handleSyclStreamType(const CXXBaseSpecifier &, QualType ArgTy) final { - // TODO do something + // TODO Implement } void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } @@ -1118,7 +1118,7 @@ class SyclKernelBodyCreator VarDecl *KernelObjClone; InitializedEntity VarEntity; CXXRecordDecl *KernelObj; - llvm::SmallVector Bases; + llvm::SmallVector MemberExprBases; FunctionDecl *KernelCallerFunc; // Using the statements/init expressions that we've created, this generates @@ -1252,7 +1252,7 @@ class SyclKernelBodyCreator BodyStmts.push_back(Call); } - // TODO get rid of kernel obj clone + // TODO Remove kernel obj clone static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC, CXXRecordDecl *KernelObj) { TypeSourceInfo *TSInfo = @@ -1268,7 +1268,7 @@ class SyclKernelBodyCreator const auto *RecordDecl = Ty->getAsCXXRecordDecl(); // TODO : we don't need all this init stuff if remove kernel obj clone // Perform initialization only if it is field of kernel object - if (Bases.size() == 1) { + if (MemberExprBases.size() == 1) { InitializedEntity Entity = InitializedEntity::InitializeMember(FD, &VarEntity); // Initialize with the default constructor. @@ -1278,7 +1278,8 @@ class SyclKernelBodyCreator ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None); InitExprs.push_back(MemberInit.get()); } - createSpecialMethodCall(RecordDecl, Bases.back(), InitMethodName, FD); + createSpecialMethodCall(RecordDecl, MemberExprBases.back(), InitMethodName, + FD); } public: @@ -1290,7 +1291,6 @@ class SyclKernelBodyCreator DC.getKernelDecl(), KernelObj)), VarEntity(InitializedEntity::InitializeVariable(KernelObjClone)), KernelObj(KernelObj), KernelCallerFunc(KernelCallerFunc) { - // TODO: Something special with the lambda when InvokeParallelForWorkGroup. markParallelWorkItemCalls(); Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), @@ -1300,7 +1300,7 @@ class SyclKernelBodyCreator S.Context, NestedNameSpecifierLoc(), SourceLocation(), KernelObjClone, false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), VK_LValue); - Bases.push_back(KernelObjCloneRef); + MemberExprBases.push_back(KernelObjCloneRef); } ~SyclKernelBodyCreator() { @@ -1308,60 +1308,60 @@ class SyclKernelBodyCreator DeclCreator.setBody(KernelBody); } - void handleSyclAccessorType(const FieldDecl *FD, QualType Ty) final { - handleSpecialType(const_cast(FD), Ty); + void handleSyclAccessorType(FieldDecl *FD, QualType Ty) final { + handleSpecialType(FD, Ty); } void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { // TODO: Creates init sequence and inits special sycl obj } - void handleSyclSamplerType(const FieldDecl *FD, QualType Ty) final { - handleSpecialType(const_cast(FD), Ty); + void handleSyclSamplerType(FieldDecl *FD, QualType Ty) final { + handleSpecialType(FD, Ty); } - void handleSyclStreamType(const FieldDecl *FD, QualType Ty) final { + void handleSyclStreamType(FieldDecl *FD, QualType Ty) final { const auto *StreamDecl = Ty->getAsCXXRecordDecl(); // Hmm okay, we created "stream" kernel argument and did copy initialization // but why? // We also don't have SEMA tests for stream, only pretty small CodeGen test - createExprForStructOrScalar(const_cast(FD)); - createSpecialMethodCall(StreamDecl, Bases.back(), InitMethodName, - const_cast(FD)); - createSpecialMethodCall(StreamDecl, Bases.back(), FinalizeMethodName, - const_cast(FD)); + createExprForStructOrScalar(FD); + createSpecialMethodCall(StreamDecl, MemberExprBases.back(), InitMethodName, + FD); + createSpecialMethodCall(StreamDecl, MemberExprBases.back(), + FinalizeMethodName, FD); } void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) final { // TODO: Creates init/finalize sequence and inits special sycl obj } - void handlePointerType(const FieldDecl *FD, QualType ArgTy) final { - createExprForStructOrScalar(const_cast(FD)); + void handlePointerType(FieldDecl *FD, QualType ArgTy) final { + createExprForStructOrScalar(FD); } - void handleStructType(const FieldDecl *FD, QualType Ty) final { - createExprForStructOrScalar(const_cast(FD)); + void handleStructType(FieldDecl *FD, QualType Ty) final { + createExprForStructOrScalar(FD); } - void handleScalarType(const FieldDecl *FD, QualType Ty) final { - createExprForStructOrScalar(const_cast(FD)); + void handleScalarType(FieldDecl *FD, QualType Ty) final { + createExprForStructOrScalar(FD); } - void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { - Bases.push_back(BuildMemberExpr(Bases.back(), const_cast(FD))); + void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { + MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); } void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { - Bases.pop_back(); + MemberExprBases.pop_back(); } void enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { - // TODO : do something here? + // TODO : Implement } void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { - // TODO : do something here? + // TODO : Implement } }; @@ -1412,7 +1412,7 @@ class SyclKernelIntHeaderCreator getOffset(BC.getType()->getAsCXXRecordDecl())); } - void handleSyclAccessorType(const FieldDecl *FD, QualType ArgTy) final { + void handleSyclAccessorType(FieldDecl *FD, QualType ArgTy) final { const auto *AccTy = cast(ArgTy->getAsRecordDecl()); assert(AccTy->getTemplateArgs().size() >= 2 && @@ -1424,7 +1424,7 @@ class SyclKernelIntHeaderCreator getOffset(FD)); } - void handleSyclSamplerType(const FieldDecl *FD, QualType ArgTy) final { + void handleSyclSamplerType(FieldDecl *FD, QualType ArgTy) final { const auto *SamplerTy = ArgTy->getAsCXXRecordDecl(); assert(SamplerTy && "Sampler type must be a C++ record type"); CXXMethodDecl *InitMethod = getMethodByName(SamplerTy, InitMethodName); @@ -1437,7 +1437,7 @@ class SyclKernelIntHeaderCreator addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler); } - void handleSyclSpecConstantType(const FieldDecl *FD, QualType ArgTy) final { + void handleSyclSpecConstantType(FieldDecl *FD, QualType ArgTy) final { const TemplateArgumentList &TemplateArgs = cast(ArgTy->getAsRecordDecl()) ->getTemplateInstantiationArgs(); @@ -1455,17 +1455,17 @@ class SyclKernelIntHeaderCreator Header.addSpecConstant(SpecConstName, SpecConstIDTy); } - void handlePointerType(const FieldDecl *FD, QualType ArgTy) final { + void handlePointerType(FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy, SYCLIntegrationHeader::kind_pointer); } - void handleStructType(const FieldDecl *FD, QualType ArgTy) final { + void handleStructType(FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); } - void handleScalarType(const FieldDecl *FD, QualType ArgTy) final { + void handleScalarType(FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); } - void handleSyclStreamType(const FieldDecl *FD, QualType ArgTy) final { + void handleSyclStreamType(FieldDecl *FD, QualType ArgTy) final { addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); } void handleSyclStreamType(const CXXBaseSpecifier &BC, QualType ArgTy) final { @@ -1473,7 +1473,7 @@ class SyclKernelIntHeaderCreator } // Keep track of the current struct offset. - void enterStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; } From cec1a358a58dbd9107dd3ed908b8d23155936322 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 20 Apr 2020 11:34:02 +0300 Subject: [PATCH 30/38] Apply comments, fixes Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 175 ++++++++++++++++++------------------ 1 file changed, 85 insertions(+), 90 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 35f983350a359..693bf31c9c64c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -648,9 +648,6 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { // kernel itself. static QualType calculateKernelNameType(ASTContext &Ctx, FunctionDecl *KernelCallerFunc) { - // TODO: Not sure what the 'fully qualified type's purpose is here, the type - // itself should have its full qualified name, so figure out what the purpose - // is. const TemplateArgumentList *TAL = KernelCallerFunc->getTemplateSpecializationArgs(); return TypeName::getFullyQualifiedType(TAL->get(0).getAsType(), Ctx, @@ -705,8 +702,10 @@ static void VisitAccessorWrapperHelper(CXXRecordDecl *Owner, RangeTy Range, } } -// poorly named Parent is the 'how we got here', basically just enough info for -// the offset adjustment to know what to do about the enter-struct info. +// Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter +// the Wrapper structure that we're currently visiting. Owner is the parent type +// (which doesn't exist in cases where it is a FieldDecl in the 'root'), and +// Wrapper is the current struct being unwrapped. template static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent, CXXRecordDecl *Wrapper, @@ -785,18 +784,16 @@ template class SyclKernelFieldHandler { virtual void handleArrayType(FieldDecl *, QualType) {} virtual void handleScalarType(FieldDecl *, QualType) {} // Most handlers shouldn't be handling this, just the field checker. - virtual void handleOtherType(const FieldDecl *, QualType) {} + virtual void handleOtherType(FieldDecl *, QualType) {} // The following are only used for keeping track of where we are in the base // class/field graph. Int Headers use this to calculate offset, most others // don't have a need for these. virtual void enterStruct(const CXXRecordDecl *, FieldDecl *) {} - virtual void leaveStruct(const CXXRecordDecl *, const FieldDecl *) {} + virtual void leaveStruct(const CXXRecordDecl *, FieldDecl *) {} virtual void enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} virtual void leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &) {} - // virtual void enterStruct(const FieldDecl *, CXXRecordDecl *Struct); - // virtual void leaveStruct(const FieldDecl *, CXXRecordDecl *Struct); }; // A type to check the valididty of all of the argument types. @@ -810,42 +807,42 @@ class SyclKernelFieldChecker : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} bool isValid() { return !IsInvalid; } - void handleReferenceType(FieldDecl *FD, QualType ArgTy) final { + void handleReferenceType(FieldDecl *FD, QualType FieldTy) final { IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) - << ArgTy; + << FieldTy; } - void handleStructType(FieldDecl *FD, QualType ArgTy) final { + void handleStructType(FieldDecl *FD, QualType FieldTy) final { if (SemaRef.getASTContext().getLangOpts().SYCLStdLayoutKernelParams && - !ArgTy->isStandardLayoutType()) + !FieldTy->isStandardLayoutType()) IsInvalid = Diag.Report(FD->getLocation(), diag::err_sycl_non_std_layout_type) - << ArgTy; + << FieldTy; else { - CXXRecordDecl *RD = ArgTy->getAsCXXRecordDecl(); + CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl(); if (!RD->hasTrivialCopyConstructor()) IsInvalid = Diag.Report(FD->getLocation(), diag::err_sycl_non_trivially_copy_ctor_dtor_type) - << 0 << ArgTy; + << 0 << FieldTy; else if (!RD->hasTrivialDestructor()) IsInvalid = Diag.Report(FD->getLocation(), diag::err_sycl_non_trivially_copy_ctor_dtor_type) - << 1 << ArgTy; + << 1 << FieldTy; } } // We should be able to handle this, so we made it part of the visitor, but // this is 'to be implemented'. - void handleArrayType(FieldDecl *FD, QualType ArgTy) final { + void handleArrayType(FieldDecl *FD, QualType FieldTy) final { IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) - << ArgTy; + << FieldTy; } - void handleOtherType(const FieldDecl *FD, QualType ArgTy) final { + void handleOtherType(FieldDecl *FD, QualType FieldTy) final { IsInvalid = Diag.Report(FD->getLocation(), diag::err_bad_kernel_param_type) - << ArgTy; + << FieldTy; } }; @@ -860,17 +857,18 @@ class SyclKernelDeclCreator // iterator as push_back invalidates iterators. size_t LastParamIndex = 0; - void addParam(const FieldDecl *FD, QualType ArgTy) { - ParamDesc newParamDesc = makeParamDesc(FD, ArgTy); - addParam(newParamDesc, ArgTy); + void addParam(const FieldDecl *FD, QualType FieldTy) { + ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); + addParam(newParamDesc, FieldTy); } - void addParam(const CXXBaseSpecifier &BS, QualType ArgTy) { - ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), BS, ArgTy); - addParam(newParamDesc, ArgTy); + void addParam(const CXXBaseSpecifier &BS, QualType FieldTy) { + ParamDesc newParamDesc = + makeParamDesc(SemaRef.getASTContext(), BS, FieldTy); + addParam(newParamDesc, FieldTy); } - void addParam(ParamDesc newParamDesc, QualType ArgTy) { + void addParam(ParamDesc newParamDesc, QualType FieldTy) { // Create a new ParmVarDecl based on the new info. auto *NewParam = ParmVarDecl::Create( SemaRef.getASTContext(), KernelDecl, SourceLocation(), SourceLocation(), @@ -888,8 +886,8 @@ class SyclKernelDeclCreator // kernel parameters from __init method parameters. We will use __init method // and kernel parameters which we build here to initialize special objects in // the kernel body. - void handleSpecialType(FieldDecl *FD, QualType ArgTy) { - const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); + void handleSpecialType(FieldDecl *FD, QualType FieldTy) { + const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method"); @@ -953,8 +951,8 @@ class SyclKernelDeclCreator } void handleSyclAccessorType(const CXXBaseSpecifier &BS, - QualType ArgTy) final { - const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); + QualType FieldTy) final { + const auto *RecordDecl = FieldTy->getAsCXXRecordDecl(); assert(RecordDecl && "The accessor/sampler must be a RecordDecl"); CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, InitMethodName); assert(InitMethod && "The accessor/sampler must have the __init method"); @@ -967,17 +965,20 @@ class SyclKernelDeclCreator LastParamIndex = ParamIndex; } - void handleSyclAccessorType(FieldDecl *FD, QualType ArgTy) final { - handleSpecialType(FD, ArgTy); + void handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { + handleSpecialType(FD, FieldTy); } - void handleSyclSamplerType(FieldDecl *FD, QualType ArgTy) final { - handleSpecialType(FD, ArgTy); + void handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { + handleSpecialType(FD, FieldTy); } - void handlePointerType(FieldDecl *FD, QualType ArgTy) final { - // TODO: Can we document what the heck this is doing?! - QualType PointeeTy = ArgTy->getPointeeType(); + void handlePointerType(FieldDecl *FD, QualType FieldTy) final { + // USM allows to use raw pointers instead of buffers/accessors, but these + // pointers point to the specially allocated memory. For pointer fields we + // add a kernel argument with the same type as field but global address + // space, because OpenCL requires it. + QualType PointeeTy = FieldTy->getPointeeType(); Qualifiers Quals = PointeeTy.getQualifiers(); Quals.setAddressSpace(LangAS::opencl_global); PointeeTy = SemaRef.getASTContext().getQualifiedType( @@ -986,24 +987,21 @@ class SyclKernelDeclCreator addParam(FD, ModTy); } - void handleScalarType(FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy); + void handleScalarType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); } - // This is implemented here because this is the only case where the recurse - // object is required. The base type is pretty cheap, so we might opt - // to just always create it (the way this one is implemented) and just put - // this implementation in the base. - void handleStructType(FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy); + void handleStructType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); } - void handleSyclStreamType(FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy); + void handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); } - void handleSyclStreamType(const CXXBaseSpecifier &, QualType ArgTy) final { - // TODO Implement + void handleSyclStreamType(const CXXBaseSpecifier &, QualType FieldTy) final { + // FIXME SYCL stream should be usable usable as a base type + // See https://github.com/intel/llvm/issues/1552 } void setBody(CompoundStmt *KB) { KernelDecl->setBody(KB); } @@ -1044,11 +1042,10 @@ class SyclKernelBodyCreator // DeclRefExpr with valid source location but with decl which is not marked // as used is invalid. KernelObjClone->setIsUsed(); - std::pair MappingPair; - MappingPair.first = KernelObjParam; - MappingPair.second = KernelObjClone; + std::pair MappingPair = + std::make_pair(KernelObjParam, KernelObjClone); - // Function scope might be empty, so we do push + // Function scope might be empty, so we do push. SemaRef.PushFunctionScope(); KernelBodyTransform KBT(MappingPair, SemaRef); Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); @@ -1159,7 +1156,8 @@ class SyclKernelBodyCreator BodyStmts.push_back(Call); } - // TODO Remove kernel obj clone + // FIXME Avoid creation of kernel obj clone. + // See https://github.com/intel/llvm/issues/1544 for details. static VarDecl *createKernelObjClone(ASTContext &Ctx, DeclContext *DC, CXXRecordDecl *KernelObj) { TypeSourceInfo *TSInfo = @@ -1173,7 +1171,6 @@ class SyclKernelBodyCreator void handleSpecialType(FieldDecl *FD, QualType Ty) { const auto *RecordDecl = Ty->getAsCXXRecordDecl(); - // TODO : we don't need all this init stuff if remove kernel obj clone // Perform initialization only if it is field of kernel object if (MemberExprBases.size() == 1) { InitializedEntity Entity = @@ -1220,7 +1217,8 @@ class SyclKernelBodyCreator } void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { - // TODO: Creates init sequence and inits special sycl obj + // FIXME SYCL accessor should be usable usable as a base type + // See https://github.com/intel/llvm/issues/28. } void handleSyclSamplerType(FieldDecl *FD, QualType Ty) final { @@ -1229,9 +1227,6 @@ class SyclKernelBodyCreator void handleSyclStreamType(FieldDecl *FD, QualType Ty) final { const auto *StreamDecl = Ty->getAsCXXRecordDecl(); - // Hmm okay, we created "stream" kernel argument and did copy initialization - // but why? - // We also don't have SEMA tests for stream, only pretty small CodeGen test createExprForStructOrScalar(FD); createSpecialMethodCall(StreamDecl, MemberExprBases.back(), InitMethodName, FD); @@ -1240,18 +1235,19 @@ class SyclKernelBodyCreator } void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) final { - // TODO: Creates init/finalize sequence and inits special sycl obj + // FIXME SYCL stream should be usable usable as a base type + // See https://github.com/intel/llvm/issues/1552 } - void handlePointerType(FieldDecl *FD, QualType ArgTy) final { + void handlePointerType(FieldDecl *FD, QualType FieldTy) final { createExprForStructOrScalar(FD); } - void handleStructType(FieldDecl *FD, QualType Ty) final { + void handleStructType(FieldDecl *FD, QualType FieldTy) final { createExprForStructOrScalar(FD); } - void handleScalarType(FieldDecl *FD, QualType Ty) final { + void handleScalarType(FieldDecl *FD, QualType FieldTy) final { createExprForStructOrScalar(FD); } @@ -1259,7 +1255,7 @@ class SyclKernelBodyCreator MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD)); } - void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + void leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { MemberExprBases.pop_back(); } @@ -1283,15 +1279,13 @@ class SyclKernelIntHeaderCreator return 0; } uint64_t getOffset(const FieldDecl *FD) const { - // TODO: Figure out how to calc lower down the structs, currently only gives - // the 'base' value. return CurOffset + SemaRef.getASTContext().getFieldOffset(FD) / 8; } - void addParam(const FieldDecl *FD, QualType ArgTy, + void addParam(const FieldDecl *FD, QualType FieldTy, SYCLIntegrationHeader::kernel_param_kind_t Kind) { uint64_t Size = - SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); + SemaRef.getASTContext().getTypeSizeInChars(FieldTy).getQuantity(); Header.addParamDesc(Kind, static_cast(Size), static_cast(getOffset(FD))); } @@ -1306,22 +1300,21 @@ class SyclKernelIntHeaderCreator } void handleSyclAccessorType(const CXXBaseSpecifier &BC, - QualType ArgTy) final { + QualType FieldTy) final { const auto *AccTy = - cast(ArgTy->getAsRecordDecl()); + cast(FieldTy->getAsRecordDecl()); assert(AccTy->getTemplateArgs().size() >= 2 && "Incorrect template args for Accessor Type"); int Dims = static_cast( AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue()); int Info = getAccessTarget(AccTy) | (Dims << 11); Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, - // TODO: is this the right way? getOffset(BC.getType()->getAsCXXRecordDecl())); } - void handleSyclAccessorType(FieldDecl *FD, QualType ArgTy) final { + void handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { const auto *AccTy = - cast(ArgTy->getAsRecordDecl()); + cast(FieldTy->getAsRecordDecl()); assert(AccTy->getTemplateArgs().size() >= 2 && "Incorrect template args for Accessor Type"); int Dims = static_cast( @@ -1331,8 +1324,8 @@ class SyclKernelIntHeaderCreator getOffset(FD)); } - void handleSyclSamplerType(FieldDecl *FD, QualType ArgTy) final { - const auto *SamplerTy = ArgTy->getAsCXXRecordDecl(); + void handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { + const auto *SamplerTy = FieldTy->getAsCXXRecordDecl(); assert(SamplerTy && "Sampler type must be a C++ record type"); CXXMethodDecl *InitMethod = getMethodByName(SamplerTy, InitMethodName); assert(InitMethod && "sampler must have __init method"); @@ -1344,9 +1337,9 @@ class SyclKernelIntHeaderCreator addParam(FD, SamplerArg->getType(), SYCLIntegrationHeader::kind_sampler); } - void handleSyclSpecConstantType(FieldDecl *FD, QualType ArgTy) final { + void handleSyclSpecConstantType(FieldDecl *FD, QualType FieldTy) final { const TemplateArgumentList &TemplateArgs = - cast(ArgTy->getAsRecordDecl()) + cast(FieldTy->getAsRecordDecl()) ->getTemplateInstantiationArgs(); assert(TemplateArgs.size() == 2 && "Incorrect template args for Accessor Type"); @@ -1362,21 +1355,23 @@ class SyclKernelIntHeaderCreator Header.addSpecConstant(SpecConstName, SpecConstIDTy); } - void handlePointerType(FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy, SYCLIntegrationHeader::kind_pointer); + void handlePointerType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_pointer); } - void handleStructType(FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); + void handleStructType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); } - void handleScalarType(FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); + void handleScalarType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); } - void handleSyclStreamType(FieldDecl *FD, QualType ArgTy) final { - addParam(FD, ArgTy, SYCLIntegrationHeader::kind_std_layout); + void handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout); } - void handleSyclStreamType(const CXXBaseSpecifier &BC, QualType ArgTy) final { - // TODO implement + void handleSyclStreamType(const CXXBaseSpecifier &BC, + QualType FieldTy) final { + // FIXME SYCL stream should be usable usable as a base type + // See https://github.com/intel/llvm/issues/1552 } // Keep track of the current struct offset. @@ -1384,7 +1379,7 @@ class SyclKernelIntHeaderCreator CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; } - void leaveStruct(const CXXRecordDecl *, const FieldDecl *FD) final { + void leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { CurOffset -= SemaRef.getASTContext().getFieldOffset(FD) / 8; } From 447fd703f83d251f90358559f4d69060b2a1dd3c Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 20 Apr 2020 15:04:40 +0300 Subject: [PATCH 31/38] Apply CR suggestions Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 15 +++------------ 1 file changed, 3 insertions(+), 12 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 693bf31c9c64c..8280de79e4702 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -767,10 +767,6 @@ template class SyclKernelFieldHandler { // Mark these virutal so that we can use override in the implementer classes, // despite virtual dispatch never being used. - //// TODO: Can these return 'bool' and we can short-circuit the handling? That - // way the field checker cna return true/false based on whether the rest - // should be still working. - // Accessor can be a base class or a field decl, so both must be handled. virtual void handleSyclAccessorType(const CXXBaseSpecifier &, QualType) {} virtual void handleSyclAccessorType(FieldDecl *, QualType) {} @@ -1045,7 +1041,7 @@ class SyclKernelBodyCreator std::pair MappingPair = std::make_pair(KernelObjParam, KernelObjClone); - // Function scope might be empty, so we do push. + // Push the Kernel function scope to ensure the scope isn't empty SemaRef.PushFunctionScope(); KernelBodyTransform KBT(MappingPair, SemaRef); Stmt *NewBody = KBT.TransformStmt(FunctionBody).get(); @@ -1259,13 +1255,8 @@ class SyclKernelBodyCreator MemberExprBases.pop_back(); } - void enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { - // TODO : Implement - } - - void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { - // TODO : Implement - } + using SyclKernelFieldHandler::enterStruct; + using SyclKernelFieldHandler::leaveStruct; }; class SyclKernelIntHeaderCreator From 1acc00b74f530d1890d1e3bd78871f5470fa88cc Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 20 Apr 2020 08:10:26 -0700 Subject: [PATCH 32/38] Implement base class offsets for int-headers Implement the base class offsets as required for the integration headers. Note the test invocation had to be changed to skip code generation, since the other stages don't play well with base classes. Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 16 +++++++-- clang/test/CodeGenSYCL/integration_header.cpp | 36 ++++++++++++++++++- 2 files changed, 48 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8280de79e4702..ee8854317dce9 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1263,11 +1263,17 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { SYCLIntegrationHeader &Header; const CXXRecordDecl *KernelLambda; + // Necessary to figure out the offset of the base class. + const CXXRecordDecl *CurStruct = nullptr; int64_t CurOffset = 0; uint64_t getOffset(const CXXRecordDecl *RD) const { // TODO: Figure this out! Offset of a base class. - return 0; + assert(CurOffset && + "Cannot have a base class without setting the active struct"); + const ASTRecordLayout &Layout = + SemaRef.getASTContext().getASTRecordLayout(CurStruct); + return CurOffset + Layout.getBaseClassOffset(RD).getQuantity(); } uint64_t getOffset(const FieldDecl *FD) const { return CurOffset + SemaRef.getASTContext().getFieldOffset(FD) / 8; @@ -1366,15 +1372,18 @@ class SyclKernelIntHeaderCreator } // Keep track of the current struct offset. - void enterStruct(const CXXRecordDecl *, FieldDecl *FD) final { + void enterStruct(const CXXRecordDecl * RD, FieldDecl *FD) final { + CurStruct = FD->getType()->getAsCXXRecordDecl(); CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; } - void leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final { + void leaveStruct(const CXXRecordDecl * RD, FieldDecl *FD) final { + CurStruct = RD; CurOffset -= SemaRef.getASTContext().getFieldOffset(FD) / 8; } void enterStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + CurStruct = BS.getType()->getAsCXXRecordDecl(); const ASTRecordLayout &Layout = SemaRef.getASTContext().getASTRecordLayout(RD); CurOffset += Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl()) @@ -1382,6 +1391,7 @@ class SyclKernelIntHeaderCreator } void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final { + CurStruct = RD; const ASTRecordLayout &Layout = SemaRef.getASTContext().getASTRecordLayout(RD); CurOffset -= Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl()) diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index ad788c646c8e8..4af3b441b3a3b 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -1,4 +1,4 @@ -// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o %T/kernel.spv +// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -fsyntax-only -o %T/kernel.spv // RUN: FileCheck -input-file=%t.h %s // // CHECK: #include @@ -21,6 +21,7 @@ // CHECK-NEXT: "_ZTSN16second_namespace13second_kernelIcEE", // CHECK-NEXT: "_ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE" // CHECK-NEXT: "_ZTS13fourth_kernelIJN15template_arg_ns14namespaced_argILi1EEEEE" +// CHECK-NEXT: "_ZTSZ4mainE16accessor_in_base" // CHECK-NEXT: }; // // CHECK: static constexpr @@ -45,6 +46,13 @@ // CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 4, 0 }, // CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 6112, 4 }, // CHECK-EMPTY: +// CHECK-NEXT: //--- _ZTSZ4mainE16accessor_in_base +// CHECK-NEXT: { kernel_param_kind_t::kind_std_layout, 64, 0 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 8 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 24 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 40 }, +// CHECK-NEXT: { kernel_param_kind_t::kind_accessor, 4062, 52 }, +// CHECK-EMPTY: // CHECK-NEXT: }; // // CHECK: template <> struct KernelInfo { @@ -77,6 +85,26 @@ struct namespaced_arg {}; template class fourth_kernel; +namespace accessor_in_base { + struct other_base{int i;}; + struct base { + int i, j; + cl::sycl::accessor acc; + }; + + struct base2 : other_base, + cl::sycl::accessor { + int i; + cl::sycl::accessor acc; + }; + + struct captured : base, base2 { + cl::sycl::accessor acc; + void use() const{} + }; + +}; + int main() { cl::sycl::accessor acc1; @@ -121,5 +149,11 @@ int main() { } }); + // FIXME: We cannot use the member-capture because non-integration headers + // don't handle these types right. + accessor_in_base::captured c; + kernel_single_task([c]() { + }); + return 0; } From 71c2b0c622734abfe928bf4e9f40c60369b95ab5 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Mon, 20 Apr 2020 18:11:15 +0300 Subject: [PATCH 33/38] Apply comments & clang-format Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 20 +++++------ clang/test/CodeGenSYCL/integration_header.cpp | 36 ++++++++++--------- 2 files changed, 28 insertions(+), 28 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ee8854317dce9..3ee9b498be5e6 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1061,14 +1061,13 @@ class SyclKernelBodyCreator assert(WGLambdaFn && "PFWG lambda not found"); // Mark the function that it "works" in a work group scope: // NOTE: In case of parallel_for_work_item the marker call itself is - // marked - // with work item scope attribute, here the '()' operator of the - // object passed as parameter is marked. This is an optimization - - // there are a lot of locals created at parallel_for_work_group - // scope before calling the lambda - it is more efficient to have - // all of them in the private address space rather then sharing via - // the local AS. See parallel_for_work_group implementation in the - // SYCL headers. + // marked with work item scope attribute, here the '()' operator of the + // object passed as parameter is marked. This is an optimization - + // there are a lot of locals created at parallel_for_work_group + // scope before calling the lambda - it is more efficient to have + // all of them in the private address space rather then sharing via + // the local AS. See parallel_for_work_group implementation in the + // SYCL headers. if (!WGLambdaFn->hasAttr()) { WGLambdaFn->addAttr(SYCLScopeAttr::CreateImplicit( SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); @@ -1268,7 +1267,6 @@ class SyclKernelIntHeaderCreator int64_t CurOffset = 0; uint64_t getOffset(const CXXRecordDecl *RD) const { - // TODO: Figure this out! Offset of a base class. assert(CurOffset && "Cannot have a base class without setting the active struct"); const ASTRecordLayout &Layout = @@ -1372,12 +1370,12 @@ class SyclKernelIntHeaderCreator } // Keep track of the current struct offset. - void enterStruct(const CXXRecordDecl * RD, FieldDecl *FD) final { + void enterStruct(const CXXRecordDecl *RD, FieldDecl *FD) final { CurStruct = FD->getType()->getAsCXXRecordDecl(); CurOffset += SemaRef.getASTContext().getFieldOffset(FD) / 8; } - void leaveStruct(const CXXRecordDecl * RD, FieldDecl *FD) final { + void leaveStruct(const CXXRecordDecl *RD, FieldDecl *FD) final { CurStruct = RD; CurOffset -= SemaRef.getASTContext().getFieldOffset(FD) / 8; } diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index 4af3b441b3a3b..1b2b8f551232e 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -86,25 +86,27 @@ template class fourth_kernel; namespace accessor_in_base { - struct other_base{int i;}; - struct base { - int i, j; - cl::sycl::accessor acc; - }; - - struct base2 : other_base, - cl::sycl::accessor { - int i; - cl::sycl::accessor acc; - }; +struct other_base { + int i; +}; +struct base { + int i, j; + cl::sycl::accessor acc; +}; - struct captured : base, base2 { - cl::sycl::accessor acc; - void use() const{} - }; +struct base2 : other_base, + cl::sycl::accessor { + int i; + cl::sycl::accessor acc; +}; +struct captured : base, base2 { + cl::sycl::accessor acc; + void use() const {} }; +}; // namespace accessor_in_base + int main() { cl::sycl::accessor acc1; @@ -149,8 +151,8 @@ int main() { } }); - // FIXME: We cannot use the member-capture because non-integration headers - // don't handle these types right. + // FIXME: We cannot use the member-capture because all the handlers except the + // integration header handler in SemaSYCL don't handle base types right. accessor_in_base::captured c; kernel_single_task([c]() { }); From a964a25be67f965266c91acd816577fae31ba4d5 Mon Sep 17 00:00:00 2001 From: Erich Keane Date: Mon, 20 Apr 2020 12:15:03 -0700 Subject: [PATCH 34/38] Refactor constructKernelName to return a pair. Since the only caller to this now uses both values, refactor it to just return a pair and unpack on the other side. Signed-off-by: Erich Keane --- clang/lib/Sema/SemaSYCL.cpp | 21 ++++++++++----------- 1 file changed, 10 insertions(+), 11 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3ee9b498be5e6..7e94980eeb59b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -656,20 +656,20 @@ static QualType calculateKernelNameType(ASTContext &Ctx, // Gets a name for the kernel caller func, calculated from the first template // argument. -static std::string constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, - MangleContext &MC, bool StableName) { +static std::pair +constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, MangleContext &MC) { QualType KernelNameType = calculateKernelNameType(S.getASTContext(), KernelCallerFunc); - if (StableName) - return PredefinedExpr::ComputeName(S.getASTContext(), - PredefinedExpr::UniqueStableNameType, - KernelNameType); SmallString<256> Result; llvm::raw_svector_ostream Out(Result); MC.mangleTypeName(KernelNameType, Out); - return std::string(Out.str()); + + return {std::string(Out.str()), + PredefinedExpr::ComputeName(S.getASTContext(), + PredefinedExpr::UniqueStableNameType, + KernelNameType)}; } // anonymous namespace so these don't get linkage. @@ -1427,10 +1427,9 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, assert(KernelLambda && "invalid kernel caller"); // Calculate both names, since Integration headers need both. - std::string CalculatedName = - constructKernelName(*this, KernelCallerFunc, MC, /*StableName*/ false); - std::string StableName = - constructKernelName(*this, KernelCallerFunc, MC, /*StableName*/ true); + std::string CalculatedName, StableName; + std::tie(CalculatedName, StableName) = + constructKernelName(*this, KernelCallerFunc, MC); StringRef KernelName(getLangOpts().SYCLUnnamedLambda ? StableName : CalculatedName); From 85687868dca0f04ad71f28c9d1e0551eac802fc1 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 21 Apr 2020 09:30:32 +0300 Subject: [PATCH 35/38] Simplify run line of the test Signed-off-by: Mariya Podchishchaeva --- clang/test/CodeGenSYCL/integration_header.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index 1b2b8f551232e..e3e2bb66bc755 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -1,4 +1,4 @@ -// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -fsyntax-only -o %T/kernel.spv +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s -fsyntax-only // RUN: FileCheck -input-file=%t.h %s // // CHECK: #include From f4ab328e2e856921e4991e2141424fc1eeb012bf Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 21 Apr 2020 09:37:01 +0300 Subject: [PATCH 36/38] Apply CR comments Signed-off-by: Mariya Podchishchaeva --- clang/lib/Sema/SemaSYCL.cpp | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7e94980eeb59b..2145bfb3bf02d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -644,18 +644,19 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) { AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue()); } -// The first template argument to the kernel function is used to identify the -// kernel itself. +// The first template argument to the kernel caller function is used to identify +// the kernel itself. static QualType calculateKernelNameType(ASTContext &Ctx, FunctionDecl *KernelCallerFunc) { const TemplateArgumentList *TAL = KernelCallerFunc->getTemplateSpecializationArgs(); + assert(TAL && "No template argument info"); return TypeName::getFullyQualifiedType(TAL->get(0).getAsType(), Ctx, /*WithGlobalNSPrefix=*/true); } -// Gets a name for the kernel caller func, calculated from the first template -// argument. +// Gets a name for the OpenCL kernel function, calculated from the first +// template argument of the kernel caller function. static std::pair constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, MangleContext &MC) { QualType KernelNameType = @@ -996,7 +997,7 @@ class SyclKernelDeclCreator } void handleSyclStreamType(const CXXBaseSpecifier &, QualType FieldTy) final { - // FIXME SYCL stream should be usable usable as a base type + // FIXME SYCL stream should be usable as a base type // See https://github.com/intel/llvm/issues/1552 } @@ -1212,7 +1213,7 @@ class SyclKernelBodyCreator } void handleSyclAccessorType(const CXXBaseSpecifier &BS, QualType Ty) final { - // FIXME SYCL accessor should be usable usable as a base type + // FIXME SYCL accessor should be usable as a base type // See https://github.com/intel/llvm/issues/28. } @@ -1230,7 +1231,7 @@ class SyclKernelBodyCreator } void handleSyclStreamType(const CXXBaseSpecifier &BS, QualType Ty) final { - // FIXME SYCL stream should be usable usable as a base type + // FIXME SYCL stream should be usable as a base type // See https://github.com/intel/llvm/issues/1552 } @@ -1365,7 +1366,7 @@ class SyclKernelIntHeaderCreator } void handleSyclStreamType(const CXXBaseSpecifier &BC, QualType FieldTy) final { - // FIXME SYCL stream should be usable usable as a base type + // FIXME SYCL stream should be usable as a base type // See https://github.com/intel/llvm/issues/1552 } From 645774bba818d2e473c46a0ac52ea60db98d5510 Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Tue, 21 Apr 2020 09:58:28 +0300 Subject: [PATCH 37/38] Apply clang-format Signed-off-by: Mariya Podchishchaeva --- 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 2145bfb3bf02d..b3ecda585cf3e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -658,7 +658,8 @@ static QualType calculateKernelNameType(ASTContext &Ctx, // Gets a name for the OpenCL kernel function, calculated from the first // template argument of the kernel caller function. static std::pair -constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, MangleContext &MC) { +constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc, + MangleContext &MC) { QualType KernelNameType = calculateKernelNameType(S.getASTContext(), KernelCallerFunc); From b0a03e62f787d6ff2b2a61a337658c7a2669363b Mon Sep 17 00:00:00 2001 From: Mariya Podchishchaeva Date: Wed, 22 Apr 2020 09:08:27 +0300 Subject: [PATCH 38/38] Try to fix LIT test on windows Signed-off-by: Mariya Podchishchaeva --- clang/test/CodeGenSYCL/integration_header.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index e3e2bb66bc755..84b35578f48e6 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s -fsyntax-only +// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only // RUN: FileCheck -input-file=%t.h %s // // CHECK: #include