diff --git a/clang/include/clang/Basic/SyclOptReportHandler.h b/clang/include/clang/Basic/SyclOptReportHandler.h index af3ba5ae1c0a..9b15c4ccd294 100644 --- a/clang/include/clang/Basic/SyclOptReportHandler.h +++ b/clang/include/clang/Basic/SyclOptReportHandler.h @@ -25,21 +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 KernelArgDecomposedField; - OptReportInfo(std::string ArgName, std::string ArgType, - SourceLocation ArgLoc) - : KernelArgName(std::move(ArgName)), KernelArgType(std::move(ArgType)), - KernelArgLoc(ArgLoc) {} + OptReportInfo(std::string ArgDescName, std::string ArgType, + SourceLocation ArgLoc, unsigned ArgSize, std::string ArgDesc, + 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, std::string ArgName, - std::string ArgType, SourceLocation ArgLoc) { - Map[FD].emplace_back(ArgName, ArgType, ArgLoc); + void AddKernelArgs(const FunctionDecl *FD, StringRef ArgDescName, + StringRef ArgType, SourceLocation ArgLoc, unsigned ArgSize, + 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 3a0b15484d80..a488ac601634 100644 --- a/clang/lib/CodeGen/CodeGenFunction.cpp +++ b/clang/lib/CodeGen/CodeGenFunction.cpp @@ -1518,14 +1518,18 @@ 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; + StringRef NameInDesc = ORI.value().KernelArgDescName; + StringRef ArgType = ORI.value().KernelArgType; + StringRef ArgDesc = ORI.value().KernelArgDesc; + unsigned ArgSize = ORI.value().KernelArgSize; + StringRef ArgDecomposedField = ORI.value().KernelArgDecomposedField; + 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()) << ":" + << 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 7cfb5bdbc5b7..bb6ab605b5dc 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,216 @@ 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 "Compiler generated argument for base class,"; + case KernelArgDescription::DecomposedMember: + return "Compiler generated argument for decomposed struct/class,"; + case KernelArgDescription::WrappedPointer: + return "Compiler generated argument for nested pointer,"; + case KernelArgDescription::WrappedArray: + return "Compiler generated argument for array,"; + case KernelArgDescription::Accessor: + return "Compiler generated argument for accessor,"; + case KernelArgDescription::AccessorBase: + return "Compiler generated argument for accessor base class,"; + case KernelArgDescription::Sampler: + return "Compiler generated argument for sampler,"; + case KernelArgDescription::Stream: + return "Compiler generated argument for stream,"; + case KernelArgDescription::KernelHandler: + return "Compiler generated argument for SYCL2020 specialization constant"; + case KernelArgDescription::None: + return ""; + } + llvm_unreachable( + "switch should cover all possible values for KernelArgDescription"); +} + +class SyclOptReportCreator : public SyclKernelFieldHandler { + SyclKernelDeclCreator &DC; + SourceLocation KernelInvocationLoc; + + 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; + + unsigned KernelArgSize = + SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); + + SemaRef.getDiagnostics().getSYCLOptReportHandler().AddKernelArgs( + DC.getKernelDecl(), NameToEmitInDescription, + isWrappedField ? "Compiler generated" : KernelArgType.getAsString(), + KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDesc), + (KernelArgDesc == KernelArgDescription::DecomposedMember) + ? ("Field:" + KernelArg->getName().str() + ", ") + : ""); + } + + 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); + } + + // Handles base classes. + void addParam(const CXXBaseSpecifier &, 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), ""); + } + + // 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) and modified + // types (arrays and pointers) + bool handleSpecialType(const FieldDecl *FD, QualType FieldTy, + KernelArgDescription Desc) { + for (const auto *Param : DC.getParamVarDeclsForCurrentField()) + 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 { + 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; + } + + 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(BS, 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) @@ -3552,18 +3757,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 fe71db3e833e..967e269dace2 100644 --- a/clang/test/SemaSYCL/kernel-arg-opt-report.cpp +++ b/clang/test/SemaSYCL/kernel-arg-opt-report.cpp @@ -1,101 +1,524 @@ // 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. #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; + int B; }; -//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 { - int a; - +struct KernelFunctor : NotDecomposedBase, DecomposedBase { + int A; + int *Ptr; + int Array[3]; + sycl::sampler Sampl; void operator()() const { } }; +struct AccessorDerived : sycl::accessor { + int B; +}; + int main() { sycl::queue q; - - q.submit([&](cl::sycl::handler &cgh) { - derived f{}; + 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; } + +// 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: ')'