From a11308ac24c669e40af906ed2eafa98cb8e5f1e9 Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Fri, 28 Dec 2018 23:46:36 +0300 Subject: [PATCH] [SYCL] SYCL integration header emission by device compiler Signed-off-by: Vladimir Lazarev --- clang/include/clang/Basic/LangOptions.h | 3 + clang/include/clang/Driver/CC1Options.td | 5 + clang/include/clang/Sema/Sema.h | 107 ++++- clang/lib/Frontend/CompilerInvocation.cpp | 2 + clang/lib/Sema/Sema.cpp | 8 +- clang/lib/Sema/SemaSYCL.cpp | 491 ++++++++++++++++++---- 6 files changed, 530 insertions(+), 86 deletions(-) diff --git a/clang/include/clang/Basic/LangOptions.h b/clang/include/clang/Basic/LangOptions.h index 9cff7c5160438..5a27c84ba7209 100644 --- a/clang/include/clang/Basic/LangOptions.h +++ b/clang/include/clang/Basic/LangOptions.h @@ -250,6 +250,9 @@ class LangOptions : public LangOptionsBase { /// input is a header file (i.e. -x c-header). bool IsHeaderFile = false; + /// SYCL integration header to be generated by the device compiler + std::string SYCLIntHeader; + LangOptions(); // Define accessors/mutators for language options of enumeration type. diff --git a/clang/include/clang/Driver/CC1Options.td b/clang/include/clang/Driver/CC1Options.td index 291b8eaaaf620..4202afbb19243 100644 --- a/clang/include/clang/Driver/CC1Options.td +++ b/clang/include/clang/Driver/CC1Options.td @@ -842,6 +842,11 @@ def fopenmp_host_ir_file_path : Separate<["-"], "fopenmp-host-ir-file-path">, def fsycl_is_device : Flag<["-"], "fsycl-is-device">, HelpText<"Generate code for SYCL device.">; +def fsycl_int_header : Separate<["-"], "fsycl-int-header">, + HelpText<"Generate SYCL integration header into this file.">; +def fsycl_int_header_EQ : Joined<["-"], "fsycl-int-header=">, + Alias; + } // let Flags = [CC1Option] //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 1d2f68aee70e0..dd20824ed918c 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -274,6 +274,101 @@ class FileNullabilityMap { } }; +// TODO SYCL Integration header approach relies on an assumption that kernel +// lambda objects created by the host compiler and any of the device compilers +// will be identical wrt to field types, order and offsets. Some verification +// mechanism should be developed to enforce that. + +// TODO FIXME SYCL Support for SYCL in FE should be refactored: +// - kernel indentification and generation should be made a separate pass over +// AST. RecursiveASTVisitor + VisitFunctionTemplateDecl + +// FunctionTemplateDecl::getSpecializations() mechanism could be used for that. +// - All SYCL stuff on Sema level should be encapsulated into a single Sema +// field +// - Move SYCL stuff into a separate header + +// Represents contents of a SYCL integration header file produced by a SYCL +// device compiler and used by SYCL host compiler (via forced inclusion into +// compiled SYCL source): +// - SYCL kernel names +// - SYCL kernel parameters and offsets of corresponding actual arguments +class SYCLIntegrationHeader { +public: + // Kind of kernel's lambda parameters as captured by the compiler in the + // kernel lambda object + enum kernel_param_kind_t { + kind_first, + kind_none = kind_first, + kind_accessor, + kind_scalar, + kind_struct, + kind_sampler, + kind_struct_padding, // can be added by the compiler to enforce alignment + kind_last = kind_struct_padding + }; + +public: + SYCLIntegrationHeader(); + + /// Emits contents of the header into given stream. + void emit(raw_ostream &Out); + + /// Emits contents of the header into a file with given name. + /// Returns true/false on success/failure. + bool emit(const StringRef &MainSrc); + + /// Signals that subsequent parameter descriptor additions will go to + /// the kernel with given name. Starts new kernel invocation descriptor. + void startKernel(StringRef KernelName); + + /// Adds a kernel parameter descriptor to current kernel invocation + /// descriptor. + void addParamDesc(kernel_param_kind_t Kind, int Info, unsigned Offset); + + /// Signals that addition of parameter descriptors to current kernel + /// invocation descriptor has finished. + void endKernel(); + +private: + // Kernel actual parameter descriptor. + struct KernelParamDesc { + // Represents a parameter kind. + kernel_param_kind_t Kind; + // If Kind is kind_scalar or kind_struct, then + // denotes parameter size in bytes (includes padding for structs) + // If Kind is kind_accessor + // denotes access target; possible access targets are defined in + // access/access.hpp + int Info; + // Offset of the captured parameter value in the lambda or function object. + unsigned Offset; + + KernelParamDesc() = default; + }; + + // Kernel invocation descriptor + struct KernelDesc { + /// Kernel name. + std::string Name; + /// Descriptor of kernel actual parameters. + SmallVector Params; + + KernelDesc() = default; + }; + + /// Returns the latest invocation descriptor started by + /// SYCLIntegrationHeader::startKernel + KernelDesc *getCurKernelDesc() { + return KernelDescs.size() > 0 ? &KernelDescs[KernelDescs.size() - 1] + : nullptr; + } + +private: + /// Keeps invocation descriptors for each kernel invocation started by + /// SYCLIntegrationHeader::startKernel + SmallVector KernelDescs; +}; + /// Sema - This implements semantic analysis and AST building for C. class Sema { Sema(const Sema &) = delete; @@ -10847,12 +10942,22 @@ class Sema { // We store SYCL Kernels here and handle separately -- which is a hack. // FIXME: It would be best to refactor this. SmallVector SyclKernel; + // SYCL integratrion header instance for current compilation unit this Sema + // is associated with. + std::unique_ptr SyclIntHeader; public: void AddSyclKernel(Decl * d) { SyclKernel.push_back(d); } SmallVector &SyclKernels() { return SyclKernel; } - void ConstructSYCLKernel(FunctionDecl* KernelHelper); + /// Lazily creates and returns SYCL integratrion header instance. + SYCLIntegrationHeader &getSyclIntegrationHeader() { + if (SyclIntHeader == nullptr) + SyclIntHeader = llvm::make_unique(); + return *SyclIntHeader.get(); + } + + void ConstructSYCLKernel(FunctionDecl *KernelCallerFunc); }; /// RAII object that enters a new expression evaluation context. diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index 9703ddaf35293..0dd85a9f7c14d 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2398,6 +2398,8 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, Opts.GPURelocatableDeviceCode = Args.hasArg(OPT_fgpu_rdc); + Opts.SYCLIntHeader = Args.getLastArgValue(OPT_fsycl_int_header); + if (Opts.ObjC) { if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) { StringRef value = arg->getValue(); diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 9fa39968625a6..7762811486d19 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -145,7 +145,8 @@ Sema::Sema(Preprocessor &pp, ASTContext &ctxt, ASTConsumer &consumer, CurrentInstantiationScope(nullptr), DisableTypoCorrection(false), TyposCorrected(0), AnalysisWarnings(*this), ThreadSafetyDeclCache(nullptr), VarDataSharingAttributesStack(nullptr), - CurScope(nullptr), Ident_super(nullptr), Ident___float128(nullptr) { + CurScope(nullptr), Ident_super(nullptr), Ident___float128(nullptr), + SyclIntHeader(nullptr) { TUScope = nullptr; LoadedExternalKnownNamespaces = false; @@ -916,6 +917,11 @@ void Sema::ActOnEndOfTranslationUnit() { PerformPendingInstantiations(); + // Emit SYCL integration header for current translation unit if needed + if (getLangOpts().SYCL && SyclIntHeader != nullptr) { + SyclIntHeader->emit(getLangOpts().SYCLIntHeader); + } + assert(LateParsedInstantiations.empty() && "end of TU template instantiation should not create more " "late-parsed templates"); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 29297dfc62c2a..a4ac7e033a99d 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -11,14 +11,19 @@ #include "TreeTransform.h" #include "clang/AST/AST.h" +#include "clang/AST/RecordLayout.h" #include "clang/AST/RecursiveASTVisitor.h" #include "clang/Sema/Sema.h" #include "llvm/ADT/SmallVector.h" +#include "llvm/Support/FileSystem.h" +#include "llvm/Support/Path.h" // used in integration header creation using namespace clang; typedef llvm::DenseMap DeclMap; +using KernelParamKind = SYCLIntegrationHeader::kernel_param_kind_t; + enum target { global_buffer = 2014, constant_buffer, @@ -29,6 +34,14 @@ enum target { image_array }; +static CXXRecordDecl *getKernelCallerLambdaArg(FunctionDecl *FD) { + auto FirstArg = (*FD->param_begin()); + if (FirstArg) + if (FirstArg->getType()->getAsCXXRecordDecl()->isLambda()) + return FirstArg->getType()->getAsCXXRecordDecl(); + return nullptr; +} + class MarkDeviceFunction : public RecursiveASTVisitor { public: MarkDeviceFunction(Sema &S) @@ -84,17 +97,10 @@ class KernelBodyTransform : public TreeTransform { Sema &SemaRef; }; -CXXRecordDecl *getBodyAsLambda(FunctionDecl *FD) { - auto FirstArg = (*FD->param_begin()); - if (FirstArg) - if (FirstArg->getType()->getAsCXXRecordDecl()->isLambda()) - return FirstArg->getType()->getAsCXXRecordDecl(); - return nullptr; -} - -FunctionDecl *CreateSYCLKernelFunction(ASTContext &Context, StringRef Name, - ArrayRef ArgTys, - ArrayRef ArgDecls) { +static FunctionDecl * +CreateSYCLKernelFunction(ASTContext &Context, StringRef Name, + ArrayRef ArgTys, + ArrayRef ArgDecls) { DeclContext *DC = Context.getTranslationUnitDecl(); FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); @@ -125,14 +131,14 @@ FunctionDecl *CreateSYCLKernelFunction(ASTContext &Context, StringRef Name, return SYCLKernel; } -CompoundStmt *CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelHelper, - DeclContext *DC) { +static CompoundStmt * +CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) { llvm::SmallVector BodyStmts; // TODO: case when kernel is functor // TODO: possible refactoring when functor case will be completed - CXXRecordDecl *LC = getBodyAsLambda(KernelHelper); + CXXRecordDecl *LC = getKernelCallerLambdaArg(KernelCallerFunc); if (LC) { // Create Lambda object auto LambdaVD = VarDecl::Create( @@ -231,9 +237,10 @@ CompoundStmt *CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelHelper, // to replace all refs to this lambda with our vardecl. // I used TreeTransform here, but I'm not sure that it is good solution // Also I used map and I'm not sure about it too. - Stmt *FunctionBody = KernelHelper->getBody(); + // TODO SYCL review the above design concerns + Stmt *FunctionBody = KernelCallerFunc->getBody(); DeclMap DMap; - ParmVarDecl *LambdaParam = *(KernelHelper->param_begin()); + ParmVarDecl *LambdaParam = *(KernelCallerFunc->param_begin()); // DeclRefExpr with valid source location but with decl which is not marked // as used is invalid. LambdaVD->setIsUsed(); @@ -248,85 +255,205 @@ CompoundStmt *CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelHelper, SourceLocation()); } -void BuildArgTys(ASTContext &Context, - llvm::SmallVector &ArgDecls, - llvm::SmallVector &NewArgDecls, - llvm::SmallVector &ArgTys) { - for (auto V : ArgDecls) { +/// Various utilities. +class Util { +public: + // TODO SYCL use AST infrastructure instead of string matching + + /// Checks whether given clang type is a sycl accessor class. + static bool isSyclAccessorType(QualType Ty) { + std::string Name = Ty.getCanonicalType().getAsString(); + return Name.find("class cl::sycl::accessor") != std::string::npos; + } + + /// Checks whether given clang type is a sycl stream class. + static bool isSyclStreamType(QualType Ty) { + std::string Name = Ty.getCanonicalType().getAsString(); + return Name == "stream"; + } +}; + +/// Identifies context of kernel lambda capture visitor function +/// invocation. +enum VisitorContext { + pre_visit, + visit_accessor, + visit_scalar, + visit_stream, + post_visit, +}; + +/// Implements visitor design pattern for lambda captures. +/// +/// Iterates over captured parameters of given lambda and invokes given +/// visitor functions at appropriate context providing information of interest. +/// \param Lambda the kernel lambda object +/// \param Vis a tuple of visitor functions, each corresponds to and is +/// invoked at a specific context. @see VisitorContext. +/// +template +static void visitKernelLambdaCaptures(const CXXRecordDecl *Lambda, + VisitorTupleTy &Vis) { + const LambdaCapture *Cpt = Lambda->captures_begin(); + RecordDecl::field_iterator Fld = Lambda->field_begin(); + const LambdaCapture *CptEnd = Lambda->captures_end(); + const RecordDecl::field_iterator FldEnd = Lambda->field_end(); + + for (; (Cpt != CptEnd) && (Fld != FldEnd); Cpt++, Fld++) { + // pre-visit context + unsigned Cnt = static_cast(std::distance(Cpt, CptEnd)); + VarDecl *V = Cpt->getCapturedVar(); QualType ArgTy = V->getType(); - QualType ActualArgType = ArgTy; - std::string Name = ArgTy.getCanonicalType().getAsString(); - if (Name.find("class cl::sycl::accessor") != std::string::npos) { - if (const auto *RecordDecl = ArgTy->getAsCXXRecordDecl()) { - const auto *TemplateDecl = - dyn_cast(RecordDecl); - if (TemplateDecl) { - // First parameter - data type - QualType PointeeType = TemplateDecl->getTemplateArgs()[0].getAsType(); - // Fourth parameter - access target - auto AccessQualifier = - TemplateDecl->getTemplateArgs()[3].getAsIntegral(); - int64_t AccessTarget = AccessQualifier.getExtValue(); - Qualifiers Quals = PointeeType.getQualifiers(); - // TODO: Support all access targets - switch (AccessTarget) { - case target::global_buffer: - Quals.setAddressSpace(LangAS::opencl_global); - break; - case target::constant_buffer: - Quals.setAddressSpace(LangAS::opencl_constant); - break; - case target::local: - Quals.setAddressSpace(LangAS::opencl_local); - break; - default: - llvm_unreachable("Unsupported access target"); - } - // TODO: get address space from accessor template parameter. - PointeeType = - Context.getQualifiedType(PointeeType.getUnqualifiedType(), Quals); - QualType PointerType = Context.getPointerType(PointeeType); - ActualArgType = - Context.getQualifiedType(PointerType.getUnqualifiedType(), Quals); - } - } - } else if (std::string(Name) == "stream") { - continue; + auto F1 = std::get(Vis); + F1(Cnt, V, *Fld); + + if (Util::isSyclAccessorType(ArgTy)) { + // accessor parameter context + const auto *RecordDecl = ArgTy->getAsCXXRecordDecl(); + assert(RecordDecl && "accessor must be of a record type"); + const auto *TemplateDecl = + dyn_cast(RecordDecl); + assert(TemplateDecl && "templated accessor type expected"); + + // First accessor template parameter - data type + QualType PointeeType = TemplateDecl->getTemplateArgs()[0].getAsType(); + // Fourth parameter - access target + auto AccessQualifier = TemplateDecl->getTemplateArgs()[3].getAsIntegral(); + int64_t AccessTarget = AccessQualifier.getExtValue(); + auto F = std::get(Vis); + F(Cnt, static_cast(AccessTarget), PointeeType, V, *Fld); + } else if (Util::isSyclStreamType(ArgTy)) { + // stream parameter context + auto F = std::get(Vis); + F(Cnt, V, *Fld); + } else if (ArgTy->isScalarType()) { + // scalar typed parameter context + auto F = std::get(Vis); + F(Cnt, V, *Fld); + } else { + llvm_unreachable("unsupported kernel parameter type"); } - DeclContext *DC = Context.getTranslationUnitDecl(); - - IdentifierInfo *VarName = 0; - SmallString<8> Str; - llvm::raw_svector_ostream OS(Str); - OS << "_arg_" << V->getIdentifier()->getName(); - VarName = &Context.Idents.get(OS.str()); - - auto NewVarDecl = VarDecl::Create( - Context, DC, SourceLocation(), SourceLocation(), VarName, ActualArgType, - Context.getTrivialTypeSourceInfo(ActualArgType), SC_None); - ArgTys.push_back(ActualArgType); - NewArgDecls.push_back(NewVarDecl); + // pos-visit context + auto F2 = std::get(Vis); + F2(Cnt, V, *Fld); } + assert((Cpt == CptEnd) && (Fld == FldEnd) && + "captures inconsistent with fields"); } -void Sema::ConstructSYCLKernel(FunctionDecl *KernelHelper) { - // TODO: Case when kernel is functor - CXXRecordDecl *LE = getBodyAsLambda(KernelHelper); - if (LE) { - - llvm::SmallVector ArgDecls; +static void BuildArgTys(ASTContext &Context, CXXRecordDecl *Lambda, + llvm::SmallVector &NewArgDecls, + llvm::SmallVector &ArgTys) { + QualType ActualArgType; // serves to transfer info between visitor lambdas + auto Vis = std::make_tuple( + // pre_visit + [&](int, VarDecl *, FieldDecl *) {}, + // visit_accessor + [&](int CaptureN, target AccTrg, QualType PointeeType, + DeclaratorDecl *CapturedVar, FieldDecl *CapturedVal) { + Qualifiers Quals = PointeeType.getQualifiers(); + // TODO: Support all access targets + switch (AccTrg) { + case target::global_buffer: + Quals.setAddressSpace(LangAS::opencl_global); + break; + case target::constant_buffer: + Quals.setAddressSpace(LangAS::opencl_constant); + break; + case target::local: + Quals.setAddressSpace(LangAS::opencl_local); + break; + default: + llvm_unreachable("Unsupported access target"); + } + // TODO: get address space from accessor template parameter. + PointeeType = + Context.getQualifiedType(PointeeType.getUnqualifiedType(), Quals); + QualType PointerType = Context.getPointerType(PointeeType); + ActualArgType = + Context.getQualifiedType(PointerType.getUnqualifiedType(), Quals); + }, + // visit_scalar + [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { + ActualArgType = CapturedVal->getType(); + }, + // visit_stream + [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { + llvm_unreachable("streams not supported yet"); + }, + // post_visit + [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { + DeclContext *DC = Context.getTranslationUnitDecl(); + + IdentifierInfo *VarName = 0; + SmallString<8> Str; + llvm::raw_svector_ostream OS(Str); + OS << "_arg_" << CapturedVar->getIdentifier()->getName(); + VarName = &Context.Idents.get(OS.str()); + + auto NewVarDecl = VarDecl::Create( + Context, DC, SourceLocation(), SourceLocation(), VarName, + ActualArgType, Context.getTrivialTypeSourceInfo(ActualArgType), + SC_None); + ArgTys.push_back(ActualArgType); + NewArgDecls.push_back(NewVarDecl); + }); + visitKernelLambdaCaptures(Lambda, Vis); +} - for (const auto &V : LE->captures()) { - ArgDecls.push_back(V.getCapturedVar()); - } +/// Adds necessary data describing given kernel to the integration header. +/// \param H the integration header object +/// \param Name kernel name +/// \param Lambda kernel lambda object +static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, + CXXRecordDecl *Lambda) { + ASTContext &Ctx = Lambda->getASTContext(); + const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(Lambda); + KernelParamKind Knd = SYCLIntegrationHeader::kind_none; + H.startKernel(Name); + unsigned Offset = 0; + int Info = 0; + + auto Vis = std::make_tuple( + // pre_visit + [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { + Offset = static_cast( + Layout.getFieldOffset(CapturedVal->getFieldIndex())); + }, + // visit_accessor + [&](int CaptureN, target AccTrg, QualType PointeeType, + DeclaratorDecl *CapturedVar, FieldDecl *CapturedVal) { + Knd = SYCLIntegrationHeader::kind_accessor; + Info = static_cast(AccTrg); + }, + // visit_scalar + [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { + Knd = SYCLIntegrationHeader::kind_scalar; + Info = static_cast( + Ctx.getTypeSizeInChars(CapturedVal->getType()).getQuantity()); + }, + // visit_stream + [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { + llvm_unreachable("streams not supported yet"); + }, + // post_visit + [&](int CaptureN, VarDecl *CapturedVar, FieldDecl *CapturedVal) { + H.addParamDesc(Knd, Info, Offset); + }); + visitKernelLambdaCaptures(Lambda, Vis); +} +void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) { + // TODO: Case when kernel is functor + CXXRecordDecl *LE = getKernelCallerLambdaArg(KernelCallerFunc); + if (LE) { llvm::SmallVector ArgTys; llvm::SmallVector NewArgDecls; - BuildArgTys(getASTContext(), ArgDecls, NewArgDecls, ArgTys); + BuildArgTys(getASTContext(), LE, NewArgDecls, ArgTys); // Get Name for our kernel. const TemplateArgumentList *TemplateArgs = - KernelHelper->getTemplateSpecializationArgs(); + KernelCallerFunc->getTemplateSpecializationArgs(); QualType KernelNameType = TemplateArgs->get(0).getAsType(); std::string Name = KernelNameType.getBaseTypeIdentifier()->getName().str(); @@ -345,12 +472,13 @@ void Sema::ConstructSYCLKernel(FunctionDecl *KernelHelper) { Name.erase(pos, ToBeErased[i].length()); } } + populateIntHeader(getSyclIntegrationHeader(), Name, LE); FunctionDecl *SYCLKernel = CreateSYCLKernelFunction(getASTContext(), Name, ArgTys, NewArgDecls); CompoundStmt *SYCLKernelBody = - CreateSYCLKernelBody(*this, KernelHelper, SYCLKernel); + CreateSYCLKernelBody(*this, KernelCallerFunc, SYCLKernel); SYCLKernel->setBody(SYCLKernelBody); AddSyclKernel(SYCLKernel); @@ -360,3 +488,198 @@ void Sema::ConstructSYCLKernel(FunctionDecl *KernelHelper) { Marker.TraverseStmt(SYCLKernelBody); } } + +// ----------------------------------------------------------------------------- +// Integration header functionality implementation +// ----------------------------------------------------------------------------- + +/// Returns a string ID of given parameter kind - used in header +/// emission. +static const char *paramKind2Str(KernelParamKind K) { +#define CASE(x) \ + case SYCLIntegrationHeader::kind_##x: \ + return "kind_" #x + switch (K) { + CASE(none); + CASE(accessor); + CASE(scalar); + CASE(struct); + CASE(sampler); + CASE(struct_padding); + default: + return ""; + } +#undef CASE +} + +// // Integration header structure: +// +// // kernel parameter kinds +// enum kernel_param_kind_t { +// kind_none, +// kind_accessor, +// kind_scalar, +// kind_struct, +// kind_sampler, +// kind_struct_padding +// }; +// +// // names of all kernels defined in the corresponding source +// const char* kernel_names[] = { +// "SimpleVadd1", +// "SimpleVadd2" +// }; +// +// // describes a kernel parameter +// struct kernel_param_desc_t { +// // parameter kind +// kernel_param_kind_t kind; +// // kind == kind_scalar, kind_struct +// // parameter size in bytes (includes padding for structs) +// // kind == kind_accessor +// // access target; possible access targets are defined in +// // access/access.hpp +// int info; +// // offset of the captured value of the parameter in the lambda or function +// // object +// int offs; +// }; +// +// // array representing signatures of all kernels defined in the +// // corresponding source +// kernel_param_desc_t kernel_signatures[] = { +// // SimpleVadd1 +// { kind_accessor, 0, 0 }, // accessorC +// { kind_accessor, 0, 64 }, // accessorA +// { kind_accessor, 0, 128 }, // accessorB +// { kind_scalar, 4, 132 }, // param +// { kind_none, 0, 0 }, // terminator +// // SimpleVadd2 +// { kind_accessor, 0, 0 }, // accessorC +// { kind_scalar, 4, 68 }, // param +// { kind_none, 0, 0 } // terminator +// }; +// +// // indices into the kernel_signatures array, each representing a start of +// // kernel signature descriptor subarray of the kernel_signature array; +// // the index order in this array corresponds to the kernel name order in the +// // kernel_names array +// unsigned kernel_signature_start[] = { +// 0, // SimpleVadd1 +// 5 // SimpleVadd2 +// }; +// +void SYCLIntegrationHeader::emit(raw_ostream &O) { + O << "// kernel parameter kinds\n"; + O << "enum kernel_param_kind_t {\n"; + + for (int I = SYCLIntegrationHeader::kind_first; + I <= SYCLIntegrationHeader::kind_last; I++) { + KernelParamKind It = static_cast(I); + O << " " << std::string(paramKind2Str(It)); + if (I < SYCLIntegrationHeader::kind_last) + O << ","; + O << "\n"; + } + O << "};\n"; + O << "\n"; + O << "// names of all kernels defined in the corresponding source\n"; + O << "const char* kernel_names[] = {\n"; + + for (unsigned I = 0; I < KernelDescs.size(); I++) { + O << " \"" << KernelDescs[I].Name << "\""; + + if (I < KernelDescs.size() - 1) + O << ","; + O << "\n"; + } + O << "};\n\n"; + + O << "// describes a kernel parameter\n"; + O << "struct kernel_param_desc_t {\n"; + O << " // parameter kind\n"; + O << " kernel_param_kind_t kind;\n"; + O << " // kind == kind_scalar, kind_struct\n"; + O << " // parameter size in bytes (includes padding for structs)\n"; + O << " // kind == kind_accessor\n"; + O << " // access target; possible access targets are defined in " + "access/access.hpp\n"; + O << " int info;\n"; + O << " // offset of the captured value of the parameter in the lambda or " + "function object\n"; + O << " int offs;\n"; + O << "};\n\n"; + + O << "// array representing signatures of all kernels defined in the\n"; + O << "// corresponding source\n"; + O << "kernel_param_desc_t kernel_signatures[] = {\n"; + + for (unsigned I = 0; I < KernelDescs.size(); I++) { + auto &K = KernelDescs[I]; + O << " //--- " << K.Name << "\n"; + + for (const auto &P : K.Params) { + std::string TyStr = paramKind2Str(P.Kind); + O << " { " << TyStr << ", " << P.Info << ", " << P.Offset << " },\n"; + } + O << " { kind_none, 0, 0 }"; + if (I < KernelDescs.size() - 1) + O << ","; + O << "\n"; + } + O << "};\n\n"; + + O << "// indices into the kernel_signatures array, each representing a start" + " of\n"; + O << "// kernel signature descriptor subarray of the kernel_signatures" + " array;\n"; + O << "// the index order in this array corresponds to the kernel name order" + " in the\n"; + O << "// kernel_names array\n"; + O << "unsigned kernel_signature_start[] = {\n"; + unsigned CurStart = 0; + + for (unsigned I = 0; I < KernelDescs.size(); I++) { + auto &K = KernelDescs[I]; + O << " " << CurStart; + if (I < KernelDescs.size() - 1) + O << ","; + O << " // " << K.Name << "\n"; + CurStart += K.Params.size() + 1; + } + O << "};\n\n"; +} + +bool SYCLIntegrationHeader::emit(const StringRef &IntHeaderName) { + if (IntHeaderName.empty()) + return false; + int IntHeaderFD = 0; + std::error_code EC = + llvm::sys::fs::openFileForWrite(IntHeaderName, IntHeaderFD); + if (EC) { + llvm::errs() << "Error: " << EC.message() << "\n"; + // compilation will fail on absent include file - don't need to fail here + return false; + } + llvm::raw_fd_ostream Out(IntHeaderFD, true /*close in destructor*/); + emit(Out); + return true; +} + +void SYCLIntegrationHeader::startKernel(StringRef KernelName) { + KernelDescs.resize(KernelDescs.size() + 1); + KernelDescs.back().Name = KernelName; +} + +void SYCLIntegrationHeader::addParamDesc(kernel_param_kind_t Kind, int Info, + unsigned Offset) { + auto *K = getCurKernelDesc(); + assert(K && "no kernels"); + K->Params.push_back(KernelParamDesc{Kind, Info, Offset}); +} + +void SYCLIntegrationHeader::endKernel() { + // nop for now +} + +SYCLIntegrationHeader::SYCLIntegrationHeader() {}