From f8f4915e118d534791e2c9b887a61afecd327ffc Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Tue, 11 May 2021 12:17:00 -0700 Subject: [PATCH 1/5] [SYCL] Imrprove output emitted in opt-report This patch attempts to make the output generated in opt-report more user-friendly by adding a description for openCL kernel arguments. It also removes compiler generated names such as '__wrapper_class'. Signed-off-by: Elizabeth Andrews --- .../clang/Basic/SyclOptReportHandler.h | 16 +- clang/lib/CodeGen/CodeGenFunction.cpp | 58 ++- clang/lib/Sema/SemaSYCL.cpp | 198 +++++++- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 421 +++++++++++++++--- 4 files changed, 604 insertions(+), 89 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index af3ba5ae1c0a8..bdac3c0eecbf8 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -28,18 +28,26 @@ class SyclOptReportHandler { std::string KernelArgName; std::string KernelArgType; SourceLocation KernelArgLoc; + unsigned KernelArgSize; + std::string KernelArgDesc; + std::string KernelArgParent; OptReportInfo(std::string ArgName, std::string ArgType, - SourceLocation ArgLoc) + SourceLocation ArgLoc, unsigned ArgSize, std::string ArgDesc, + std::string ArgParent) : KernelArgName(std::move(ArgName)), KernelArgType(std::move(ArgType)), - KernelArgLoc(ArgLoc) {} + KernelArgLoc(ArgLoc), KernelArgSize(ArgSize), + KernelArgDesc(std::move(ArgDesc)), + KernelArgParent(std::move(ArgParent)) {} }; llvm::DenseMap> Map; public: void AddKernelArgs(const FunctionDecl *FD, std::string ArgName, - std::string ArgType, SourceLocation ArgLoc) { - Map[FD].emplace_back(ArgName, ArgType, ArgLoc); + std::string ArgType, SourceLocation ArgLoc, + unsigned ArgSize, std::string ArgDesc, + std::string ArgParent) { + Map[FD].emplace_back(ArgName, ArgType, ArgLoc, ArgSize, ArgDesc, ArgParent); } SmallVector &GetInfo(const FunctionDecl *FD) { auto It = Map.find(FD); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index aabfb649174b1..5f5acddac4154 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1472,6 +1472,33 @@ QualType CodeGenFunction::BuildFunctionArgList(GlobalDecl GD, return ResTy; } +enum class ArgDescEncoding { + BaseClass, + DecomposedMember, + WrappedPointer, + WrappedArray, + Accessor, + AccessorBase, + Sampler, + Stream, + KernelHandler, + None +}; + +static ArgDescEncoding getArgDescEncoding(std::string KernelArgDesc) { + return llvm::StringSwitch(KernelArgDesc) + .Case("base class", ArgDescEncoding::BaseClass) + .Case("decomposed struct/class", ArgDescEncoding::DecomposedMember) + .Case("nested pointer", ArgDescEncoding::WrappedPointer) + .Case("array", ArgDescEncoding::WrappedArray) + .Case("accessor", ArgDescEncoding::Accessor) + .Case("accessor base class", ArgDescEncoding::AccessorBase) + .Case("sampler", ArgDescEncoding::Sampler) + .Case("stream", ArgDescEncoding::Stream) + .Case("SYCL2020 specialization constant", ArgDescEncoding::KernelHandler) + .Default(ArgDescEncoding::None); +} + void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, const CGFunctionInfo &FnInfo) { const FunctionDecl *FD = cast(GD.getDecl()); @@ -1529,14 +1556,33 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, for (auto ORI : llvm::enumerate(OptReportHandler.GetInfo(FD))) { llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.value().KernelArgLoc); - std::string KAN = ORI.value().KernelArgName; + std::string ArgName = ORI.value().KernelArgName; + std::string ArgType = ORI.value().KernelArgType; + std::string ArgDesc = ORI.value().KernelArgDesc; + unsigned ArgSize = ORI.value().KernelArgSize; + std::string ArgParent = ORI.value().KernelArgParent; + + ArgDescEncoding ArgDescEnc = getArgDescEncoding(ArgDesc); + bool isWrappedField = (ArgDescEnc == ArgDescEncoding::WrappedPointer || + ArgDescEnc == ArgDescEncoding::WrappedArray) + ? true + : false; + llvm::OptimizationRemark Remark("sycl", "Region", DL, &Fn->getEntryBlock()); - Remark << "Argument " << llvm::ore::NV("Argument", ORI.index()) - << " for function kernel: " - << llvm::ore::NV(KAN.empty() ? "&" : "") << " " << Fn->getName() - << "." << llvm::ore::NV(KAN.empty() ? " " : KAN) << "(" - << ORI.value().KernelArgType << ")"; + Remark << "Arg " << llvm::ore::NV("Argument", ORI.index()) << ":" + << ((ArgDescEnc != ArgDescEncoding::None) + ? ("Compiler generated argument for " + ArgDesc + ",") + : "") + << ((ArgDescEnc == ArgDescEncoding::DecomposedMember) ? ArgParent + : ArgName) + << " (" + << ((ArgDescEnc == ArgDescEncoding::DecomposedMember) + ? ("Field:" + ArgName + ", ") + : "") + << "Type:" << ((isWrappedField) ? "Compiler generated" : ArgType) + << ", " + << "Size: " << llvm::ore::NV("Argument", ArgSize) << ")"; ORE.emit(Remark); } } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b0459d97c01f0..9c7297d16bcb2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1768,9 +1768,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *FD, QualType FieldTy) { ParamDesc newParamDesc = makeParamDesc(FD, FieldTy); - SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( - KernelDecl, FD->getName().data(), FieldTy.getAsString(), - FD->getLocation()); addParam(newParamDesc, FieldTy); } @@ -1781,8 +1778,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { StringRef Name = "_arg__base"; ParamDesc newParamDesc = makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); - SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( - KernelDecl, "", FieldTy.getAsString(), BS.getBaseTypeLoc()); addParam(newParamDesc, FieldTy); } // Add a parameter with specified name and type @@ -2230,6 +2225,193 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { using SyclKernelFieldHandler::handleSyclHalfType; }; +enum class KernelArgDescription { + BaseClass, + DecomposedMember, + WrappedPointer, + WrappedArray, + Accessor, + AccessorBase, + Sampler, + Stream, + KernelHandler, + None +}; + +StringRef getKernelArgDesc(KernelArgDescription Desc) { + switch (Desc) { + case KernelArgDescription::BaseClass: + return "base class"; + break; + case KernelArgDescription::DecomposedMember: + return "decomposed struct/class"; + break; + case KernelArgDescription::WrappedPointer: + return "nested pointer"; + break; + case KernelArgDescription::WrappedArray: + return "array"; + break; + case KernelArgDescription::Accessor: + return "accessor"; + break; + case KernelArgDescription::AccessorBase: + return "accessor base class"; + break; + case KernelArgDescription::Sampler: + return "sampler"; + break; + case KernelArgDescription::Stream: + return "stream"; + break; + case KernelArgDescription::KernelHandler: + return "SYCL2020 specialization constant"; + break; + default: + return ""; + } +} + +class SyclOptReportCreator : public SyclKernelFieldHandler { + SyclKernelDeclCreator &DC; + SourceLocation KernelInvocationLoc; + + void addParam(const FieldDecl *KernelArg, QualType KernelArgType, + KernelArgDescription KernelArgDesc) { + unsigned KernelArgSize = + SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); + const RecordDecl *KernelArgParent = KernelArg->getParent(); + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( + DC.getKernelDecl(), KernelArg->getName().data(), + KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, + getKernelArgDesc(KernelArgDesc).data(), + KernelArgParent ? KernelArgParent->getName().data() : ""); + } + + void addParam(const FieldDecl *FD, QualType FieldTy) { + KernelArgDescription Desc = KernelArgDescription::None; + const RecordDecl *RD = FD->getParent(); + if (RD && RD->hasAttr()) + Desc = KernelArgDescription::DecomposedMember; + + addParam(FD, FieldTy, Desc); + } + + void addParam(QualType KernelArgType, KernelArgDescription KernelArgDesc) { + unsigned KernelArgSize = + SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( + DC.getKernelDecl(), KernelArgType.getAsString(), + KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, + getKernelArgDesc(KernelArgDesc).data(), ""); + } + + // Handles SYCL special types - accessor, sampler and stream + // Also handles modified types - arrays and pointers + bool handleSpecialType(const FieldDecl *FD, QualType FieldTy, + KernelArgDescription Desc) { + llvm::ArrayRef KernelParameters = + DC.getParamVarDeclsForCurrentField(); + for (auto *Param : KernelParameters) + addParam(FD, Param->getType(), Desc); + return true; + } + +public: + static constexpr const bool VisitInsideSimpleContainers = false; + SyclOptReportCreator(Sema &S, SyclKernelDeclCreator &DC, SourceLocation Loc) + : SyclKernelFieldHandler(S), DC(DC), KernelInvocationLoc(Loc) {} + + bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final { + return handleSpecialType( + FD, FieldTy, KernelArgDescription(KernelArgDescription::Accessor)); + } + + bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS, + QualType FieldTy) final { + llvm::ArrayRef KernelParameters = + DC.getParamVarDeclsForCurrentField(); + for (auto *Param : KernelParameters) + addParam(Param->getType(), + KernelArgDescription(KernelArgDescription::AccessorBase)); + return true; + } + + bool handleSyclSamplerType(FieldDecl *FD, QualType FieldTy) final { + return handleSpecialType( + FD, FieldTy, KernelArgDescription(KernelArgDescription::Sampler)); + } + + bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { + KernelArgDescription Desc = KernelArgDescription::None; + ParmVarDecl *KernelParameter = DC.getParamVarDeclsForCurrentField()[0]; + // Compiler generated openCL kernel argument for current pointer field + // is not a pointer. This means we are processing a nested pointer and + // the openCL kernel argument is of type __wrapper_class. + if (!KernelParameter->getType()->isPointerType()) + Desc = KernelArgDescription::WrappedPointer; + return handleSpecialType(FD, FieldTy, Desc); + } + + bool handleScalarType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); + return true; + } + + bool handleSimpleArrayType(FieldDecl *FD, QualType FieldTy) final { + // Simple arrays are always wrapped. + handleSpecialType(FD, FieldTy, + KernelArgDescription(KernelArgDescription::WrappedArray)); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *, FieldDecl *FD, + QualType Ty) final { + addParam(FD, Ty); + return true; + } + + bool handleNonDecompStruct(const CXXRecordDecl *Base, + const CXXBaseSpecifier &BS, QualType Ty) final { + addParam(Ty, KernelArgDescription(KernelArgDescription::BaseClass)); + return true; + } + + bool handleUnionType(FieldDecl *FD, QualType FieldTy) final { + return handleScalarType(FD, FieldTy); + } + + bool handleSyclHalfType(FieldDecl *FD, QualType FieldTy) final { + addParam(FD, FieldTy); + return true; + } + + bool handleSyclStreamType(FieldDecl *FD, QualType FieldTy) final { + // For the current implementation of stream class, the Visitor 'handles' + // stream argument and then visits each accessor field in stream. Therefore + // handleSpecialType in this case only adds a single argument for stream. + // The arguments corresponding to accessors in stream are handled in + // handleSyclAccessorType. The opt-report therefore does not diffrentiate + // between the accessors in streams and accessors captured by SYCL kernel. + // Once stream API is modified to use __init(), the visitor will no longer + // visit the stream object and opt-report output for stream class will be + // similar to that of other special types. + return handleSpecialType( + FD, FieldTy, KernelArgDescription(KernelArgDescription::Stream)); + } + + void handleSyclKernelHandlerType() { + ASTContext &Context = SemaRef.getASTContext(); + if (isDefaultSPIRArch(Context)) + return; + addParam(DC.getParamVarDeclsForCurrentField()[0]->getType(), + KernelArgDescription(KernelArgDescription::KernelHandler)); + } + using SyclKernelFieldHandler::handleSyclHalfType; + using SyclKernelFieldHandler::handleSyclSamplerType; + using SyclKernelFieldHandler::handleSyclStreamType; +}; + static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) { for (auto *MD : Rec->methods()) { if (MD->getOverloadedOperator() == OO_Call) @@ -3563,18 +3745,20 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, StableName, KernelCallerFunc); SyclKernelIntFooterCreator int_footer(*this, getSyclIntegrationFooter()); + SyclOptReportCreator opt_report(*this, kernel_decl, KernelObj->getLocation()); KernelObjVisitor Visitor{*this}; Visitor.VisitRecordBases(KernelObj, kernel_decl, kernel_body, int_header, - int_footer); + int_footer, opt_report); Visitor.VisitRecordFields(KernelObj, kernel_decl, kernel_body, int_header, - int_footer); + int_footer, opt_report); if (ParmVarDecl *KernelHandlerArg = getSyclKernelHandlerArg(KernelCallerFunc)) { kernel_decl.handleSyclKernelHandlerType(); kernel_body.handleSyclKernelHandlerType(KernelHandlerArg); int_header.handleSyclKernelHandlerType(KernelHandlerArg->getType()); + opt_report.handleSyclKernelHandlerType(); } } diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp index fe71db3e833e3..760f1a9614a80 100644 --- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -6,85 +6,26 @@ #include "Inputs/sycl.hpp" -class second_base { -public: - int *e; -}; +sycl::handler H; -class InnerFieldBase { +class decomposedbase { public: - int d; -}; -class InnerField : public InnerFieldBase { - int c; + float decompvar; + int *decompptr; + sycl::accessor decompAcc; + sycl::stream decompStream{0, 0, H}; }; -struct base { +struct notdecomposedbase { public: int b; - InnerField obj; }; -//CHECK: --- !Passed -//CHECK: Pass:{{.*}}sycl -//CHECK: Name:{{.*}}Region -//CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -//CHECK: Line: 85, Column: 18 } -//CHECK: Function: _ZTS7derived -//CHECK: Args: -//CHECK-NEXT: String: 'Argument ' -//CHECK-NEXT: Argument: '0' -//CHECK-NEXT: String: ' for function kernel: ' -//CHECK-NEXT: String: '&' -//CHECK-NEXT: String: ' ' -//CHECK-NEXT: String: _ZTS7derived -//CHECK-NEXT: String: . -//CHECK-NEXT: String: ' ' -//CHECK-NEXT: String: '(' -//CHECK-NEXT: String: struct base -//CHECK-NEXT: String: ')' - -//CHECK: --- !Passed -//CHECK: Pass:{{.*}}sycl -//CHECK: Name:{{.*}}Region -//CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -//CHECK: Line: 11, Column: 8 } -//CHECK: Function: _ZTS7derived -//CHECK: Args: -//CHECK-NEXT: String: 'Argument ' -//CHECK-NEXT: Argument: '1' -//CHECK-NEXT: String: ' for function kernel: ' -//CHECK-NEXT: String: '' -//CHECK-NEXT: String: ' ' -//CHECK-NEXT: String: _ZTS7derived -//CHECK-NEXT: String: . -//CHECK-NEXT: String: e -//CHECK-NEXT: String: '(' -//CHECK-NEXT: String: struct __wrapper_class -//CHECK-NEXT: String: ')' - -//CHECK: --- !Passed -//CHECK: Pass:{{.*}}sycl -//CHECK: Name:{{.*}}Region -//CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -//CHECK: Line: 86, Column: 7 } -//CHECK: Function: _ZTS7derived -//CHECK: Args: -//CHECK-NEXT: String: 'Argument ' -//CHECK-NEXT: Argument: '2' -//CHECK-NEXT: String: ' for function kernel: ' -//CHECK-NEXT: String: '' -//CHECK-NEXT: String: ' ' -//CHECK-NEXT: String: _ZTS7derived -//CHECK-NEXT: String: . -//CHECK-NEXT: String: a -//CHECK-NEXT: String: '(' -//CHECK-NEXT: String: int -//CHECK-NEXT: String: ')' - -struct derived : base, second_base { +struct kernelfunctor : notdecomposedbase, decomposedbase { int a; - + int *ptr; + int array[3]; + sycl::sampler sampl; void operator()() const { } }; @@ -92,10 +33,346 @@ struct derived : base, second_base { int main() { sycl::queue q; - q.submit([&](cl::sycl::handler &cgh) { - derived f{}; + q.submit([&](sycl::handler &cgh) { + kernelfunctor f{}; cgh.single_task(f); }); return 0; } + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '0' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for base class, +// CHECK-NEXT: String: struct notdecomposedbase +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: struct notdecomposedbase +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '4' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '1' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: 'Compiler generated argument for decomposed struct/class,' +// CHECK-NEXT: String: decomposedbase +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: 'Field:decompvar, ' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: float +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '4' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '2' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for nested pointer, +// CHECK-NEXT: String: decompptr +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: Compiler generated +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '8' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '3' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: decompAcc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: '__global char *' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '8' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '4' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: decompAcc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: 'struct sycl::range<1>' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '1' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '5' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: decompAcc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: 'struct sycl::range<1>' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '1' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '6' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: decompAcc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: 'struct sycl::id<1>' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '1' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '7' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for stream, +// CHECK-NEXT: String: decompStream +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: 'sycl::stream' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '3' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '8' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: acc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: '__global int *' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '8' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '9' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: acc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: 'struct sycl::range<1>' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '1' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '10' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: acc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: 'struct sycl::range<1>' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '1' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '11' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for accessor, +// CHECK-NEXT: String: acc +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: 'struct sycl::id<1>' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '1' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '12' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: a +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: int +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '4' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '13' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: ptr +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: '__global int *' +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '8' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '14' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for array, +// CHECK-NEXT: String: array +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: Compiler generated +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '12' +// CHECK-NEXT: String: ')' + +// CHECK: --- !Passed +// CHECK: Pass:{{.*}}sycl +// CHECK: Name:{{.*}}Region +// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// CHECK-NEXT: Line: 24, Column: 8 } +// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Args: +// CHECK-NEXT: String: 'Arg ' +// CHECK-NEXT: Argument: '15' +// CHECK-NEXT: String: ':' +// CHECK-NEXT: String: Compiler generated argument for sampler, +// CHECK-NEXT: String: sampl +// CHECK-NEXT: String: ' (' +// CHECK-NEXT: String: '' +// CHECK-NEXT: String: 'Type:' +// CHECK-NEXT: String: sampler_t +// CHECK-NEXT: String: ', ' +// CHECK-NEXT: String: 'Size: ' +// CHECK-NEXT: Argument: '8' +// CHECK-NEXT: String: ')' From f82497315176bbe48c7c393d1f27785937b29543 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Wed, 12 May 2021 16:26:36 -0700 Subject: [PATCH 2/5] Implement review comments Signed-off-by: Elizabeth Andrews --- .../clang/Basic/SyclOptReportHandler.h | 10 +-- clang/lib/CodeGen/CodeGenFunction.cpp | 15 ++-- clang/lib/Sema/SemaSYCL.cpp | 28 ++----- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 76 +++++++++---------- 4 files changed, 58 insertions(+), 71 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index bdac3c0eecbf8..490e957259775 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -43,11 +43,11 @@ class SyclOptReportHandler { llvm::DenseMap> Map; public: - void AddKernelArgs(const FunctionDecl *FD, std::string ArgName, - std::string ArgType, SourceLocation ArgLoc, - unsigned ArgSize, std::string ArgDesc, - std::string ArgParent) { - Map[FD].emplace_back(ArgName, ArgType, ArgLoc, ArgSize, ArgDesc, ArgParent); + void AddKernelArgs(const FunctionDecl *FD, StringRef ArgName, + StringRef ArgType, SourceLocation ArgLoc, unsigned ArgSize, + StringRef ArgDesc, StringRef ArgParent) { + Map[FD].emplace_back(ArgName.data(), ArgType.data(), ArgLoc, ArgSize, + ArgDesc.data(), ArgParent.data()); } SmallVector &GetInfo(const FunctionDecl *FD) { auto It = Map.find(FD); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index ccd0ec236d34a..7e3026a2dcf61 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1474,7 +1474,7 @@ enum class ArgDescEncoding { None }; -static ArgDescEncoding getArgDescEncoding(std::string KernelArgDesc) { +static ArgDescEncoding getArgDescEncoding(StringRef KernelArgDesc) { return llvm::StringSwitch(KernelArgDesc) .Case("base class", ArgDescEncoding::BaseClass) .Case("decomposed struct/class", ArgDescEncoding::DecomposedMember) @@ -1545,11 +1545,11 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, for (auto ORI : llvm::enumerate(OptReportHandler.GetInfo(FD))) { llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.value().KernelArgLoc); - std::string ArgName = ORI.value().KernelArgName; - std::string ArgType = ORI.value().KernelArgType; - std::string ArgDesc = ORI.value().KernelArgDesc; + StringRef ArgName = ORI.value().KernelArgName; + StringRef ArgType = ORI.value().KernelArgType; + StringRef ArgDesc = ORI.value().KernelArgDesc; unsigned ArgSize = ORI.value().KernelArgSize; - std::string ArgParent = ORI.value().KernelArgParent; + StringRef ArgParent = ORI.value().KernelArgParent; ArgDescEncoding ArgDescEnc = getArgDescEncoding(ArgDesc); bool isWrappedField = (ArgDescEnc == ArgDescEncoding::WrappedPointer || @@ -1561,13 +1561,14 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, &Fn->getEntryBlock()); Remark << "Arg " << llvm::ore::NV("Argument", ORI.index()) << ":" << ((ArgDescEnc != ArgDescEncoding::None) - ? ("Compiler generated argument for " + ArgDesc + ",") + ? ("Compiler generated argument for " + ArgDesc.str() + + ",") : "") << ((ArgDescEnc == ArgDescEncoding::DecomposedMember) ? ArgParent : ArgName) << " (" << ((ArgDescEnc == ArgDescEncoding::DecomposedMember) - ? ("Field:" + ArgName + ", ") + ? ("Field:" + ArgName.str() + ", ") : "") << "Type:" << ((isWrappedField) ? "Compiler generated" : ArgType) << ", " diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index ef6c5eaff6720..36bb4519d87c8 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2242,32 +2242,23 @@ StringRef getKernelArgDesc(KernelArgDescription Desc) { switch (Desc) { case KernelArgDescription::BaseClass: return "base class"; - break; case KernelArgDescription::DecomposedMember: return "decomposed struct/class"; - break; case KernelArgDescription::WrappedPointer: return "nested pointer"; - break; case KernelArgDescription::WrappedArray: return "array"; - break; case KernelArgDescription::Accessor: return "accessor"; - break; case KernelArgDescription::AccessorBase: return "accessor base class"; - break; case KernelArgDescription::Sampler: return "sampler"; - break; case KernelArgDescription::Stream: return "stream"; - break; case KernelArgDescription::KernelHandler: return "SYCL2020 specialization constant"; - break; - default: + case KernelArgDescription::None: return ""; } } @@ -2282,10 +2273,9 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); const RecordDecl *KernelArgParent = KernelArg->getParent(); SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( - DC.getKernelDecl(), KernelArg->getName().data(), - KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, - getKernelArgDesc(KernelArgDesc).data(), - KernelArgParent ? KernelArgParent->getName().data() : ""); + DC.getKernelDecl(), KernelArg->getName(), KernelArgType.getAsString(), + KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDesc), + KernelArgParent ? KernelArgParent->getName() : ""); } void addParam(const FieldDecl *FD, QualType FieldTy) { @@ -2303,16 +2293,14 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( DC.getKernelDecl(), KernelArgType.getAsString(), KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, - getKernelArgDesc(KernelArgDesc).data(), ""); + getKernelArgDesc(KernelArgDesc), ""); } // Handles SYCL special types - accessor, sampler and stream // Also handles modified types - arrays and pointers bool handleSpecialType(const FieldDecl *FD, QualType FieldTy, KernelArgDescription Desc) { - llvm::ArrayRef KernelParameters = - DC.getParamVarDeclsForCurrentField(); - for (auto *Param : KernelParameters) + for (const auto *Param : DC.getParamVarDeclsForCurrentField()) addParam(FD, Param->getType(), Desc); return true; } @@ -2329,9 +2317,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS, QualType FieldTy) final { - llvm::ArrayRef KernelParameters = - DC.getParamVarDeclsForCurrentField(); - for (auto *Param : KernelParameters) + for (const auto *Param : DC.getParamVarDeclsForCurrentField()) addParam(Param->getType(), KernelArgDescription(KernelArgDescription::AccessorBase)); return true; diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp index 760f1a9614a80..76771ce23690b 100644 --- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -8,24 +8,24 @@ sycl::handler H; -class decomposedbase { +class DecomposedBase { public: - float decompvar; - int *decompptr; + float DecompVar; + int *DecompPtr; sycl::accessor decompAcc; - sycl::stream decompStream{0, 0, H}; + sycl::stream DecompStream{0, 0, H}; }; -struct notdecomposedbase { +struct NotDecomposedBase { public: - int b; + int B; }; -struct kernelfunctor : notdecomposedbase, decomposedbase { - int a; - int *ptr; - int array[3]; - sycl::sampler sampl; +struct KernelFunctor : NotDecomposedBase, DecomposedBase { + int A; + int *Ptr; + int Array[3]; + sycl::sampler Sampl; void operator()() const { } }; @@ -34,7 +34,7 @@ int main() { sycl::queue q; q.submit([&](sycl::handler &cgh) { - kernelfunctor f{}; + KernelFunctor f{}; cgh.single_task(f); }); @@ -46,17 +46,17 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '0' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: Compiler generated argument for base class, -// CHECK-NEXT: String: struct notdecomposedbase +// CHECK-NEXT: String: struct NotDecomposedBase // CHECK-NEXT: String: ' (' // CHECK-NEXT: String: '' // CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: struct notdecomposedbase +// CHECK-NEXT: String: struct NotDecomposedBase // CHECK-NEXT: String: ', ' // CHECK-NEXT: String: 'Size: ' // CHECK-NEXT: Argument: '4' @@ -67,15 +67,15 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '1' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: 'Compiler generated argument for decomposed struct/class,' -// CHECK-NEXT: String: decomposedbase +// CHECK-NEXT: String: DecomposedBase // CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: 'Field:decompvar, ' +// CHECK-NEXT: String: 'Field:DecompVar, ' // CHECK-NEXT: String: 'Type:' // CHECK-NEXT: String: float // CHECK-NEXT: String: ', ' @@ -88,13 +88,13 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '2' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: Compiler generated argument for nested pointer, -// CHECK-NEXT: String: decompptr +// CHECK-NEXT: String: DecompPtr // CHECK-NEXT: String: ' (' // CHECK-NEXT: String: '' // CHECK-NEXT: String: 'Type:' @@ -109,7 +109,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '3' @@ -130,7 +130,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '4' @@ -151,7 +151,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '5' @@ -172,7 +172,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '6' @@ -193,13 +193,13 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '7' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: Compiler generated argument for stream, -// CHECK-NEXT: String: decompStream +// CHECK-NEXT: String: DecompStream // CHECK-NEXT: String: ' (' // CHECK-NEXT: String: '' // CHECK-NEXT: String: 'Type:' @@ -214,7 +214,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '8' @@ -235,7 +235,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '9' @@ -256,7 +256,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '10' @@ -277,7 +277,7 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '11' @@ -298,13 +298,13 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '12' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: '' -// CHECK-NEXT: String: a +// CHECK-NEXT: String: A // CHECK-NEXT: String: ' (' // CHECK-NEXT: String: '' // CHECK-NEXT: String: 'Type:' @@ -319,13 +319,13 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '13' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: '' -// CHECK-NEXT: String: ptr +// CHECK-NEXT: String: Ptr // CHECK-NEXT: String: ' (' // CHECK-NEXT: String: '' // CHECK-NEXT: String: 'Type:' @@ -340,13 +340,13 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '14' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: Compiler generated argument for array, -// CHECK-NEXT: String: array +// CHECK-NEXT: String: Array // CHECK-NEXT: String: ' (' // CHECK-NEXT: String: '' // CHECK-NEXT: String: 'Type:' @@ -361,13 +361,13 @@ int main() { // CHECK: Name:{{.*}}Region // CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', // CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13kernelfunctor +// CHECK-NEXT: Function: _ZTS13KernelFunctor // CHECK-NEXT: Args: // CHECK-NEXT: String: 'Arg ' // CHECK-NEXT: Argument: '15' // CHECK-NEXT: String: ':' // CHECK-NEXT: String: Compiler generated argument for sampler, -// CHECK-NEXT: String: sampl +// CHECK-NEXT: String: Sampl // CHECK-NEXT: String: ' (' // CHECK-NEXT: String: '' // CHECK-NEXT: String: 'Type:' From 7d2e1bb078084c30bd6e8097b698c26604266479 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Thu, 13 May 2021 15:40:28 -0700 Subject: [PATCH 3/5] Construct description in Sema instead of CodeGen. Also added test for accessor as base class and SYCL 2020 spec constants. Signed-off-by: Elizabeth Andrews --- .../clang/Basic/SyclOptReportHandler.h | 26 +- clang/lib/CodeGen/CodeGenFunction.cpp | 51 +- clang/lib/Sema/SemaSYCL.cpp | 71 +- clang/test/SemaSYCL/kernel-arg-opt-report.cpp | 820 +++++++++++------- 4 files changed, 555 insertions(+), 413 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index 490e957259775..681caa89de06f 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -25,29 +25,31 @@ class FunctionDecl; class SyclOptReportHandler { private: struct OptReportInfo { - std::string KernelArgName; + std::string KernelArgDescName; // Kernel argument name itself, or the name + // of the parent class if the kernel argument + // is a decomposed member std::string KernelArgType; SourceLocation KernelArgLoc; unsigned KernelArgSize; std::string KernelArgDesc; - std::string KernelArgParent; + std::string KernelArgDecomposedField; - OptReportInfo(std::string ArgName, std::string ArgType, + OptReportInfo(std::string ArgDescName, std::string ArgType, SourceLocation ArgLoc, unsigned ArgSize, std::string ArgDesc, - std::string ArgParent) - : KernelArgName(std::move(ArgName)), KernelArgType(std::move(ArgType)), - KernelArgLoc(ArgLoc), KernelArgSize(ArgSize), - KernelArgDesc(std::move(ArgDesc)), - KernelArgParent(std::move(ArgParent)) {} + std::string ArgDecomposedField) + : KernelArgDescName(std::move(ArgDescName)), + KernelArgType(std::move(ArgType)), KernelArgLoc(ArgLoc), + KernelArgSize(ArgSize), KernelArgDesc(std::move(ArgDesc)), + KernelArgDecomposedField(std::move(ArgDecomposedField)) {} }; llvm::DenseMap> Map; public: - void AddKernelArgs(const FunctionDecl *FD, StringRef ArgName, + void AddKernelArgs(const FunctionDecl *FD, StringRef ArgDescName, StringRef ArgType, SourceLocation ArgLoc, unsigned ArgSize, - StringRef ArgDesc, StringRef ArgParent) { - Map[FD].emplace_back(ArgName.data(), ArgType.data(), ArgLoc, ArgSize, - ArgDesc.data(), ArgParent.data()); + StringRef ArgDesc, StringRef ArgDecomposedField) { + Map[FD].emplace_back(ArgDescName.data(), ArgType.data(), ArgLoc, ArgSize, + ArgDesc.data(), ArgDecomposedField.data()); } SmallVector &GetInfo(const FunctionDecl *FD) { auto It = Map.find(FD); diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp index 7e3026a2dcf61..a488ac6016349 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1461,33 +1461,6 @@ QualType CodeGenFunction::BuildFunctionArgList(GlobalDecl GD, return ResTy; } -enum class ArgDescEncoding { - BaseClass, - DecomposedMember, - WrappedPointer, - WrappedArray, - Accessor, - AccessorBase, - Sampler, - Stream, - KernelHandler, - None -}; - -static ArgDescEncoding getArgDescEncoding(StringRef KernelArgDesc) { - return llvm::StringSwitch(KernelArgDesc) - .Case("base class", ArgDescEncoding::BaseClass) - .Case("decomposed struct/class", ArgDescEncoding::DecomposedMember) - .Case("nested pointer", ArgDescEncoding::WrappedPointer) - .Case("array", ArgDescEncoding::WrappedArray) - .Case("accessor", ArgDescEncoding::Accessor) - .Case("accessor base class", ArgDescEncoding::AccessorBase) - .Case("sampler", ArgDescEncoding::Sampler) - .Case("stream", ArgDescEncoding::Stream) - .Case("SYCL2020 specialization constant", ArgDescEncoding::KernelHandler) - .Default(ArgDescEncoding::None); -} - void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, const CGFunctionInfo &FnInfo) { const FunctionDecl *FD = cast(GD.getDecl()); @@ -1545,33 +1518,17 @@ void CodeGenFunction::GenerateCode(GlobalDecl GD, llvm::Function *Fn, for (auto ORI : llvm::enumerate(OptReportHandler.GetInfo(FD))) { llvm::DiagnosticLocation DL = SourceLocToDebugLoc(ORI.value().KernelArgLoc); - StringRef ArgName = ORI.value().KernelArgName; + StringRef NameInDesc = ORI.value().KernelArgDescName; StringRef ArgType = ORI.value().KernelArgType; StringRef ArgDesc = ORI.value().KernelArgDesc; unsigned ArgSize = ORI.value().KernelArgSize; - StringRef ArgParent = ORI.value().KernelArgParent; - - ArgDescEncoding ArgDescEnc = getArgDescEncoding(ArgDesc); - bool isWrappedField = (ArgDescEnc == ArgDescEncoding::WrappedPointer || - ArgDescEnc == ArgDescEncoding::WrappedArray) - ? true - : false; + StringRef ArgDecomposedField = ORI.value().KernelArgDecomposedField; llvm::OptimizationRemark Remark("sycl", "Region", DL, &Fn->getEntryBlock()); Remark << "Arg " << llvm::ore::NV("Argument", ORI.index()) << ":" - << ((ArgDescEnc != ArgDescEncoding::None) - ? ("Compiler generated argument for " + ArgDesc.str() + - ",") - : "") - << ((ArgDescEnc == ArgDescEncoding::DecomposedMember) ? ArgParent - : ArgName) - << " (" - << ((ArgDescEnc == ArgDescEncoding::DecomposedMember) - ? ("Field:" + ArgName.str() + ", ") - : "") - << "Type:" << ((isWrappedField) ? "Compiler generated" : ArgType) - << ", " + << ArgDesc << NameInDesc << " (" << ArgDecomposedField + << "Type:" << ArgType << ", " << "Size: " << llvm::ore::NV("Argument", ArgSize) << ")"; ORE.emit(Remark); } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 36bb4519d87c8..b4c66ede2bf20 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2241,23 +2241,23 @@ enum class KernelArgDescription { StringRef getKernelArgDesc(KernelArgDescription Desc) { switch (Desc) { case KernelArgDescription::BaseClass: - return "base class"; + return "Compiler generated argument for base class,"; case KernelArgDescription::DecomposedMember: - return "decomposed struct/class"; + return "Compiler generated argument for decomposed struct/class,"; case KernelArgDescription::WrappedPointer: - return "nested pointer"; + return "Compiler generated argument for nested pointer,"; case KernelArgDescription::WrappedArray: - return "array"; + return "Compiler generated argument for array,"; case KernelArgDescription::Accessor: - return "accessor"; + return "Compiler generated argument for accessor,"; case KernelArgDescription::AccessorBase: - return "accessor base class"; + return "Compiler generated argument for accessor base class,"; case KernelArgDescription::Sampler: - return "sampler"; + return "Compiler generated argument for sampler,"; case KernelArgDescription::Stream: - return "stream"; + return "Compiler generated argument for stream,"; case KernelArgDescription::KernelHandler: - return "SYCL2020 specialization constant"; + return "Compiler generated argument for SYCL2020 specialization constant"; case KernelArgDescription::None: return ""; } @@ -2269,13 +2269,29 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { void addParam(const FieldDecl *KernelArg, QualType KernelArgType, KernelArgDescription KernelArgDesc) { + StringRef NameToEmitInDescription = KernelArg->getName(); + const RecordDecl *KernelArgParent = KernelArg->getParent(); + if (KernelArgParent && + KernelArgDesc == KernelArgDescription::DecomposedMember) { + NameToEmitInDescription = KernelArgParent->getName(); + } + + bool isWrappedField = + (KernelArgDesc == KernelArgDescription::WrappedPointer || + KernelArgDesc == KernelArgDescription::WrappedArray) + ? true + : false; + unsigned KernelArgSize = SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); - const RecordDecl *KernelArgParent = KernelArg->getParent(); + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( - DC.getKernelDecl(), KernelArg->getName(), KernelArgType.getAsString(), + DC.getKernelDecl(), NameToEmitInDescription, + (isWrappedField) ? "Compiler generated" : KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDesc), - KernelArgParent ? KernelArgParent->getName() : ""); + (KernelArgDesc == KernelArgDescription::DecomposedMember) + ? ("Field:" + KernelArg->getName().str() + ", ") + : ""); } void addParam(const FieldDecl *FD, QualType FieldTy) { @@ -2287,7 +2303,9 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { addParam(FD, FieldTy, Desc); } - void addParam(QualType KernelArgType, KernelArgDescription KernelArgDesc) { + // Handles base classes + void addParam(const CXXBaseSpecifier &, QualType KernelArgType, + KernelArgDescription KernelArgDesc) { unsigned KernelArgSize = SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( @@ -2296,6 +2314,16 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { getKernelArgDesc(KernelArgDesc), ""); } + // Handles specialization constants + void addParam(QualType KernelArgType, KernelArgDescription KernelArgDesc) { + unsigned KernelArgSize = + SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( + DC.getKernelDecl(), "", KernelArgType.getAsString(), + KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDesc), + ""); + } + // Handles SYCL special types - accessor, sampler and stream // Also handles modified types - arrays and pointers bool handleSpecialType(const FieldDecl *FD, QualType FieldTy, @@ -2317,9 +2345,18 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS, QualType FieldTy) final { - for (const auto *Param : DC.getParamVarDeclsForCurrentField()) - addParam(Param->getType(), - KernelArgDescription(KernelArgDescription::AccessorBase)); + for (const auto *Param : DC.getParamVarDeclsForCurrentField()) { + QualType KernelArgType = Param->getType(); + unsigned KernelArgSize = SemaRef.getASTContext() + .getTypeSizeInChars(KernelArgType) + .getQuantity(); + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( + DC.getKernelDecl(), FieldTy.getAsString(), + KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, + getKernelArgDesc( + KernelArgDescription(KernelArgDescription::AccessorBase)), + ""); + } return true; } @@ -2359,7 +2396,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { bool handleNonDecompStruct(const CXXRecordDecl *Base, const CXXBaseSpecifier &BS, QualType Ty) final { - addParam(Ty, KernelArgDescription(KernelArgDescription::BaseClass)); + addParam(BS, Ty, KernelArgDescription(KernelArgDescription::BaseClass)); return true; } diff --git a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp index 76771ce23690b..967e269dace2d 100644 --- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -1,6 +1,10 @@ // RUN: %clang_cc1 -triple spir64-unknown-unknown-sycldevice -fsycl-is-device \ // RUN: -Wno-sycl-2017-compat -emit-llvm-bc %s -o %t-host.bc -opt-record-file %t-host.yaml -// RUN: FileCheck -check-prefix=CHECK --input-file %t-host.yaml %s +// RUN: FileCheck -check-prefix=SPIR --input-file %t-host.yaml %s + +// RUN: %clang_cc1 -triple nvptx64-unknown-unknown -fsycl-is-device \ +// RUN: -Wno-sycl-2017-compat -emit-llvm-bc %s -o %t-host.bc -opt-record-file %t-host.yaml +// RUN: FileCheck -check-prefix=NVPTX --input-file %t-host.yaml %s // The test generates remarks about the kernel argument, their location and type // in the resulting yaml file. @@ -30,349 +34,491 @@ struct KernelFunctor : NotDecomposedBase, DecomposedBase { } }; +struct AccessorDerived : sycl::accessor { + int B; +}; + int main() { sycl::queue q; - q.submit([&](sycl::handler &cgh) { KernelFunctor f{}; cgh.single_task(f); }); + AccessorDerived DerivedObject; + q.submit([&](sycl::handler &cgh) { + sycl::kernel_handler kh; + + cgh.single_task( + [=](auto) { + DerivedObject.use(); + }, + kh); + }); + return 0; } -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '0' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for base class, -// CHECK-NEXT: String: struct NotDecomposedBase -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: struct NotDecomposedBase -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '4' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '1' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: 'Compiler generated argument for decomposed struct/class,' -// CHECK-NEXT: String: DecomposedBase -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: 'Field:DecompVar, ' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: float -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '4' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '2' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for nested pointer, -// CHECK-NEXT: String: DecompPtr -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: Compiler generated -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '8' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '3' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: decompAcc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: '__global char *' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '8' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '4' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: decompAcc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: 'struct sycl::range<1>' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '1' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '5' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: decompAcc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: 'struct sycl::range<1>' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '1' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '6' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: decompAcc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: 'struct sycl::id<1>' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '1' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '7' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for stream, -// CHECK-NEXT: String: DecompStream -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: 'sycl::stream' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '3' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '8' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: acc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: '__global int *' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '8' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '9' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: acc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: 'struct sycl::range<1>' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '1' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '10' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: acc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: 'struct sycl::range<1>' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '1' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '11' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for accessor, -// CHECK-NEXT: String: acc -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: 'struct sycl::id<1>' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '1' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '12' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: A -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: int -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '4' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '13' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: Ptr -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: '__global int *' -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '8' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '14' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for array, -// CHECK-NEXT: String: Array -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: Compiler generated -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '12' -// CHECK-NEXT: String: ')' - -// CHECK: --- !Passed -// CHECK: Pass:{{.*}}sycl -// CHECK: Name:{{.*}}Region -// CHECK: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', -// CHECK-NEXT: Line: 24, Column: 8 } -// CHECK-NEXT: Function: _ZTS13KernelFunctor -// CHECK-NEXT: Args: -// CHECK-NEXT: String: 'Arg ' -// CHECK-NEXT: Argument: '15' -// CHECK-NEXT: String: ':' -// CHECK-NEXT: String: Compiler generated argument for sampler, -// CHECK-NEXT: String: Sampl -// CHECK-NEXT: String: ' (' -// CHECK-NEXT: String: '' -// CHECK-NEXT: String: 'Type:' -// CHECK-NEXT: String: sampler_t -// CHECK-NEXT: String: ', ' -// CHECK-NEXT: String: 'Size: ' -// CHECK-NEXT: Argument: '8' -// CHECK-NEXT: String: ')' +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '0' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for base class, +// SPIR-NEXT: String: struct NotDecomposedBase +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: struct NotDecomposedBase +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '4' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: 'Compiler generated argument for decomposed struct/class,' +// SPIR-NEXT: String: DecomposedBase +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: 'Field:DecompVar, ' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: float +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '4' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '2' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for nested pointer, +// SPIR-NEXT: String: DecompPtr +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: Compiler generated +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '3' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: decompAcc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: '__global char *' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '4' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: decompAcc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::range<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '5' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: decompAcc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::range<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '6' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: decompAcc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::id<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '7' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for stream, +// SPIR-NEXT: String: DecompStream +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'sycl::stream' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '3' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: acc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: '__global int *' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '9' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: acc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::range<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '10' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: acc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::range<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '11' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor, +// SPIR-NEXT: String: acc +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::id<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '12' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: A +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: int +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '4' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '13' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: Ptr +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: '__global int *' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '14' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for array, +// SPIR-NEXT: String: Array +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: Compiler generated +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '12' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 28, Column: 8 } +// SPIR-NEXT: Function: _ZTS13KernelFunctor +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '15' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for sampler, +// SPIR-NEXT: String: Sampl +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: sampler_t +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' + +// Output for kernel XYZ + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Function: '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ' +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '0' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor base class, +// SPIR-NEXT: String: 'sycl::accessor' +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: '__global char *' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '8' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Function: '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ' +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor base class, +// SPIR-NEXT: String: 'sycl::accessor' +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::range<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Function: '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ' +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '2' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor base class, +// SPIR-NEXT: String: 'sycl::accessor' +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::range<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Function: '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ' +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '3' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: Compiler generated argument for accessor base class, +// SPIR-NEXT: String: 'sycl::accessor' +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: '' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: 'struct sycl::id<1>' +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '1' +// SPIR-NEXT: String: ')' + +// SPIR: --- !Passed +// SPIR: Pass:{{.*}}sycl +// SPIR: Name:{{.*}}Region +// SPIR: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// SPIR-NEXT: Line: 53, Column: 9 } +// SPIR-NEXT: Function: '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ' +// SPIR-NEXT: Args: +// SPIR-NEXT: String: 'Arg ' +// SPIR-NEXT: Argument: '4' +// SPIR-NEXT: String: ':' +// SPIR-NEXT: String: 'Compiler generated argument for decomposed struct/class,' +// SPIR-NEXT: String: AccessorDerived +// SPIR-NEXT: String: ' (' +// SPIR-NEXT: String: 'Field:B, ' +// SPIR-NEXT: String: 'Type:' +// SPIR-NEXT: String: int +// SPIR-NEXT: String: ', ' +// SPIR-NEXT: String: 'Size: ' +// SPIR-NEXT: Argument: '4' +// SPIR-NEXT: String: ')' + +// NVPTX: --- !Passed +// NVPTX: Pass:{{.*}}sycl +// NVPTX: Name:{{.*}}Region +// NVPTX: DebugLoc:{{.*}} { File: '{{.*}}kernel-arg-opt-report.cpp', +// NVPTX: Line: 53, Column: 9 } +// NVPTX-NEXT: Function: '_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE3XYZ' +// NVPTX-NEXT: Args: +// NVPTX-NEXT: String: 'Arg ' +// NVPTX: Argument: '5' +// NVPTX-NEXT: String: ':' +// NVPTX-NEXT: String: Compiler generated argument for SYCL2020 specialization constant +// NVPTX-NEXT: String: '' +// NVPTX-NEXT: String: ' (' +// NVPTX-NEXT: String: '' +// NVPTX-NEXT: String: 'Type:' +// NVPTX-NEXT: String: '__global char *' +// NVPTX-NEXT: String: ', ' +// NVPTX-NEXT: String: 'Size: ' +// NVPTX-NEXT: Argument: '8' +// NVPTX-NEXT: String: ')' From de373082bce4e74e80843c892c5d7adda3a6df56 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 14 May 2021 06:09:38 -0700 Subject: [PATCH 4/5] Implement review comments Signed-off-by: Elizabeth Andrews --- clang/include/clang/Basic/SyclOptReportHandler.h | 2 +- clang/lib/Sema/SemaSYCL.cpp | 16 +++++++--------- 2 files changed, 8 insertions(+), 10 deletions(-) diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index 681caa89de06f..9b15c4ccd294a 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -27,7 +27,7 @@ class SyclOptReportHandler { struct OptReportInfo { std::string KernelArgDescName; // Kernel argument name itself, or the name // of the parent class if the kernel argument - // is a decomposed member + // is a decomposed member. std::string KernelArgType; SourceLocation KernelArgLoc; unsigned KernelArgSize; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b4c66ede2bf20..38f9970b4b6e7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2277,17 +2277,15 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { } bool isWrappedField = - (KernelArgDesc == KernelArgDescription::WrappedPointer || - KernelArgDesc == KernelArgDescription::WrappedArray) - ? true - : false; + KernelArgDesc == KernelArgDescription::WrappedPointer || + KernelArgDesc == KernelArgDescription::WrappedArray; unsigned KernelArgSize = SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( DC.getKernelDecl(), NameToEmitInDescription, - (isWrappedField) ? "Compiler generated" : KernelArgType.getAsString(), + isWrappedField ? "Compiler generated" : KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDesc), (KernelArgDesc == KernelArgDescription::DecomposedMember) ? ("Field:" + KernelArg->getName().str() + ", ") @@ -2303,7 +2301,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { addParam(FD, FieldTy, Desc); } - // Handles base classes + // Handles base classes. void addParam(const CXXBaseSpecifier &, QualType KernelArgType, KernelArgDescription KernelArgDesc) { unsigned KernelArgSize = @@ -2314,7 +2312,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { getKernelArgDesc(KernelArgDesc), ""); } - // Handles specialization constants + // Handles specialization constants. void addParam(QualType KernelArgType, KernelArgDescription KernelArgDesc) { unsigned KernelArgSize = SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); @@ -2324,8 +2322,8 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { ""); } - // Handles SYCL special types - accessor, sampler and stream - // Also handles modified types - arrays and pointers + // Handles SYCL special types (accessor, sampler and stream) and modified + // types (arrays and pointers) bool handleSpecialType(const FieldDecl *FD, QualType FieldTy, KernelArgDescription Desc) { for (const auto *Param : DC.getParamVarDeclsForCurrentField()) From 7452f9f6f87487f02812d01d47a4366c1c153190 Mon Sep 17 00:00:00 2001 From: Elizabeth Andrews Date: Fri, 14 May 2021 09:42:46 -0700 Subject: [PATCH 5/5] Fix compiler warning about reaching end of non-void function Signed-off-by: Elizabeth Andrews --- clang/lib/Sema/SemaSYCL.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 38f9970b4b6e7..bb6ab605b5dcf 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2261,6 +2261,8 @@ StringRef getKernelArgDesc(KernelArgDescription Desc) { case KernelArgDescription::None: return ""; } + llvm_unreachable( + "switch should cover all possible values for KernelArgDescription"); } class SyclOptReportCreator : public SyclKernelFieldHandler {