diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3505c1c05e370..664d26e8a9cb6 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9532,6 +9532,7 @@ def err_sycl_restrict : Error< "|use rtti" "|use a non-const static data variable" "|call a virtual function" + "|call a recursive function" "|call through a function pointer" "|allocate storage" "|use exceptions" @@ -9539,6 +9540,7 @@ def err_sycl_restrict : Error< def err_sycl_virtual_types : Error< "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">; def note_sycl_used_here : Note<"used here">; +def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">; def err_sycl_non_std_layout_type : Error< "kernel parameter has non-standard layout class/struct type">; } // end of sema component. diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index e4a7988967a6d..5862c3bf5c6ca 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -216,6 +216,7 @@ LANGOPT(CUDADeviceApproxTranscendentals, 1, 0, "using approximate transcendental LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code") LANGOPT(SYCL , 1, 0, "Generate code for SYCL device") +LANGOPT(SYCLUseBitcode , 1, 0, "Generate bitcode for SYCL") LANGOPT(SizedDeallocation , 1, 0, "sized deallocation") LANGOPT(AlignedAllocation , 1, 0, "aligned allocation") LANGOPT(AlignedAllocationUnavailable, 1, 0, "aligned allocation functions are unavailable") diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index f48702718637c..e149448de0e3c 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -1701,6 +1701,10 @@ def fsycl_add_targets_EQ : CommaJoined<["-"], "fsycl-add-targets=">, Flags<[Driv HelpText<"Specify comma-separated list of triple and device binary image pairs to add to the final SYCL binary">; def fsycl_link_targets_EQ : CommaJoined<["-"], "fsycl-link-targets=">, Flags<[DriverOption, CC1Option]>, HelpText<"Specify comma-separated list of triples SYCL offloading targets to produce linked device images">; +def fsycl_use_bitcode : Flag<["-"], "fsycl-use-bitcode">, + Flags<[CC1Option]>, HelpText<"Use LLVM bitcode instead of SPIR-V in fat objects">; +def fno_sycl_use_bitcode : Flag<["-"], "fno-sycl-use-bitcode">, + Flags<[CC1Option]>, HelpText<"Use SPIR-V instead of LLVM bitcode in fat objects">; def fsyntax_only : Flag<["-"], "fsyntax-only">, Flags<[DriverOption,CoreOption,CC1Option]>, Group; def ftabstop_EQ : Joined<["-"], "ftabstop=">, Group; diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index b6ff9044f8251..365d96ead65e8 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -855,6 +855,7 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action, case Backend_EmitSPIRV: if (LangOpts.SYCL) { + // TODO: SPIRVNoDerefAttr is not modeled when using the bitcode pass SPIRV::SPIRVNoDerefAttr = true; // TODO: this pass added to work around missing linkonce_odr in SPIR-V PerModulePasses.add( @@ -862,7 +863,11 @@ void EmitAssemblyHelper::EmitAssembly(BackendAction Action, PerModulePasses.add(createASFixerPass()); PerModulePasses.add(createDeadCodeEliminationPass()); } - PerModulePasses.add(createSPIRVWriterPass(*OS)); + if (LangOpts.SYCLUseBitcode) + PerModulePasses.add( + createBitcodeWriterPass(*OS, CodeGenOpts.EmitLLVMUseLists, false)); + else + PerModulePasses.add(createSPIRVWriterPass(*OS)); break; diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 11bf7ab7d433a..dd2e214175f73 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -3523,6 +3523,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back("-aux-triple"); CmdArgs.push_back(Args.MakeArgString(NormalizedTriple)); CmdArgs.push_back("-disable-llvm-passes"); + if (Args.hasFlag(options::OPT_fsycl_use_bitcode, + options::OPT_fno_sycl_use_bitcode, true)) { + CmdArgs.push_back("-fsycl-use-bitcode"); + } } if (IsOpenMPDevice) { diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index d0b0ca40bfb8d..e7f7b38551524 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1,9 +1,8 @@ //===--- SYCL.cpp - SYCL Tool and ToolChain Implementations -----*- C++ -*-===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -40,6 +39,7 @@ const char *SYCL::Linker::constructLLVMSpirvCommand(Compilation &C, CmdArgs.push_back("-o"); CmdArgs.push_back(OutputFileName); } else { + CmdArgs.push_back("-spirv-no-deref-attr"); CmdArgs.push_back("-o"); CmdArgs.push_back(Output.getFilename()); } @@ -87,7 +87,8 @@ void SYCL::Linker::constructLlcCommand(Compilation &C, const JobAction &JA, } // For SYCL the inputs of the linker job are SPIR-V binaries and output is -// a single SPIR-V binary. +// a single SPIR-V binary. Input can also be bitcode when specified by +// the user void SYCL::Linker::ConstructJob(Compilation &C, const JobAction &JA, const InputInfo &Output, const InputInfoList &Inputs, @@ -110,9 +111,15 @@ void SYCL::Linker::ConstructJob(Compilation &C, const JobAction &JA, for (const auto &II : Inputs) { if (!II.isFilename()) continue; - const char *LLVMSpirvOutputFile = - constructLLVMSpirvCommand(C, JA, Output, Prefix, true, II.getFilename()); - SpirvInputs.push_back(LLVMSpirvOutputFile); + if (Args.hasFlag(options::OPT_fsycl_use_bitcode, + options::OPT_fno_sycl_use_bitcode, true)) + SpirvInputs.push_back(II.getFilename()); + else { + const char *LLVMSpirvOutputFile = + constructLLVMSpirvCommand(C, JA, Output, Prefix, true, + II.getFilename()); + SpirvInputs.push_back(LLVMSpirvOutputFile); + } } const char *LLVMLinkOutputFile = constructLLVMLinkCommand(C, JA, SubArchName, Prefix, SpirvInputs); diff --git a/clang/lib/Driver/ToolChains/SYCL.h b/clang/lib/Driver/ToolChains/SYCL.h index c18d3eaad0ded..a67c6ea2dbfa1 100644 --- a/clang/lib/Driver/ToolChains/SYCL.h +++ b/clang/lib/Driver/ToolChains/SYCL.h @@ -1,9 +1,8 @@ //===--- SYCL.h - SYCL ToolChain Implementations -----------------*- C++ -*-===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index d85a05ca5b3ca..a30c034f5eea6 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2883,6 +2883,8 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, } Opts.SYCL = Args.hasArg(options::OPT_fsycl_is_device); + Opts.SYCLUseBitcode = Args.hasFlag(options::OPT_fsycl_use_bitcode, + options::OPT_fno_sycl_use_bitcode, false); // Set CUDA mode for OpenMP target NVPTX if specified in options Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && T.isNVPTX() && diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index 9b2e0d9e4c6d2..1e1690c120083 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -125,6 +125,7 @@ set(cuda_wrapper_files set(sycl_wrapper_files sycl_wrappers/algorithm sycl_wrappers/random + sycl_wrappers/__config ) set(output_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include) diff --git a/clang/lib/Headers/sycl_wrappers/__config b/clang/lib/Headers/sycl_wrappers/__config new file mode 100644 index 0000000000000..35670334636b3 --- /dev/null +++ b/clang/lib/Headers/sycl_wrappers/__config @@ -0,0 +1,36 @@ +// -*- C++ -*- +//===--------------------------- __config ---------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef _LIBCPP_STD_VER +# if __cplusplus <= 201103L +# define _LIBCPP_STD_VER 11 +# elif __cplusplus <= 201402L +# define _LIBCPP_STD_VER 14 +# elif __cplusplus <= 201703L +# define _LIBCPP_STD_VER 17 +# endif +#endif + +#if _LIBCPP_STD_VER > 11 +# define _LIBCPP_CONSTEXPR_AFTER_CXX11 constexpr +#else +# define _LIBCPP_CONSTEXPR_AFTER_CXX11 +#endif + +#if _LIBCPP_STD_VER > 14 +# define _LIBCPP_CONSTEXPR_AFTER_CXX14 constexpr +#else +# define _LIBCPP_CONSTEXPR_AFTER_CXX14 +#endif + +#if _LIBCPP_STD_VER > 17 +# define _LIBCPP_CONSTEXPR_AFTER_CXX17 constexpr +#else +# define _LIBCPP_CONSTEXPR_AFTER_CXX17 +#endif diff --git a/clang/lib/Headers/sycl_wrappers/algorithm b/clang/lib/Headers/sycl_wrappers/algorithm index 94375c9b3979c..bff60089a903a 100644 --- a/clang/lib/Headers/sycl_wrappers/algorithm +++ b/clang/lib/Headers/sycl_wrappers/algorithm @@ -1,13 +1,13 @@ // -*- C++ -*- //===-------------------------- algorithm ---------------------------------===// // -// The LLVM Compiler Infrastructure -// -// This file is dual licensed under the MIT and the University of Illinois Open -// Source Licenses. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// +#include <__config> #include #include @@ -495,59 +495,77 @@ RandomAccessIterator is_heap_until(RandomAccessIterator first, RandomAccessIterator last, Compare comp); template -ForwardIterator min_element(ForwardIterator first, ForwardIterator last); +_LIBCPP_CONSTEXPR_AFTER_CXX11 ForwardIterator min_element(ForwardIterator first, + ForwardIterator last); template -ForwardIterator min_element(ForwardIterator first, ForwardIterator last, - Compare comp); +_LIBCPP_CONSTEXPR_AFTER_CXX11 ForwardIterator min_element(ForwardIterator first, + ForwardIterator last, + Compare comp); -template const T &min(const T &a, const T &b); +template +_LIBCPP_CONSTEXPR_AFTER_CXX11 const T &min(const T &a, const T &b); template -const T &min(const T &a, const T &b, Compare comp); +_LIBCPP_CONSTEXPR_AFTER_CXX11 const T &min(const T &a, const T &b, + Compare comp); -template T min(initializer_list t); +template _LIBCPP_CONSTEXPR_AFTER_CXX11 T min(initializer_list t); -template T min(initializer_list t, Compare comp); +template +_LIBCPP_CONSTEXPR_AFTER_CXX11 T min(initializer_list t, Compare comp); -template const T &clamp(const T &v, const T &lo, const T &hi); +#if _LIBCPP_STD_VER > 14 +template +const T constexpr &clamp(const T &v, const T &lo, const T &hi); template -const T &clamp(const T &v, const T &lo, const T &hi, Compare comp); +const T constexpr &clamp(const T &v, const T &lo, const T &hi, Compare comp); +#endif template -ForwardIterator max_element(ForwardIterator first, ForwardIterator last); +_LIBCPP_CONSTEXPR_AFTER_CXX11 ForwardIterator max_element(ForwardIterator first, + ForwardIterator last); template -ForwardIterator max_element(ForwardIterator first, ForwardIterator last, - Compare comp); +_LIBCPP_CONSTEXPR_AFTER_CXX11 ForwardIterator max_element(ForwardIterator first, + ForwardIterator last, + Compare comp); -template const T &max(const T &a, const T &b); +template +_LIBCPP_CONSTEXPR_AFTER_CXX11 const T &max(const T &a, const T &b); template -const T &max(const T &a, const T &b, Compare comp); +_LIBCPP_CONSTEXPR_AFTER_CXX11 const T &max(const T &a, const T &b, + Compare comp); -template T max(initializer_list t); +template _LIBCPP_CONSTEXPR_AFTER_CXX11 T max(initializer_list t); -template T max(initializer_list t, Compare comp); +template +_LIBCPP_CONSTEXPR_AFTER_CXX11 T max(initializer_list t, Compare comp); template -pair minmax_element(ForwardIterator first, - ForwardIterator last); +_LIBCPP_CONSTEXPR_AFTER_CXX11 pair +minmax_element(ForwardIterator first, ForwardIterator last); template -pair +pair _LIBCPP_CONSTEXPR_AFTER_CXX11 minmax_element(ForwardIterator first, ForwardIterator last, Compare comp); -template pair minmax(const T &a, const T &b); +template +_LIBCPP_CONSTEXPR_AFTER_CXX11 pair minmax(const T &a, + const T &b); template -pair minmax(const T &a, const T &b, Compare comp); +_LIBCPP_CONSTEXPR_AFTER_CXX11 pair +minmax(const T &a, const T &b, Compare comp); -template pair minmax(initializer_list t); +template +_LIBCPP_CONSTEXPR_AFTER_CXX11 pair minmax(initializer_list t); template -pair minmax(initializer_list t, Compare comp); +_LIBCPP_CONSTEXPR_AFTER_CXX11 pair minmax(initializer_list t, + Compare comp); template bool lexicographical_compare(InputIterator1 first1, InputIterator1 last1, diff --git a/clang/lib/Headers/sycl_wrappers/random b/clang/lib/Headers/sycl_wrappers/random index b03602c8a68ca..797a36693e180 100644 --- a/clang/lib/Headers/sycl_wrappers/random +++ b/clang/lib/Headers/sycl_wrappers/random @@ -1,10 +1,9 @@ // -*- C++ -*- //===--------------------------- random -----------------------------------===// // -// The LLVM Compiler Infrastructure -// -// This file is dual licensed under the MIT and the University of Illinois Open -// Source Licenses. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7ef90f00a3a8c..c7301d5970b6e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1,9 +1,8 @@ //===- SemaSYCL.cpp - Semantic Analysis for SYCL constructs ---------------===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// // This implements Semantic Analysis for SYCL constructs. @@ -21,6 +20,7 @@ #include "llvm/Support/FileSystem.h" #include "llvm/Support/Path.h" #include "llvm/Support/raw_ostream.h" +#include "clang/Analysis/CallGraph.h" #include @@ -45,6 +45,7 @@ enum RestrictKind { KernelRTTI, KernelNonConstStaticDataVariable, KernelCallVirtualFunction, + KernelCallRecursiveFunction, KernelCallFunctionPointer, KernelAllocateStorage, KernelUseExceptions, @@ -85,20 +86,25 @@ class MarkDeviceFunction : public RecursiveASTVisitor { bool VisitCallExpr(CallExpr *e) { for (const auto &Arg : e->arguments()) - CheckTypeForVirtual(Arg->getType(), Arg->getSourceRange()); + CheckSYCLType(Arg->getType(), Arg->getSourceRange()); if (FunctionDecl *Callee = e->getDirectCallee()) { + Callee = Callee->getCanonicalDecl(); // Remember that all SYCL kernel functions have deferred // instantiation as template functions. It means that // all functions used by kernel have already been parsed and have // definitions. + llvm::SmallPtrSet VisitedSet; + if (IsRecursive(Callee, Callee, VisitedSet)) + SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << + KernelCallRecursiveFunction; if (const CXXMethodDecl *Method = dyn_cast(Callee)) if (Method->isVirtual()) SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) << KernelCallVirtualFunction; - CheckTypeForVirtual(Callee->getReturnType(), Callee->getSourceRange()); + CheckSYCLType(Callee->getReturnType(), Callee->getSourceRange()); if (FunctionDecl *Def = Callee->getDefinition()) { if (!Def->hasAttr()) { @@ -116,7 +122,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { bool VisitCXXConstructExpr(CXXConstructExpr *E) { for (const auto &Arg : E->arguments()) - CheckTypeForVirtual(Arg->getType(), Arg->getSourceRange()); + CheckSYCLType(Arg->getType(), Arg->getSourceRange()); CXXConstructorDecl *Ctor = E->getConstructor(); @@ -150,22 +156,22 @@ class MarkDeviceFunction : public RecursiveASTVisitor { } bool VisitTypedefNameDecl(TypedefNameDecl *TD) { - CheckTypeForVirtual(TD->getUnderlyingType(), TD->getLocation()); + CheckSYCLType(TD->getUnderlyingType(), TD->getLocation()); return true; } bool VisitRecordDecl(RecordDecl *RD) { - CheckTypeForVirtual(QualType{RD->getTypeForDecl(), 0}, RD->getLocation()); + CheckSYCLType(QualType{RD->getTypeForDecl(), 0}, RD->getLocation()); return true; } bool VisitParmVarDecl(VarDecl *VD) { - CheckTypeForVirtual(VD->getType(), VD->getLocation()); + CheckSYCLType(VD->getType(), VD->getLocation()); return true; } bool VisitVarDecl(VarDecl *VD) { - CheckTypeForVirtual(VD->getType(), VD->getLocation()); + CheckSYCLType(VD->getType(), VD->getLocation()); return true; } @@ -180,7 +186,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { } bool VisitDeclRefExpr(DeclRefExpr *E) { - CheckTypeForVirtual(E->getType(), E->getSourceRange()); + CheckSYCLType(E->getType(), E->getSourceRange()); if (VarDecl *VD = dyn_cast(E->getDecl())) { bool IsConst = VD->getType().getNonReferenceType().isConstQualified(); if (!IsConst && VD->hasGlobalStorage() && !VD->isStaticLocal() && @@ -199,12 +205,17 @@ class MarkDeviceFunction : public RecursiveASTVisitor { // storage are disallowed in a SYCL kernel. The placement // new operator and any user-defined overloads that // do not allocate storage are permitted. - const FunctionDecl *FD = E->getOperatorNew(); - if (FD && !FD->isReservedGlobalPlacementOperator()) { - OverloadedOperatorKind Kind = FD->getOverloadedOperator(); - if (Kind == OO_New || Kind == OO_Array_New) - SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) << - KernelAllocateStorage; + if (FunctionDecl *FD = E->getOperatorNew()) { + if (FD->isReplaceableGlobalAllocationFunction()) { + SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) << + KernelAllocateStorage; + } else if (FunctionDecl *Def = FD->getDefinition()) { + if (!Def->hasAttr()) { + Def->addAttr(SYCLDeviceAttr::CreateImplicit(SemaRef.Context)); + this->TraverseStmt(Def->getBody()); + SemaRef.AddSyclKernel(Def); + } + } } return true; } @@ -245,8 +256,42 @@ class MarkDeviceFunction : public RecursiveASTVisitor { return true; } + // The call graph for this translation unit. + CallGraph SYCLCG; private: - bool CheckTypeForVirtual(QualType Ty, SourceRange Loc) { + // Determines whether the function FD is recursive. + // CalleeNode is a function which is called either directly + // or indirectly from FD. If recursion is detected then create + // diagnostic notes on each function as the callstack is unwound. + bool IsRecursive(FunctionDecl *CalleeNode, FunctionDecl *FD, + llvm::SmallPtrSet VisitedSet) { + // We're currently checking CalleeNode on a different + // trace through the CallGraph, we avoid infinite recursion + // by using VisitedSet to keep track of this. + if (!VisitedSet.insert(CalleeNode).second) + return false; + if (CallGraphNode *N = SYCLCG.getNode(CalleeNode)) { + for (const CallGraphNode *CI : *N) { + if (FunctionDecl *Callee = dyn_cast(CI->getDecl())) { + Callee = Callee->getCanonicalDecl(); + if (Callee == FD) + return SemaRef.Diag(FD->getSourceRange().getBegin(), + diag::note_sycl_recursive_function_declared_here) + << KernelCallRecursiveFunction; + else if (IsRecursive(Callee, FD, VisitedSet)) + return true; + } + } + } + return false; + } + + bool CheckSYCLType(QualType Ty, SourceRange Loc) { + if (Ty->isVariableArrayType()) { + SemaRef.Diag(Loc.getBegin(), diag::err_vla_unsupported); + return false; + } + while (Ty->isAnyPointerType() || Ty->isArrayType()) Ty = QualType{Ty->getPointeeOrArrayElementType(), 0}; @@ -264,25 +309,25 @@ class MarkDeviceFunction : public RecursiveASTVisitor { } for (const auto &Field : CRD->fields()) { - if (!CheckTypeForVirtual(Field->getType(), Field->getSourceRange())) { + if (!CheckSYCLType(Field->getType(), Field->getSourceRange())) { SemaRef.Diag(Loc.getBegin(), diag::note_sycl_used_here); return false; } } } else if (const auto *RD = Ty->getAsRecordDecl()) { for (const auto &Field : RD->fields()) { - if (!CheckTypeForVirtual(Field->getType(), Field->getSourceRange())) { + if (!CheckSYCLType(Field->getType(), Field->getSourceRange())) { SemaRef.Diag(Loc.getBegin(), diag::note_sycl_used_here); return false; } } } else if (const auto *FPTy = dyn_cast(Ty)) { for (const auto &ParamTy : FPTy->param_types()) - if (!CheckTypeForVirtual(ParamTy, Loc)) + if (!CheckSYCLType(ParamTy, Loc)) return false; - return CheckTypeForVirtual(FPTy->getReturnType(), Loc); + return CheckSYCLType(FPTy->getReturnType(), Loc); } else if (const auto *FTy = dyn_cast(Ty)) { - return CheckTypeForVirtual(FTy->getReturnType(), Loc); + return CheckSYCLType(FTy->getReturnType(), Loc); } return true; } @@ -347,6 +392,8 @@ static FunctionDecl *CreateSYCLKernelFunction(ASTContext &Context, SYCLKernel->addAttr(SYCLDeviceAttr::CreateImplicit(Context)); SYCLKernel->addAttr(OpenCLKernelAttr::CreateImplicit(Context)); SYCLKernel->addAttr(AsmLabelAttr::CreateImplicit(Context, Name)); + SYCLKernel->addAttr(ArtificialAttr::CreateImplicit(Context)); + // To see kernel in AST-dump. DC->addDecl(SYCLKernel); return SYCLKernel; @@ -726,6 +773,10 @@ void Sema::ConstructSYCLKernel(FunctionDecl *KernelCallerFunc) { AddSyclKernel(SYCLKernel); // Let's mark all called functions with SYCL Device attribute. MarkDeviceFunction Marker(*this); + // Create the call graph so we can detect recursion and check the validity + // of new operator overrides. Add the kernel function itself in case + // it is recursive. + Marker.SYCLCG.addToCallGraph(getASTContext().getTranslationUnitDecl()); Marker.TraverseStmt(SYCLKernelBody); } diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp new file mode 100644 index 0000000000000..2d2d8cc7dc4b9 --- /dev/null +++ b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp @@ -0,0 +1,29 @@ +// RUN: %clang --sycl %s -S -emit-llvm -g -o - | FileCheck %s +// +// Verify the SYCL kernel routine is marked artificial. +// +// Since it has no source correlation of its own, the SYCL kernel needs to be +// marked artificial or it will inherit source correlation from the surrounding +// code. +// + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); +} + +int main() { + int value = 0; + int* captured = &value; + kernel_single_task([=]() { + *captured = 1; + }); + return 0; +} + +// CHECK: define{{.*}} spir_kernel {{.*}}void @_ZTSZ4mainE15kernel_function(i32*{{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{ +// CHECK: [[FILE:![0-9]+]] = !DIFile(filename: "{{.*}}debug-info-srcpos-kernel.cpp"{{.*}}) +// CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "_ZTSZ4mainE15kernel_function" +// CHECK-SAME: scope: [[FILE]] +// CHECK-SAME: file: [[FILE]] +// CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped diff --git a/clang/test/CodeGenSYCL/spir-no-deref-attr.cpp b/clang/test/CodeGenSYCL/spir-no-deref-attr.cpp index 0fe8af83a6ad8..c87230d281525 100644 --- a/clang/test/CodeGenSYCL/spir-no-deref-attr.cpp +++ b/clang/test/CodeGenSYCL/spir-no-deref-attr.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-spirv -x c++ %s -o %t.spv +// RUN: %clang_cc1 -fno-sycl-use-bitcode -triple spir64-unknown-linux-sycldevice -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-spirv -x c++ %s -o %t.spv // RUN: llvm-spirv %t.spv -to-text -o %t.txt // RUN: FileCheck < %t.txt %s --check-prefix=CHECK diff --git a/clang/test/Driver/sycl-offload.c b/clang/test/Driver/sycl-offload.c index 24587615b854f..8708569f7e4ef 100644 --- a/clang/test/Driver/sycl-offload.c +++ b/clang/test/Driver/sycl-offload.c @@ -69,9 +69,12 @@ /// We should have an offload action joining the host compile and device /// preprocessor and another one joining the device linking outputs to the host /// action. The same graph should be generated when no -fsycl-targets is used +/// The same phase graph will be used with -fsycl-use-bitcode // RUN: %clang -ccc-print-phases -target x86_64-unknown-linux-gnu -fsycl -fsycl-targets=spir64-unknown-linux-sycldevice %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-PHASES %s -// RUN: %clang -ccc-print-phases -target x86_64-unknown-linux-gnu -fsycl %s 2>&1 \ +// RUN: %clang -ccc-print-phases -target x86_64-unknown-linux-gnu -fsycl -fno-sycl-use-bitcode %s 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-PHASES %s +// RUN: %clang -ccc-print-phases -target x86_64-unknown-linux-gnu -fsycl -fsycl-use-bitcode %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-PHASES %s // CHK-PHASES: 0: input, "[[INPUT:.+\.c]]", c, (host-sycl) // CHK-PHASES: 1: preprocessor, {0}, cpp-output, (host-sycl) @@ -214,10 +217,20 @@ /// ########################################################################### /// Check -fsycl-is-device is passed when compiling for the device. +/// also check for SPIR-V binary creation // RUN: %clang -### -no-canonical-prefixes -fsycl -fsycl-targets=spir64-unknown-linux-sycldevice %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-FSYCL-IS-DEVICE %s -// CHK-FSYCL-IS-DEVICE: clang{{.*}} "-fsycl-is-device" {{.*}}.c +// CHK-FSYCL-IS-DEVICE: clang{{.*}} "-fsycl-is-device" {{.*}} "-emit-spirv" {{.*}}.c + +/// ########################################################################### + +/// Check -fsycl-is-device and emitting to .spv when compiling for the device +/// when using -fno-sycl-use-bitcode +// RUN: %clang -### -fno-sycl-use-bitcode -fsycl -fsycl-targets=spir64-unknown-linux-sycldevice %s 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-FSYCL-IS-DEVICE-NO-BITCODE %s + +// CHK-FSYCL-IS-DEVICE-NO-BITCODE: clang{{.*}} "-fsycl-is-device" {{.*}} "-emit-spirv" {{.*}}.c /// ########################################################################### @@ -257,3 +270,4 @@ // CHK-ADD-TARGETS-UB: 3: input, "dummy.spv", sycl-fatbin, (device-sycl) // CHK-ADD-TARGETS-UB: 4: clang-offload-wrapper, {3}, object, (device-sycl) // CHK-ADD-TARGETS-UB: 5: offload, "host-sycl (x86_64-unknown-linux-gnu)" {2}, "device-sycl (spir64-unknown-linux-sycldevice)" {4}, image + diff --git a/clang/test/SemaSYCL/restrict-recursion.cpp b/clang/test/SemaSYCL/restrict-recursion.cpp new file mode 100644 index 0000000000000..e68f85789bf71 --- /dev/null +++ b/clang/test/SemaSYCL/restrict-recursion.cpp @@ -0,0 +1,107 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fsycl-is-device -Wno-return-type -verify -fsyntax-only -x c++ -emit-llvm-only -std=c++17 %s + +// This recursive function is not called from sycl kernel, +// so it should not be diagnosed. +int fib(int n) +{ + if (n <= 1) + return n; + return fib(n-1) + fib(n-2); +} + +typedef struct S { +template + // expected-note@+1 2{{function implemented using recursion declared here}} +T factT(T i, T j) +{ + // expected-error@+1 1{{SYCL kernel cannot call a recursive function}} + return factT(j,i); +} + +int fact(unsigned i) +{ + if (i==0) return 1; + // expected-error@+1 1{{SYCL kernel cannot call a recursive function}} + else return factT(i-1, i); +} +} S_type; + + + // expected-note@+1 2{{function implemented using recursion declared here}} +int fact(unsigned i); + // expected-note@+1 2{{function implemented using recursion declared here}} +int fact1(unsigned i) +{ + if (i==0) return 1; + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + else return fact(i-1) * i; +} +int fact(unsigned i) +{ + if (i==0) return 1; + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + else return fact1(i-1) * i; +} + +bool isa_B(void) { + S_type s; + + unsigned f = s.fact(3); + // expected-error@+1 1{{SYCL kernel cannot call a recursive function}} + unsigned f1 = s.factT(3,4); + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + unsigned g = fact(3); + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + unsigned g1 = fact1(3); + return 0; +} + +__attribute__((sycl_kernel)) void kernel1(void) { + isa_B(); +} + // expected-note@+1 2{{function implemented using recursion declared here}} +__attribute__((sycl_kernel)) void kernel2(void) { + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + kernel2(); +} +__attribute__((sycl_kernel)) void kernel3(void) { + ; +} + +using myFuncDef = int(int,int); + +void usage( myFuncDef functionPtr ) { + kernel1(); +} +void usage2( myFuncDef functionPtr ) { + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + kernel2(); +} +void usage3( myFuncDef functionPtr ) { + kernel3(); +} + +int addInt(int n, int m) { + return n+m; +} + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); +} + +template + // expected-note@+1 2{{function implemented using recursion declared here}} +__attribute__((sycl_kernel)) void kernel_single_task2(Func kernelFunc) { + kernelFunc(); + // expected-error@+1 2{{SYCL kernel cannot call a recursive function}} + kernel_single_task2(kernelFunc); +} + +int main() { + kernel_single_task([]() { usage( &addInt ); }); + kernel_single_task([]() { usage2( &addInt ); }); + kernel_single_task2([]() { usage3( &addInt ); }); + return fib(5); +} + diff --git a/clang/test/SemaSYCL/restrict-recursion2.cpp b/clang/test/SemaSYCL/restrict-recursion2.cpp new file mode 100644 index 0000000000000..3dff0a91ada3e --- /dev/null +++ b/clang/test/SemaSYCL/restrict-recursion2.cpp @@ -0,0 +1,81 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fsycl-is-device -Wno-return-type -verify -fsyntax-only -x c++ -emit-llvm-only -std=c++17 %s + +// This recursive function is not called from sycl kernel, +// so it should not be diagnosed. +int fib(int n) +{ + if (n <= 1) + return n; + return fib(n-1) + fib(n-2); +} + +typedef struct S { +template + // expected-note@+1 2{{function implemented using recursion declared here}} +T factT(T i, T j) +{ + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + return factT(j,i); +} + +int fact(unsigned i) +{ + if (i==0) return 1; + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + else return factT(i-1, i); +} +} S_type; + + + // expected-note@+1 2{{function implemented using recursion declared here}} +int fact(unsigned i); + // expected-note@+1 2{{function implemented using recursion declared here}} +int fact1(unsigned i) +{ + if (i==0) return 1; + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + else return fact(i-1) * i; +} +int fact(unsigned i) +{ + if (i==0) return 1; + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + else return fact1(i-1) * i; +} + +bool isa_B(void) { + S_type s; + + unsigned f = s.fact(3); + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + unsigned f1 = s.factT(3,4); + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + unsigned g = fact(3); + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + unsigned g1 = fact1(3); + return 0; +} + +void kernel1(void) { + isa_B(); +} + +using myFuncDef = int(int,int); + +void usage(myFuncDef functionPtr) { + kernel1(); +} + +int addInt(int n, int m) { + return n+m; +} + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); +} + +int main() { + kernel_single_task([]() {usage(&addInt);}); + return fib(5); +} diff --git a/clang/test/SemaSYCL/restrict-recursion3.cpp b/clang/test/SemaSYCL/restrict-recursion3.cpp new file mode 100644 index 0000000000000..c8af218965a92 --- /dev/null +++ b/clang/test/SemaSYCL/restrict-recursion3.cpp @@ -0,0 +1,39 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fsycl-is-device -Wno-return-type -verify -fsyntax-only -x c++ -emit-llvm-only -std=c++17 %s + +// This recursive function is not called from sycl kernel, +// so it should not be diagnosed. +int fib(int n) +{ + if (n <= 1) + return n; + return fib(n-1) + fib(n-2); +} + +void kernel3(void) { + ; +} + +using myFuncDef = int(int,int); + +void usage3(myFuncDef functionPtr) { + kernel3(); +} + +int addInt(int n, int m) { + return n+m; +} + +template + // expected-note@+1 2{{function implemented using recursion declared here}} +__attribute__((sycl_kernel)) void kernel_single_task2(Func kernelFunc) { + kernelFunc(); + // expected-error@+1 2{{SYCL kernel cannot allocate storage}} + int *ip = new int; + // expected-error@+1 2{{SYCL kernel cannot call a recursive function}} + kernel_single_task2(kernelFunc); +} + +int main() { + kernel_single_task2([]() { usage3( &addInt ); }); + return fib(5); +} diff --git a/clang/test/SemaSYCL/restrict-recursion4.cpp b/clang/test/SemaSYCL/restrict-recursion4.cpp new file mode 100644 index 0000000000000..ec52e8c1fdf1c --- /dev/null +++ b/clang/test/SemaSYCL/restrict-recursion4.cpp @@ -0,0 +1,39 @@ +// RUN: %clang_cc1 -fcxx-exceptions -fsycl-is-device -Wno-return-type -verify -fsyntax-only -x c++ -emit-llvm-only -std=c++17 %s + +// This recursive function is not called from sycl kernel, +// so it should not be diagnosed. +int fib(int n) +{ + if (n <= 1) + return n; + return fib(n-1) + fib(n-2); +} + + // expected-note@+1 2{{function implemented using recursion declared here}} +void kernel2(void) { + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + kernel2(); +} + +using myFuncDef = int(int,int); + +void usage2(myFuncDef functionPtr) { + // expected-error@+1 {{SYCL kernel cannot call a recursive function}} + kernel2(); +} + +int addInt(int n, int m) { + return n+m; +} + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + // expected-error@+1 {{SYCL kernel cannot allocate storage}} + int *ip = new int; + kernelFunc(); +} + +int main() { + kernel_single_task([]() {usage2(&addInt);}); + return fib(5); +} diff --git a/clang/test/SemaSYCL/sycl-restrict.cpp b/clang/test/SemaSYCL/sycl-restrict.cpp index c0cb0ae8af82e..bea75fa2c5431 100644 --- a/clang/test/SemaSYCL/sycl-restrict.cpp +++ b/clang/test/SemaSYCL/sycl-restrict.cpp @@ -5,6 +5,17 @@ namespace std { class type_info; typedef __typeof__(sizeof(int)) size_t; } + +namespace Check_VLA_Restriction { +void no_restriction(int p) { + int index[p+2]; +} +void restriction(int p) { + // expected-error@+1 {{variable length arrays are not supported for the current target}} + int index[p+2]; +} +} + void* operator new (std::size_t size, void* ptr) throw() { return ptr; }; namespace Check_RTTI_Restriction { // expected-error@+1 5{{No class with a vtable can be used in a SYCL kernel or any code included in the kernel}} @@ -17,8 +28,12 @@ struct B : public A { }; struct OverloadedNewDelete { - // These overloads do not allocate. - void *operator new(std::size_t size) throw() {return 0;} + // This overload allocates storage, give diagnostic. + void *operator new(std::size_t size) throw() { + // expected-error@+1 {{SYCL kernel cannot allocate storage}} + float *pt = new float; + return 0;} + // This overload does not allocate: no diagnostic. void *operator new[](std::size_t size) throw() {return 0;} void operator delete(void *){}; void operator delete[](void *){}; @@ -26,12 +41,12 @@ struct OverloadedNewDelete { bool isa_B(A *a) { + Check_VLA_Restriction::restriction(7); // expected-error@+1 {{SYCL kernel cannot allocate storage}} int *ip = new int; int i; int *p3 = new(&i) int; // no error on placement new - //FIXME call to overloaded new should not get error message - // expected-error@+1 {{SYCL kernel cannot allocate storage}} OverloadedNewDelete *x = new( struct OverloadedNewDelete ); + auto y = new struct OverloadedNewDelete [5]; // expected-error@+1 {{SYCL kernel cannot use rtti}} (void)typeid(int); // expected-error@+2 {{SYCL kernel cannot use rtti}} diff --git a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp index c5b215a73e78a..16bd85fd02849 100644 --- a/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp +++ b/clang/tools/clang-offload-wrapper/ClangOffloadWrapper.cpp @@ -1,9 +1,8 @@ //===-- clang-offload-wrapper/ClangOffloadWrapper.cpp ---------------------===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// /// diff --git a/llvm-spirv/LICENSE2.TXT b/llvm-spirv/LICENSE2.TXT deleted file mode 100644 index f9dc50615d7ec..0000000000000 --- a/llvm-spirv/LICENSE2.TXT +++ /dev/null @@ -1,219 +0,0 @@ - Apache License - Version 2.0, January 2004 - http://www.apache.org/licenses/ - - TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION - - 1. Definitions. - - "License" shall mean the terms and conditions for use, reproduction, - and distribution as defined by Sections 1 through 9 of this document. - - "Licensor" shall mean the copyright owner or entity authorized by - the copyright owner that is granting the License. - - "Legal Entity" shall mean the union of the acting entity and all - other entities that control, are controlled by, or are under common - control with that entity. For the purposes of this definition, - "control" means (i) the power, direct or indirect, to cause the - direction or management of such entity, whether by contract or - otherwise, or (ii) ownership of fifty percent (50%) or more of the - outstanding shares, or (iii) beneficial ownership of such entity. - - "You" (or "Your") shall mean an individual or Legal Entity - exercising permissions granted by this License. - - "Source" form shall mean the preferred form for making modifications, - including but not limited to software source code, documentation - source, and configuration files. - - "Object" form shall mean any form resulting from mechanical - transformation or translation of a Source form, including but - not limited to compiled object code, generated documentation, - and conversions to other media types. - - "Work" shall mean the work of authorship, whether in Source or - Object form, made available under the License, as indicated by a - copyright notice that is included in or attached to the work - (an example is provided in the Appendix below). - - "Derivative Works" shall mean any work, whether in Source or Object - form, that is based on (or derived from) the Work and for which the - editorial revisions, annotations, elaborations, or other modifications - represent, as a whole, an original work of authorship. For the purposes - of this License, Derivative Works shall not include works that remain - separable from, or merely link (or bind by name) to the interfaces of, - the Work and Derivative Works thereof. - - "Contribution" shall mean any work of authorship, including - the original version of the Work and any modifications or additions - to that Work or Derivative Works thereof, that is intentionally - submitted to Licensor for inclusion in the Work by the copyright owner - or by an individual or Legal Entity authorized to submit on behalf of - the copyright owner. For the purposes of this definition, "submitted" - means any form of electronic, verbal, or written communication sent - to the Licensor or its representatives, including but not limited to - communication on electronic mailing lists, source code control systems, - and issue tracking systems that are managed by, or on behalf of, the - Licensor for the purpose of discussing and improving the Work, but - excluding communication that is conspicuously marked or otherwise - designated in writing by the copyright owner as "Not a Contribution." - - "Contributor" shall mean Licensor and any individual or Legal Entity - on behalf of whom a Contribution has been received by Licensor and - subsequently incorporated within the Work. - - 2. Grant of Copyright License. Subject to the terms and conditions of - this License, each Contributor hereby grants to You a perpetual, - worldwide, non-exclusive, no-charge, royalty-free, irrevocable - copyright license to reproduce, prepare Derivative Works of, - publicly display, publicly perform, sublicense, and distribute the - Work and such Derivative Works in Source or Object form. - - 3. Grant of Patent License. Subject to the terms and conditions of - this License, each Contributor hereby grants to You a perpetual, - worldwide, non-exclusive, no-charge, royalty-free, irrevocable - (except as stated in this section) patent license to make, have made, - use, offer to sell, sell, import, and otherwise transfer the Work, - where such license applies only to those patent claims licensable - by such Contributor that are necessarily infringed by their - Contribution(s) alone or by combination of their Contribution(s) - with the Work to which such Contribution(s) was submitted. If You - institute patent litigation against any entity (including a - cross-claim or counterclaim in a lawsuit) alleging that the Work - or a Contribution incorporated within the Work constitutes direct - or contributory patent infringement, then any patent licenses - granted to You under this License for that Work shall terminate - as of the date such litigation is filed. - - 4. Redistribution. You may reproduce and distribute copies of the - Work or Derivative Works thereof in any medium, with or without - modifications, and in Source or Object form, provided that You - meet the following conditions: - - (a) You must give any other recipients of the Work or - Derivative Works a copy of this License; and - - (b) You must cause any modified files to carry prominent notices - stating that You changed the files; and - - (c) You must retain, in the Source form of any Derivative Works - that You distribute, all copyright, patent, trademark, and - attribution notices from the Source form of the Work, - excluding those notices that do not pertain to any part of - the Derivative Works; and - - (d) If the Work includes a "NOTICE" text file as part of its - distribution, then any Derivative Works that You distribute must - include a readable copy of the attribution notices contained - within such NOTICE file, excluding those notices that do not - pertain to any part of the Derivative Works, in at least one - of the following places: within a NOTICE text file distributed - as part of the Derivative Works; within the Source form or - documentation, if provided along with the Derivative Works; or, - within a display generated by the Derivative Works, if and - wherever such third-party notices normally appear. The contents - of the NOTICE file are for informational purposes only and - do not modify the License. You may add Your own attribution - notices within Derivative Works that You distribute, alongside - or as an addendum to the NOTICE text from the Work, provided - that such additional attribution notices cannot be construed - as modifying the License. - - You may add Your own copyright statement to Your modifications and - may provide additional or different license terms and conditions - for use, reproduction, or distribution of Your modifications, or - for any such Derivative Works as a whole, provided Your use, - reproduction, and distribution of the Work otherwise complies with - the conditions stated in this License. - - 5. Submission of Contributions. Unless You explicitly state otherwise, - any Contribution intentionally submitted for inclusion in the Work - by You to the Licensor shall be under the terms and conditions of - this License, without any additional terms or conditions. - Notwithstanding the above, nothing herein shall supersede or modify - the terms of any separate license agreement you may have executed - with Licensor regarding such Contributions. - - 6. Trademarks. This License does not grant permission to use the trade - names, trademarks, service marks, or product names of the Licensor, - except as required for reasonable and customary use in describing the - origin of the Work and reproducing the content of the NOTICE file. - - 7. Disclaimer of Warranty. Unless required by applicable law or - agreed to in writing, Licensor provides the Work (and each - Contributor provides its Contributions) on an "AS IS" BASIS, - WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or - implied, including, without limitation, any warranties or conditions - of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A - PARTICULAR PURPOSE. You are solely responsible for determining the - appropriateness of using or redistributing the Work and assume any - risks associated with Your exercise of permissions under this License. - - 8. Limitation of Liability. In no event and under no legal theory, - whether in tort (including negligence), contract, or otherwise, - unless required by applicable law (such as deliberate and grossly - negligent acts) or agreed to in writing, shall any Contributor be - liable to You for damages, including any direct, indirect, special, - incidental, or consequential damages of any character arising as a - result of this License or out of the use or inability to use the - Work (including but not limited to damages for loss of goodwill, - work stoppage, computer failure or malfunction, or any and all - other commercial damages or losses), even if such Contributor - has been advised of the possibility of such damages. - - 9. Accepting Warranty or Additional Liability. While redistributing - the Work or Derivative Works thereof, You may choose to offer, - and charge a fee for, acceptance of support, warranty, indemnity, - or other liability obligations and/or rights consistent with this - License. However, in accepting such obligations, You may act only - on Your own behalf and on Your sole responsibility, not on behalf - of any other Contributor, and only if You agree to indemnify, - defend, and hold each Contributor harmless for any liability - incurred by, or claims asserted against, such Contributor by reason - of your accepting any such warranty or additional liability. - - END OF TERMS AND CONDITIONS - - APPENDIX: How to apply the Apache License to your work. - - To apply the Apache License to your work, attach the following - boilerplate notice, with the fields enclosed by brackets "[]" - replaced with your own identifying information. (Don't include - the brackets!) The text should be enclosed in the appropriate - comment syntax for the file format. We also recommend that a - file or class name and description of purpose be included on the - same "printed page" as the copyright notice for easier - identification within third-party archives. - - Copyright [yyyy] [name of copyright owner] - - Licensed under the Apache License, Version 2.0 (the "License"); - you may not use this file except in compliance with the License. - You may obtain a copy of the License at - - http://www.apache.org/licenses/LICENSE-2.0 - - Unless required by applicable law or agreed to in writing, software - distributed under the License is distributed on an "AS IS" BASIS, - WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. - See the License for the specific language governing permissions and - limitations under the License. - - ---- LLVM Exceptions to the Apache 2.0 License ---- - -As an exception, if, as a result of your compiling your source code, portions -of this Software are embedded into an Object form of such source code, you -may redistribute such embedded portions in such Object form without complying -with the conditions of Sections 4(a), 4(b) and 4(d) of the License. - -In addition, if you combine or link compiled forms of this Software with -software that is licensed under the GPLv2 ("Combined Software") and if a -court of competent jurisdiction determines that the patent provision (Section -3), the indemnity provision (Section 9) or other Section of the License -conflicts with the conditions of the GPLv2, you may retroactively and -prospectively choose to deem waived or otherwise exclude such Section(s) of -the License, but only in their entirety and only with respect to the Combined -Software. - diff --git a/llvm/include/llvm/SYCL/ASFixer.h b/llvm/include/llvm/SYCL/ASFixer.h index 3b515b2e68c94..9dcf322df581a 100644 --- a/llvm/include/llvm/SYCL/ASFixer.h +++ b/llvm/include/llvm/SYCL/ASFixer.h @@ -1,9 +1,8 @@ //===- ASFixer.h - SYCL address spaces fixer pass -------------------------===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// // diff --git a/llvm/lib/SYCL/ASFixer.cpp b/llvm/lib/SYCL/ASFixer.cpp index ac67c28c656c1..235f2d539f1b7 100644 --- a/llvm/lib/SYCL/ASFixer.cpp +++ b/llvm/lib/SYCL/ASFixer.cpp @@ -1,10 +1,8 @@ //===- ASFixer.cpp - Address spaces fixer pass implementation -------------===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. -// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception// //===----------------------------------------------------------------------===// // // This implements the address spaces fixer pass. diff --git a/sycl/.gitignore b/sycl/.gitignore new file mode 100644 index 0000000000000..6f61542267633 --- /dev/null +++ b/sycl/.gitignore @@ -0,0 +1 @@ +include/CL/sycl/version.hpp diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 9e1fe82ebc106..35879344cbd2f 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -34,6 +34,12 @@ find_package(OpenCL REQUIRED) include_directories(${OpenCL_INCLUDE_DIRS}) link_libraries(OpenCL) +# Configure SYCL version macro +set(sycl_inc_dir ${CMAKE_CURRENT_SOURCE_DIR}/include/CL) +string(TIMESTAMP __SYCL_COMPILER_VERSION "%Y%m%d") +set(version_header "${sycl_inc_dir}/sycl/version.hpp") +configure_file("${version_header}.in" "${version_header}") + # Copy SYCL headers set(sycl_inc_dir ${CMAKE_CURRENT_SOURCE_DIR}/include/CL) set(dst_dir ${LLVM_LIBRARY_OUTPUT_INTDIR}/clang/${CLANG_VERSION}/include/CL) @@ -68,6 +74,7 @@ set(OPENCL_INCLUDE "${OpenCL_INCLUDE_DIRS}") add_library("${SYCLLibrary}" SHARED "${includeRootPath}/CL/sycl.hpp" "${sourceRootPath}/detail/common.cpp" + "${sourceRootPath}/detail/context_impl.cpp" "${sourceRootPath}/detail/device_info.cpp" "${sourceRootPath}/detail/event_impl.cpp" "${sourceRootPath}/detail/force_device.cpp" diff --git a/sycl/doc/GetStartedWithSYCLCompiler.md b/sycl/doc/GetStartedWithSYCLCompiler.md index 8fe8613a342f2..2b749cd62c6a9 100644 --- a/sycl/doc/GetStartedWithSYCLCompiler.md +++ b/sycl/doc/GetStartedWithSYCLCompiler.md @@ -132,7 +132,7 @@ The SYCL Compiler supports two types of compilation: a. Compile the device code from the C++ file into the SPIR-V file: ```bash - clang++ --sycl -Xclang -fsycl-int-header=simple-sycl-app-int-header.h -c simple-sycl-app.cpp -o kernel.spv + clang++ --sycl -fno-sycl-use-bitcode -Xclang -fsycl-int-header=simple-sycl-app-int-header.h -c simple-sycl-app.cpp -o kernel.spv # NOTE: The section "-Xclang -fsycl-int-header=simple-sycl-app-int-header.h" # generates `integration header` file. # This file must be included for the host side compilation. diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 3197620f41e6e..6e7cb9077c568 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -1,9 +1,8 @@ //==---------- spirv_ops.hpp --- SPIRV operations -------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -110,6 +109,54 @@ __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned long long) __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Min) __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max) +extern bool OpGroupAll(int32_t Scope, bool Predicate) noexcept; + +extern bool OpGroupAny(int32_t Scope, bool Predicate) noexcept; + +template +extern dataT OpGroupBroadcast(int32_t Scope, dataT Value, + uint32_t LocalId) noexcept; + +template +extern dataT OpGroupIAdd(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpGroupFAdd(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpGroupUMin(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpGroupSMin(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpGroupFMin(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpGroupUMax(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpGroupSMax(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpGroupFMax(int32_t Scope, int32_t Op, dataT Value) noexcept; +template +extern dataT OpSubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept; +template +extern dataT OpSubgroupShuffleDownINTEL(dataT Current, dataT Next, + uint32_t Delta) noexcept; +template +extern dataT OpSubgroupShuffleUpINTEL(dataT Previous, dataT Current, + uint32_t Delta) noexcept; +template +extern dataT OpSubgroupShuffleXorINTEL(dataT Data, uint32_t Value) noexcept; + +template +extern dataT OpSubgroupBlockReadINTEL(const __global uint16_t *Ptr) noexcept; + +template +extern void OpSubgroupBlockWriteINTEL(__global uint16_t *Ptr, + dataT Data) noexcept; + +template +extern dataT OpSubgroupBlockReadINTEL(const __global uint32_t *Ptr) noexcept; + +template +extern void OpSubgroupBlockWriteINTEL(__global uint32_t *Ptr, + dataT Data) noexcept; #else template diff --git a/sycl/include/CL/__spirv/spirv_types.hpp b/sycl/include/CL/__spirv/spirv_types.hpp index 8e2d6bfa357b2..b8e38fbd6d688 100644 --- a/sycl/include/CL/__spirv/spirv_types.hpp +++ b/sycl/include/CL/__spirv/spirv_types.hpp @@ -1,9 +1,8 @@ //===----------- spirv_types.hpp --- SPIRV types -------------------------===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -44,5 +43,6 @@ enum MemorySemantics { // Only in such cases the class is recognized as SPIRV type OpTypeEvent. class OpTypeEvent; +enum GroupOperation { Reduce = 0, InclusiveScan = 1, ExclusiveScan = 2 }; } // namespace __spirv } // namespace cl diff --git a/sycl/include/CL/__spirv/spirv_vars.hpp b/sycl/include/CL/__spirv/spirv_vars.hpp new file mode 100644 index 0000000000000..7c9a6496dba10 --- /dev/null +++ b/sycl/include/CL/__spirv/spirv_vars.hpp @@ -0,0 +1,48 @@ +//==---------- spirv_vars.hpp --- SPIRV variables -------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// ===-------------------------------------------------------------------=== // + +#pragma once + +#ifdef __SYCL_DEVICE_ONLY__ + +namespace cl { +namespace __spirv { +typedef size_t size_t_vec __attribute__((ext_vector_type(3))); + +extern const __constant size_t_vec VarGlobalSize; +extern const __constant size_t_vec VarGlobalInvocationId; +extern const __constant size_t_vec VarWorkgroupSize; +extern const __constant size_t_vec VarLocalInvocationId; +extern const __constant size_t_vec VarWorkgroupId; +extern const __constant size_t_vec VarGlobalOffset; + +#define DEFINE_INT_ID_TO_XYZ_CONVERTER(POSTFIX) \ + template static size_t get##POSTFIX(); \ + template <> static size_t get##POSTFIX<0>() { return Var##POSTFIX.x; } \ + template <> static size_t get##POSTFIX<1>() { return Var##POSTFIX.y; } \ + template <> static size_t get##POSTFIX<2>() { return Var##POSTFIX.z; } + +DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalSize); +DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalInvocationId) +DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupSize) +DEFINE_INT_ID_TO_XYZ_CONVERTER(LocalInvocationId) +DEFINE_INT_ID_TO_XYZ_CONVERTER(WorkgroupId) +DEFINE_INT_ID_TO_XYZ_CONVERTER(GlobalOffset) + +#undef DEFINE_INT_ID_TO_XYZ_CONVERTER + +extern const __constant uint32_t VarSubgroupSize; +extern const __constant uint32_t VarSubgroupMaxSize; +extern const __constant uint32_t VarNumSubgroups; +extern const __constant uint32_t VarNumEnqueuedSubgroups; +extern const __constant uint32_t VarSubgroupId; +extern const __constant uint32_t VarSubgroupLocalInvocationId; + +} // namespace __spirv +} // namespace cl +#endif // __SYCL_DEVICE_ONLY__ diff --git a/sycl/include/CL/sycl.hpp b/sycl/include/CL/sycl.hpp index 131d1863ba6e3..6de33e8410364 100644 --- a/sycl/include/CL/sycl.hpp +++ b/sycl/include/CL/sycl.hpp @@ -1,9 +1,8 @@ //==------------ sycl.hpp - SYCL standard header file ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -35,6 +34,7 @@ #include #include #include +#include // Do not include RT only function implementations for device code as it leads // to problem. Should be finally fixed when we introduce library. diff --git a/sycl/include/CL/sycl/access/access.hpp b/sycl/include/CL/sycl/access/access.hpp index 2c844e22285a5..ae3944ca427a9 100644 --- a/sycl/include/CL/sycl/access/access.hpp +++ b/sycl/include/CL/sycl/access/access.hpp @@ -1,9 +1,8 @@ //==---------------- access.hpp --- SYCL access ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #pragma once diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp index 0f865ccd40aa9..413196a9cf076 100644 --- a/sycl/include/CL/sycl/accessor.hpp +++ b/sycl/include/CL/sycl/accessor.hpp @@ -1,9 +1,8 @@ //==--------- accessor.hpp --- SYCL accessor -------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -43,7 +42,7 @@ class subscript_obj { subscript_obj(const accessor_t &acc, cl::sycl::id &indexes) : accRef(acc), ids(indexes) {} - INLINE_IF_DEVICE subscript_obj operator[](size_t index) { ids[accessorDim - dimensions] = index; @@ -66,7 +65,7 @@ class subscript_obj &indexes) : accRef(acc), ids(indexes) {} - INLINE_IF_DEVICE dataT &operator[](size_t index) { + dataT &operator[](size_t index) { ids[accessorDim - 1] = index; return accRef.__impl()->Data[getOffsetForId( accRef.__impl()->Range, ids, accRef.__impl()->Offset)]; @@ -87,7 +86,7 @@ class subscript_obj &indexes) : accRef(acc), ids(indexes) {} - INLINE_IF_DEVICE typename detail::remove_AS::type + typename detail::remove_AS::type operator[](size_t index) { ids[accessorDim - 1] = index; return accRef.__impl()->Data[getOffsetForId( @@ -118,7 +117,7 @@ SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions == 0) { accessor_impl(dataT *Data) : Data(Data) {} // Returns the number of accessed elements. - INLINE_IF_DEVICE size_t get_count() const { return 1; } + size_t get_count() const { return 1; } }; /// Implementation of host accessor. @@ -133,7 +132,7 @@ SYCL_ACCESSOR_IMPL(isTargetHostAccess(accessTarget) && dimensions > 0) { : Data(Data), Range(Range), Offset(Offset) {} // Returns the number of accessed elements. - INLINE_IF_DEVICE size_t get_count() const { return Range.size(); } + size_t get_count() const { return Range.size(); } }; /// Implementation of device (kernel) accessor providing access to a single @@ -163,7 +162,7 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) && {} // Returns the number of accessed elements. - INLINE_IF_DEVICE size_t get_count() const { return 1; } + size_t get_count() const { return 1; } static_assert( std::is_same::type, @@ -201,7 +200,7 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) && {} // Returns the number of accessed elements. - INLINE_IF_DEVICE size_t get_count() const { return Range.size(); } + size_t get_count() const { return Range.size(); } static_assert( std::is_same::type, @@ -240,7 +239,7 @@ SYCL_ACCESSOR_IMPL(accessTarget == access::target::local && } // Returns the number of accessed elements. - INLINE_IF_DEVICE size_t get_count() const { return 1; } + size_t get_count() const { return 1; } static_assert( std::is_same::type, @@ -285,7 +284,7 @@ SYCL_ACCESSOR_IMPL(accessTarget == access::target::local && } // Returns the number of accessed elements. - INLINE_IF_DEVICE size_t get_count() const { return Range.size(); } + size_t get_count() const { return Range.size(); } static_assert( std::is_same::type, @@ -307,11 +306,11 @@ class accessor_base { using _ImplT = accessor_impl; - INLINE_IF_DEVICE const _ImplT *__impl() const { + const _ImplT *__impl() const { return reinterpret_cast(this); } - INLINE_IF_DEVICE _ImplT *__impl() { return reinterpret_cast<_ImplT *>(this); } + _ImplT *__impl() { return reinterpret_cast<_ImplT *>(this); } static_assert( std::is_same::type, @@ -339,21 +338,21 @@ class accessor_base { SYCL_ACCESSOR_SUBCLASS(accessor_common, accessor_base, true /* always */) { // Returns true if the current accessor is a placeholder accessor. - INLINE_IF_DEVICE constexpr bool is_placeholder() const { + constexpr bool is_placeholder() const { return isPlaceholder == access::placeholder::true_t; } // Returns the size of the accessed memory in bytes. - INLINE_IF_DEVICE size_t get_size() const { return this->get_count() * sizeof(dataT); } + size_t get_size() const { return this->get_count() * sizeof(dataT); } // Returns the number of accessed elements. - INLINE_IF_DEVICE size_t get_count() const { return this->__impl()->get_count(); } + size_t get_count() const { return this->__impl()->get_count(); } - template INLINE_IF_DEVICE + template typename std::enable_if<(Dimensions > 0), range>::type get_range() const { return this->__impl()->Range; } - template INLINE_IF_DEVICE + template typename std::enable_if<(Dimensions > 0), id>::type get_offset() const { return this->__impl()->Offset; } }; @@ -364,7 +363,7 @@ SYCL_ACCESSOR_SUBCLASS(accessor_opdata_w, accessor_common, accessMode == access::mode::discard_write || accessMode == access::mode::discard_read_write) && dimensions == 0) { - INLINE_IF_DEVICE operator dataT &() const { + operator dataT &() const { return this->__impl()->Data[0]; } }; @@ -382,7 +381,7 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_wn, accessor_opdata_w, subscript_obj - INLINE_IF_DEVICE operator[](size_t index) const { + operator[](size_t index) const { id ids; ids[0] = index; return subscript_obj) once again. - INLINE_IF_DEVICE dataT &operator[](id index) const { + dataT &operator[](id index) const { return this->operator[]( getOffsetForId(this->get_range(), index, this->get_offset())); } - INLINE_IF_DEVICE dataT &operator[](size_t index) const { + dataT &operator[](size_t index) const { return this->__impl()->Data[index]; } }; @@ -560,7 +559,7 @@ class accessor // implementation. _ImplT __impl; - INLINE_IF_DEVICE void __init(_ValueType *Ptr, range Range, + void __init(_ValueType *Ptr, range Range, id Offset) { __impl.Data = Ptr; __impl.Range = Range; diff --git a/sycl/include/CL/sycl/atomic.hpp b/sycl/include/CL/sycl/atomic.hpp index ab06af9233302..fffff6ca2d8e6 100644 --- a/sycl/include/CL/sycl/atomic.hpp +++ b/sycl/include/CL/sycl/atomic.hpp @@ -1,9 +1,8 @@ //==---------------- atomic.hpp - SYCL atomics -----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp index 5cd1e5109ced5..f06a4aff59dfe 100644 --- a/sycl/include/CL/sycl/buffer.hpp +++ b/sycl/include/CL/sycl/buffer.hpp @@ -1,9 +1,8 @@ //==----------- buffer.hpp --- SYCL buffer ---------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/context.hpp b/sycl/include/CL/sycl/context.hpp index a53933042a832..aceaa4cb2d9ec 100644 --- a/sycl/include/CL/sycl/context.hpp +++ b/sycl/include/CL/sycl/context.hpp @@ -1,34 +1,32 @@ //==---------------- context.hpp - SYCL context ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #pragma once -#include -#include -#include -#include +#include +#include +#include +#include #include #include -#include // 4.6.2 Context class namespace cl { namespace sycl { +// Forward declarations +class device; +class platform; class context { public: - explicit context(const async_handler &asyncHandler = {}) - : context(default_selector().select_device(), asyncHandler) {} + explicit context(const async_handler &asyncHandler = {}); - context(const device &dev, async_handler asyncHandler = {}) - : context(vector_class(1, dev), asyncHandler) {} + context(const device &dev, async_handler asyncHandler = {}); - context(const platform &plt, async_handler asyncHandler = {}) - : context(plt.get_devices(), asyncHandler) {} + context(const platform &plt, async_handler asyncHandler = {}); context(const vector_class &deviceList, async_handler asyncHandler = {}); @@ -37,9 +35,7 @@ class context { template typename info::param_traits::return_type - get_info() const { - return impl->get_info(); - } + get_info() const; context(const context &rhs) = default; @@ -49,17 +45,17 @@ class context { context &operator=(context &&rhs) = default; - bool operator==(const context &rhs) const { return impl == rhs.impl; } + bool operator==(const context &rhs) const; - bool operator!=(const context &rhs) const { return !(*this == rhs); } + bool operator!=(const context &rhs) const; - cl_context get() const { return impl->get(); } + cl_context get() const; - bool is_host() const { return impl->is_host(); } + bool is_host() const; - platform get_platform() const { return impl->get_platform(); } + platform get_platform() const; - vector_class get_devices() const { return impl->get_devices(); } + vector_class get_devices() const; private: std::shared_ptr impl; diff --git a/sycl/include/CL/sycl/detail/array.hpp b/sycl/include/CL/sycl/detail/array.hpp index b8cf259dd0bd7..0659fbdaffe9d 100644 --- a/sycl/include/CL/sycl/detail/array.hpp +++ b/sycl/include/CL/sycl/detail/array.hpp @@ -1,9 +1,8 @@ //==-------- array.hpp --- SYCL common iteration object ---------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -21,29 +20,29 @@ namespace detail { template class array { public: - INLINE_IF_DEVICE array() : common_array{0} {} + array() : common_array{0} {} /* The following constructor is only available in the array struct * specialization where: dimensions==1 */ - template INLINE_IF_DEVICE + template array(typename std::enable_if<(N == 1), size_t>::type dim0) : common_array{dim0} {} /* The following constructor is only available in the array struct * specialization where: dimensions==2 */ - template INLINE_IF_DEVICE + template array(typename std::enable_if<(N == 2), size_t>::type dim0, size_t dim1) : common_array{dim0, dim1} {} /* The following constructor is only available in the array struct * specialization where: dimensions==3 */ - template INLINE_IF_DEVICE + template array(typename std::enable_if<(N == 3), size_t>::type dim0, size_t dim1, size_t dim2) : common_array{dim0, dim1, dim2} {} // Conversion operators to derived classes - INLINE_IF_DEVICE operator cl::sycl::id() const { + operator cl::sycl::id() const { cl::sycl::id result; for (int i = 0; i < dimensions; ++i) { result[i] = common_array[i]; @@ -51,7 +50,7 @@ template class array { return result; } - INLINE_IF_DEVICE operator cl::sycl::range() const { + operator cl::sycl::range() const { cl::sycl::range result; for (int i = 0; i < dimensions; ++i) { result[i] = common_array[i]; @@ -59,29 +58,29 @@ template class array { return result; } - INLINE_IF_DEVICE size_t get(int dimension) const { + size_t get(int dimension) const { check_dimension(dimension); return common_array[dimension]; } - INLINE_IF_DEVICE size_t &operator[](int dimension) { + size_t &operator[](int dimension) { check_dimension(dimension); return common_array[dimension]; } - INLINE_IF_DEVICE size_t operator[](int dimension) const { + size_t operator[](int dimension) const { check_dimension(dimension); return common_array[dimension]; } - INLINE_IF_DEVICE array(const array &rhs) = default; - INLINE_IF_DEVICE array(array &&rhs) = default; - INLINE_IF_DEVICE array &operator=(const array &rhs) = default; - INLINE_IF_DEVICE array &operator=(array &&rhs) = default; + array(const array &rhs) = default; + array(array &&rhs) = default; + array &operator=(const array &rhs) = default; + array &operator=(array &&rhs) = default; // Returns true iff all elements in 'this' are equal to // the corresponding elements in 'rhs'. - INLINE_IF_DEVICE bool operator==(const array &rhs) const { + bool operator==(const array &rhs) const { for (int i = 0; i < dimensions; ++i) { if (this->common_array[i] != rhs.common_array[i]) { return false; @@ -92,7 +91,7 @@ template class array { // Returns true iff there is at least one element in 'this' // which is not equal to the corresponding element in 'rhs'. - INLINE_IF_DEVICE bool operator!=(const array &rhs) const { + bool operator!=(const array &rhs) const { for (int i = 0; i < dimensions; ++i) { if (this->common_array[i] != rhs.common_array[i]) { return true; diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp index f763c06fe4c96..73064e988538f 100644 --- a/sycl/include/CL/sycl/detail/buffer_impl.hpp +++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp @@ -1,9 +1,8 @@ //==---------- buffer_impl.hpp --- SYCL buffer ----------------*- C++-*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -29,6 +28,7 @@ namespace cl { namespace sycl { using QueueImplPtr = std::shared_ptr; using EventImplPtr = std::shared_ptr; +using ContextImplPtr = std::shared_ptr; // Forward declarations template @@ -120,8 +120,9 @@ class buffer_impl { "allowed"); CHECK_OCL_CODE(clGetMemObjectInfo(MemObject, CL_MEM_CONTEXT, - sizeof(OpenCLContext), &OpenCLContext, nullptr)); - if (SyclContext.get() != OpenCLContext) + sizeof(OpenCLContext), &OpenCLContext, + nullptr)); + if (detail::getSyclObjImpl(SyclContext)->getHandleRef() != OpenCLContext) throw cl::sycl::invalid_parameter_error( "Input context must be the same as the context of cl_mem"); OCLState.Mem = MemObject; @@ -268,7 +269,7 @@ void buffer_impl::fill( size_t Offset = OffsetArr[0]; size_t Size = RangeArr[0] * PatternSize; - cl::sycl::context Context = Queue->get_context(); + ContextImplPtr Context = detail::getSyclObjImpl(Queue->get_context()); OCLState.Queue = std::move(Queue); Event->setIsHostEvent(false); @@ -278,9 +279,9 @@ void buffer_impl::fill( detail::getOrWaitEvents(std::move(DepEvents), Context); cl_command_queue CommandQueue = OCLState.Queue->get(); - cl_int Error = clEnqueueFillBuffer( - CommandQueue, OCLState.Mem, Pattern, PatternSize, Offset, Size, - CLEvents.size(), CLEvents.data(), &BufEvent); + cl_int Error = clEnqueueFillBuffer(CommandQueue, OCLState.Mem, Pattern, + PatternSize, Offset, Size, CLEvents.size(), + CLEvents.data(), &BufEvent); CHECK_OCL_CODE(Error); CHECK_OCL_CODE(clReleaseCommandQueue(CommandQueue)); @@ -299,7 +300,7 @@ void buffer_impl::copy( size_t SizeTyDest = sizeof(T); const int DimDest = dimensions; - cl::sycl::context Context = Queue->get_context(); + ContextImplPtr Context = detail::getSyclObjImpl(Queue->get_context()); cl_event &BufEvent = Event->getHandleRef(); std::vector CLEvents = @@ -314,11 +315,11 @@ void buffer_impl::copy( CLEvents.data(), &BufEvent); } else { size_t SrcOrigin[3] = {SrcOffset[0] * SizeTySrc, - (1 == DimSrc) ? 0 : SrcOffset[1], - (3 == DimSrc) ? SrcOffset[2] : 0}; + (1 == DimSrc) ? 0 : SrcOffset[1], + (3 == DimSrc) ? SrcOffset[2] : 0}; size_t DstOrigin[3] = {DestOffset[0] * SizeTyDest, - (1 == DimDest) ? 0 : DestOffset[1], - (3 == DimDest) ? DestOffset[2] : 0}; + (1 == DimDest) ? 0 : DestOffset[1], + (3 == DimDest) ? DestOffset[2] : 0}; size_t Region[3] = {SrcRange[0] * SizeTySrc, (1 == DimSrc) ? 1 : SrcRange[1], (3 == DimSrc) ? SrcRange[2] : 1}; @@ -345,9 +346,9 @@ void buffer_impl::moveMemoryTo( QueueImplPtr Queue, std::vector DepEvents, EventImplPtr Event) { - cl::sycl::context Context = Queue->get_context(); + ContextImplPtr Context = detail::getSyclObjImpl(Queue->get_context()); - if (OpenCLInterop && (Context.get() != OpenCLContext)) + if (OpenCLInterop && (Context->getHandleRef() != OpenCLContext)) throw cl::sycl::runtime_error( "Interoperability buffer could not be used in a context other than the " "context associated with the OpenCL memory object."); @@ -364,7 +365,7 @@ void buffer_impl::moveMemoryTo( return; } - assert(OCLState.Queue->get_context() != Context || + assert(OCLState.Queue->get_context() != Queue->get_context() || OCLState.Queue->get_device() != Queue->get_device() && "Attempt to move to the same env"); @@ -394,8 +395,9 @@ void buffer_impl::moveMemoryTo( if (OCLState.Queue->is_host() && !Queue->is_host()) { const size_t ByteSize = get_size(); cl_int Error; - cl_mem Mem = clCreateBuffer(Context.get(), CL_MEM_READ_WRITE, ByteSize, - /*host_ptr=*/nullptr, &Error); + cl_mem Mem = + clCreateBuffer(Context->getHandleRef(), CL_MEM_READ_WRITE, ByteSize, + /*host_ptr=*/nullptr, &Error); CHECK_OCL_CODE(Error); OCLState.Queue = std::move(Queue); @@ -461,9 +463,9 @@ void buffer_impl::allocate( detail::waitEvents(DepEvents); - cl::sycl::context Context = Queue->get_context(); + ContextImplPtr Context = detail::getSyclObjImpl(Queue->get_context()); - if (OpenCLInterop && (Context.get() != OpenCLContext)) + if (OpenCLInterop && (Context->getHandleRef() != OpenCLContext)) throw cl::sycl::runtime_error( "Interoperability buffer could not be used in a context other than the " "context associated with the OpenCL memory object."); @@ -479,8 +481,9 @@ void buffer_impl::allocate( size_t ByteSize = get_size(); cl_int Error; - cl_mem Mem = clCreateBuffer(Context.get(), convertSycl2OCLMode(mode), - ByteSize, nullptr, &Error); + cl_mem Mem = + clCreateBuffer(Context->getHandleRef(), convertSycl2OCLMode(mode), + ByteSize, nullptr, &Error); CHECK_OCL_CODE(Error); cl_event &WriteBufEvent = Event->getHandleRef(); diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp index 7241f7fe12118..412485b9ad6d1 100644 --- a/sycl/include/CL/sycl/detail/common.hpp +++ b/sycl/include/CL/sycl/detail/common.hpp @@ -1,9 +1,8 @@ //==---------- common.hpp ----- Common declarations ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -74,16 +73,6 @@ const char *stringifyErrorCode(cl_int error); #define ALWAYS_INLINE #endif -// TODO this macro is introduced to workaround SPIRV translator problem with -// dropping linkonce_odr attribute leading to duplicated symbol errors in -// the bitcode linker for functions defined in the headers. Remove once fixed. -#ifdef __SYCL_DEVICE_ONLY__ -#define INLINE_IF_DEVICE ALWAYS_INLINE -#else -#define INLINE_IF_DEVICE -#endif // __SYCL_DEVICE_ONLY__ - - namespace cl { namespace sycl { namespace detail { diff --git a/sycl/include/CL/sycl/detail/common_info.hpp b/sycl/include/CL/sycl/detail/common_info.hpp index 636dd9a428196..814707f0bb4ea 100644 --- a/sycl/include/CL/sycl/detail/common_info.hpp +++ b/sycl/include/CL/sycl/detail/common_info.hpp @@ -1,9 +1,8 @@ //==------- common_info.hpp ----- Common SYCL info methods------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/context_host.hpp b/sycl/include/CL/sycl/detail/context_host.hpp deleted file mode 100644 index 0731a22c54f8d..0000000000000 --- a/sycl/include/CL/sycl/detail/context_host.hpp +++ /dev/null @@ -1,46 +0,0 @@ -//==------------- context_host.hpp - SYCL host context ---------------------==// -// -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. -// -//===----------------------------------------------------------------------===// - -#pragma once -#include -#include -#include -#include -#include -#include -// 4.6.2 Context class - -namespace cl { -namespace sycl { -namespace detail { -class context_host : public context_impl { -public: - context_host(const device &rhs, async_handler asyncHandler) - : context_impl(asyncHandler), dev(rhs) {} - - cl_context get() const override { - throw invalid_object_error("This instance of context is a host instance"); - } - - bool is_host() const override { return true; } - - platform get_platform() const override { return platform(); } - - vector_class get_devices() const override { - return vector_class(1, dev); - } - - template - typename info::param_traits::return_type get_info() const; -private: - device dev; -}; -} // namespace detail -} // namespace sycl -} // namespace cl diff --git a/sycl/include/CL/sycl/detail/context_impl.hpp b/sycl/include/CL/sycl/detail/context_impl.hpp index 087d09779f7a8..334091b759e23 100644 --- a/sycl/include/CL/sycl/detail/context_impl.hpp +++ b/sycl/include/CL/sycl/detail/context_impl.hpp @@ -1,9 +1,8 @@ -//==---------------- context.hpp - SYCL context ----------------------------==// +//==---------------- context_impl.hpp - SYCL context -----------*- C++-*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -13,71 +12,51 @@ #include #include #include + +#include // 4.6.2 Context class namespace cl { namespace sycl { // Forward declaration -class platform; class device; namespace detail { -template struct get_context_info_cl { - using RetType = - typename info::param_traits::return_type; - - static RetType _(cl_context ctx) { - RetType Result = 0; - // TODO catch an exception and put it to list of asynchronous exceptions - CHECK_OCL_CODE(clGetContextInfo(ctx, cl_context_info(param), sizeof(Result), - &Result, nullptr)); - return Result; - } -}; - class context_impl { public: - context_impl(async_handler asyncHandler) : m_AsyncHandler(asyncHandler) {} + context_impl(const device &Device, async_handler AsyncHandler); - template - inline typename info::param_traits::return_type - get_info() const; + context_impl(const vector_class Devices, + async_handler AsyncHandler); + + context_impl(cl_context ClContext, async_handler AsyncHandler); + + ~context_impl(); - const async_handler& get_async_handler() const { return m_AsyncHandler; } + cl_context get() const; - virtual cl_context get() const = 0; + bool is_host() const; - virtual bool is_host() const = 0; + platform get_platform() const; - virtual platform get_platform() const = 0; + vector_class get_devices() const; - virtual vector_class get_devices() const = 0; + const async_handler &get_async_handler() const; + + template + typename info::param_traits::return_type + get_info() const; - virtual ~context_impl() = default; + // Warning. Returned reference will be invalid if context_impl was destroyed. + cl_context &getHandleRef(); private: async_handler m_AsyncHandler; + vector_class m_Devices; + cl_context m_ClContext; + platform m_Platform; + bool m_OpenCLInterop; + bool m_HostContext; }; -template <> -inline typename info::param_traits::return_type -context_impl::get_info() const { - if (is_host()) { - return 0; - } - return get_context_info_cl::_(this->get()); -} -template <> -inline typename info::param_traits::return_type -context_impl::get_info() const { - return get_platform(); -} -template <> -inline typename info::param_traits::return_type -context_impl::get_info() const { - return get_devices(); -} } // namespace detail } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/context_info.hpp b/sycl/include/CL/sycl/detail/context_info.hpp new file mode 100644 index 0000000000000..e7ae8cafcb3cc --- /dev/null +++ b/sycl/include/CL/sycl/detail/context_info.hpp @@ -0,0 +1,33 @@ +//==---------------- context_info.hpp - SYCL context -----------*- C++ -*---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// ===--------------------------------------------------------------------=== // + +#pragma once + +#include +#include + +namespace cl { +namespace sycl { +namespace detail { + +template struct get_context_info_cl { + using RetType = + typename info::param_traits::return_type; + + static RetType _(cl_context ctx) { + RetType Result = 0; + // TODO catch an exception and put it to list of asynchronous exceptions + CHECK_OCL_CODE(clGetContextInfo(ctx, cl_context_info(param), sizeof(Result), + &Result, nullptr)); + return Result; + } +}; + +} // namespace detail +} // namespace sycl +} // namespace cl diff --git a/sycl/include/CL/sycl/detail/context_opencl.hpp b/sycl/include/CL/sycl/detail/context_opencl.hpp deleted file mode 100644 index e5755902e29b9..0000000000000 --- a/sycl/include/CL/sycl/detail/context_opencl.hpp +++ /dev/null @@ -1,90 +0,0 @@ -//==------------ context_opencl.hpp - SYCL OpenCL context ------------------==// -// -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. -// -//===----------------------------------------------------------------------===// - -#pragma once -#include -#include -#include -#include -#include -#include - -// 4.6.2 Context class - -namespace cl { -namespace sycl { -// Forward declaration -class platform; -namespace detail { -class context_opencl : public context_impl { -public: - context_opencl(const vector_class devices, - async_handler asyncHandler) - : context_impl(asyncHandler) { - dev_list = devices; - plt = dev_list[0].get_platform(); - vector_class dev_ids; - for (const auto &d : dev_list) - dev_ids.push_back(d.get()); - cl_int error; - id = clCreateContext(0, dev_ids.size(), dev_ids.data(), 0, 0, &error); - // TODO catch an exception and put it to list of asynchronous exceptions - CHECK_OCL_CODE(error); - } - - context_opencl(cl_context clContext, async_handler asyncHandler) - : context_impl(asyncHandler) { - id = clContext; - vector_class dev_ids; - size_t devicesBuffer = 0; - // TODO catch an exception and put it to list of asynchronous exceptions - CHECK_OCL_CODE( - clGetContextInfo(id, CL_CONTEXT_DEVICES, 0, nullptr, &devicesBuffer)); - dev_ids.resize(devicesBuffer / sizeof(cl_device_id)); - // TODO catch an exception and put it to list of asynchronous exceptions - CHECK_OCL_CODE(clGetContextInfo(id, CL_CONTEXT_DEVICES, devicesBuffer, - &dev_ids[0], nullptr)); - - for (auto dev : dev_ids) { - dev_list.emplace_back(dev); - } - // TODO What if dev_list if empty? dev_list[0].get_platform() - plt = platform(dev_list[0].get_platform()); - // TODO catch an exception and put it to list of asynchronous exceptions - CHECK_OCL_CODE(clRetainContext(id)); - } - - cl_context get() const override { - // TODO catch an exception and put it to list of asynchronous exceptions - CHECK_OCL_CODE(clRetainContext(id)); - return id; - } - - bool is_host() const override { return false; } - - platform get_platform() const override { return plt; } - - vector_class get_devices() const override { return dev_list; } - - ~context_opencl() { - // TODO replace CHECK_OCL_CODE_NO_EXC to CHECK_OCL_CODE and - // TODO catch an exception and put it to list of asynchronous exceptions - CHECK_OCL_CODE_NO_EXC(clReleaseContext(id)); - } - // TODO: implement param traits - // template - // typename param_traits::type get_info() const; -private: - vector_class dev_list; - cl_context id; - platform plt; -}; -} // namespace detail -} // namespace sycl -} // namespace cl diff --git a/sycl/include/CL/sycl/detail/device_host.hpp b/sycl/include/CL/sycl/detail/device_host.hpp index 6ce172ac435e8..8f17c4c880b6e 100644 --- a/sycl/include/CL/sycl/detail/device_host.hpp +++ b/sycl/include/CL/sycl/detail/device_host.hpp @@ -1,9 +1,8 @@ //==--------------- device_host.hpp - SYCL host device --------------------== // // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/device_impl.hpp b/sycl/include/CL/sycl/detail/device_impl.hpp index 671a9ed187d79..1238ce058d976 100644 --- a/sycl/include/CL/sycl/detail/device_impl.hpp +++ b/sycl/include/CL/sycl/detail/device_impl.hpp @@ -1,9 +1,8 @@ //==----------------- device_impl.hpp - SYCL device ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/device_info.hpp b/sycl/include/CL/sycl/detail/device_info.hpp index dc581c31160fc..03c9b8a7face6 100644 --- a/sycl/include/CL/sycl/detail/device_info.hpp +++ b/sycl/include/CL/sycl/detail/device_info.hpp @@ -1,9 +1,8 @@ //==-------- device_info.hpp - SYCL device info methods --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/device_opencl.hpp b/sycl/include/CL/sycl/detail/device_opencl.hpp index a61206b8b760d..31808fb3d3a9b 100644 --- a/sycl/include/CL/sycl/detail/device_opencl.hpp +++ b/sycl/include/CL/sycl/detail/device_opencl.hpp @@ -1,9 +1,8 @@ //==------------ device_opencl.hpp - SYCL OpenCL device --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/event_impl.hpp b/sycl/include/CL/sycl/detail/event_impl.hpp index 833cc335bb326..0e12972753674 100644 --- a/sycl/include/CL/sycl/detail/event_impl.hpp +++ b/sycl/include/CL/sycl/detail/event_impl.hpp @@ -1,9 +1,8 @@ //==---------------- event_impl.hpp - SYCL event ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -33,6 +32,8 @@ class event_impl { // Self is needed in order to pass shared_ptr to Scheduler. void wait(std::shared_ptr Self) const; + void wait_and_throw(std::shared_ptr Self); + template typename info::param_traits::return_type get_profiling_info() const; @@ -44,6 +45,7 @@ class event_impl { void waitInternal() const; + // Warning. Returned reference will be invalid if event_impl was destroyed. cl_event &getHandleRef(); void setIsHostEvent(bool Value); diff --git a/sycl/include/CL/sycl/detail/event_info.hpp b/sycl/include/CL/sycl/detail/event_info.hpp index 56725642d80ed..5a5740cfba9b7 100644 --- a/sycl/include/CL/sycl/detail/event_info.hpp +++ b/sycl/include/CL/sycl/detail/event_info.hpp @@ -1,14 +1,14 @@ //==---------------- event_info.hpp - SYCL event ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #pragma once +#include #include namespace cl { diff --git a/sycl/include/CL/sycl/detail/helpers.hpp b/sycl/include/CL/sycl/detail/helpers.hpp index f8e95977eee45..03a151a861171 100644 --- a/sycl/include/CL/sycl/detail/helpers.hpp +++ b/sycl/include/CL/sycl/detail/helpers.hpp @@ -1,15 +1,16 @@ //==---------------- helpers.hpp - SYCL helpers ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #pragma once #include + +#include #include #include #include @@ -24,12 +25,12 @@ template class range; template class id; template class nd_item; namespace detail { - +class context_impl; // The function returns list of events that can be passed to OpenCL API as // dependency list and waits for others. std::vector getOrWaitEvents(std::vector DepEvents, - cl::sycl::context Context); + std::shared_ptr Context); void waitEvents(std::vector DepEvents); diff --git a/sycl/include/CL/sycl/detail/image_impl.hpp b/sycl/include/CL/sycl/detail/image_impl.hpp index a91eaf78ce53e..6ba8241731b65 100644 --- a/sycl/include/CL/sycl/detail/image_impl.hpp +++ b/sycl/include/CL/sycl/detail/image_impl.hpp @@ -1,9 +1,8 @@ //==------------ image_impl.hpp --------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index 25862ab7dea30..d01a9a9ab6773 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -1,9 +1,8 @@ //==----------------------- kernel_desc.hpp --------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===//// diff --git a/sycl/include/CL/sycl/detail/kernel_impl.hpp b/sycl/include/CL/sycl/detail/kernel_impl.hpp index ea1a9ca0a91e1..db5a6af110cd8 100644 --- a/sycl/include/CL/sycl/detail/kernel_impl.hpp +++ b/sycl/include/CL/sycl/detail/kernel_impl.hpp @@ -1,9 +1,8 @@ //==------- kernel_impl.hpp --- SYCL kernel implementation -----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/kernel_info.hpp b/sycl/include/CL/sycl/detail/kernel_info.hpp index cbae1fb42edb8..2d90ee2338c0b 100644 --- a/sycl/include/CL/sycl/detail/kernel_info.hpp +++ b/sycl/include/CL/sycl/detail/kernel_info.hpp @@ -1,9 +1,8 @@ //==-------- kernel_info.hpp - SYCL kernel info methods --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/platform_host.hpp b/sycl/include/CL/sycl/detail/platform_host.hpp index ceed82b0a5bd8..efb612576e02e 100644 --- a/sycl/include/CL/sycl/detail/platform_host.hpp +++ b/sycl/include/CL/sycl/detail/platform_host.hpp @@ -1,9 +1,8 @@ //==------------ platform_host.hpp - SYCL host platform --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/platform_impl.hpp b/sycl/include/CL/sycl/detail/platform_impl.hpp index 2da2488a1f385..8d1831bea6323 100644 --- a/sycl/include/CL/sycl/detail/platform_impl.hpp +++ b/sycl/include/CL/sycl/detail/platform_impl.hpp @@ -1,9 +1,8 @@ //==-------------- platform_impl.hpp - SYCL platform -----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/platform_info.hpp b/sycl/include/CL/sycl/detail/platform_info.hpp index 04d97f712c948..4e2336015df63 100644 --- a/sycl/include/CL/sycl/detail/platform_info.hpp +++ b/sycl/include/CL/sycl/detail/platform_info.hpp @@ -1,9 +1,8 @@ //==------ platform_info.hpp - SYCL platform info methods ------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/platform_opencl.hpp b/sycl/include/CL/sycl/detail/platform_opencl.hpp index 47aaf1d459d5e..a88fd4552d0ad 100644 --- a/sycl/include/CL/sycl/detail/platform_opencl.hpp +++ b/sycl/include/CL/sycl/detail/platform_opencl.hpp @@ -1,9 +1,8 @@ //==-------- platform_opencl.hpp - SYCL OpenCL platform --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/program_impl.hpp b/sycl/include/CL/sycl/detail/program_impl.hpp index d10f1c636beea..eec71118ec610 100644 --- a/sycl/include/CL/sycl/detail/program_impl.hpp +++ b/sycl/include/CL/sycl/detail/program_impl.hpp @@ -1,19 +1,17 @@ //==----- program_impl.hpp --- SYCL program implementation -----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// - #pragma once -#include +#include #include #include #include -#include +#include #include #include #include @@ -66,10 +64,10 @@ class program_impl { ClPrograms.push_back(Prg->ClProgram); } cl_int Err; - ClProgram = - clLinkProgram(Context.get(), ClDevices.size(), ClDevices.data(), - LinkOptions.c_str(), ProgramList.size(), - ClPrograms.data(), nullptr, nullptr, &Err); + ClProgram = clLinkProgram(detail::getSyclObjImpl(Context)->getHandleRef(), + ClDevices.size(), ClDevices.data(), + LinkOptions.c_str(), ProgramList.size(), + ClPrograms.data(), nullptr, nullptr, &Err); CHECK_OCL_CODE_THROW(Err, compile_program_error); } } @@ -179,9 +177,10 @@ class program_impl { if (!is_host()) { vector_class ClDevices(get_cl_devices()); cl_int Err; - ClProgram = clLinkProgram(Context.get(), ClDevices.size(), - ClDevices.data(), LinkOptions.c_str(), 1, - &ClProgram, nullptr, nullptr, &Err); + ClProgram = + clLinkProgram(detail::getSyclObjImpl(Context)->getHandleRef(), + ClDevices.size(), ClDevices.data(), LinkOptions.c_str(), + 1, &ClProgram, nullptr, nullptr, &Err); CHECK_OCL_CODE_THROW(Err, compile_program_error); LinkOptions = LinkOptions; } @@ -274,7 +273,7 @@ class program_impl { private: void create_cl_program_with_il() { assert(!ClProgram && "This program already has an encapsulated cl_program"); - ClProgram = ProgramManager::getInstance().getBuiltOpenCLProgram(Context); + ClProgram = ProgramManager::getInstance().createOpenCLProgram(Context); } void create_cl_program_with_source(const string_class &Source) { @@ -282,7 +281,8 @@ class program_impl { cl_int Err; const char *Src = Source.c_str(); size_t Size = Source.size(); - ClProgram = clCreateProgramWithSource(Context.get(), 1, &Src, &Size, &Err); + ClProgram = clCreateProgramWithSource( + detail::getSyclObjImpl(Context)->getHandleRef(), 1, &Src, &Size, &Err); CHECK_OCL_CODE(Err); } diff --git a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp index 56376d7848f88..3eaa39f15188c 100644 --- a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp +++ b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp @@ -1,9 +1,8 @@ //==------ program_manager.hpp --- SYCL program manager---------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -51,6 +50,7 @@ namespace detail { class ProgramManager { public: static ProgramManager &getInstance(); + cl_program createOpenCLProgram(const context &Context); cl_program getBuiltOpenCLProgram(const context &Context); cl_kernel getOrCreateKernel(const context &Context, const char *KernelName); cl_program getClProgramFromClKernel(cl_kernel ClKernel); diff --git a/sycl/include/CL/sycl/detail/queue_impl.hpp b/sycl/include/CL/sycl/detail/queue_impl.hpp old mode 100644 new mode 100755 index 1c2782bbd53a1..e71c3d8755d0c --- a/sycl/include/CL/sycl/detail/queue_impl.hpp +++ b/sycl/include/CL/sycl/detail/queue_impl.hpp @@ -1,9 +1,8 @@ //==------------------ queue_impl.hpp - SYCL queue -------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -20,34 +19,22 @@ namespace cl { namespace sycl { namespace detail { +// Set max number of queues supported by FPGA RT. +const size_t MaxNumQueues = 256; + class queue_impl { public: queue_impl(const device &SyclDevice, async_handler AsyncHandler, const property_list &PropList) - : m_Device(SyclDevice), m_Context(m_Device), m_AsyncHandler(AsyncHandler), + : queue_impl(SyclDevice, context(SyclDevice), AsyncHandler, PropList){}; + + queue_impl(const device &SyclDevice, const context &Context, + async_handler AsyncHandler, const property_list &PropList) + : m_Device(SyclDevice), m_Context(Context), m_AsyncHandler(AsyncHandler), m_PropList(PropList), m_HostQueue(m_Device.is_host()) { m_OpenCLInterop = !m_HostQueue; if (!m_HostQueue) { - cl_command_queue_properties CreationFlags = - CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE; - - if (m_PropList.has_property()) { - CreationFlags |= CL_QUEUE_PROFILING_ENABLE; - } - - cl_int Error = CL_SUCCESS; -#ifdef CL_VERSION_2_0 - vector_class CreationFlagProperties = { - CL_QUEUE_PROPERTIES, CreationFlags, 0}; - m_CommandQueue = clCreateCommandQueueWithProperties( - m_Context.get(), m_Device.get(), CreationFlagProperties.data(), - &Error); -#else - m_CommandQueue = clCreateCommandQueue(m_Context.get(), m_Device.get(), - CreationFlags, &Error); -#endif - CHECK_OCL_CODE(Error); - // TODO catch an exception and put it to list of asynchronous exceptions + m_CommandQueue = createQueue(); } } @@ -89,8 +76,9 @@ class queue_impl { template typename info::param_traits::return_type get_info() const; - template event submit(T cgf, std::shared_ptr self, - std::shared_ptr second_queue) { + template + event submit(T cgf, std::shared_ptr self, + std::shared_ptr second_queue) { event Event; try { Event = submit_impl(cgf, self); @@ -105,7 +93,7 @@ class queue_impl { event Event; try { Event = submit_impl(cgf, self); - } catch(...) { + } catch (...) { m_Exceptions.push_back(std::current_exception()); } return Event; @@ -132,7 +120,62 @@ class queue_impl { m_Exceptions.clear(); } - cl_command_queue &getHandleRef() { return m_CommandQueue; } + cl_command_queue createQueue() const { + cl_command_queue_properties CreationFlags = 0; + + // FPGA RT can't handle out of order queue - create in order queue instead + if (!m_Device.is_accelerator()) { + CreationFlags = CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE; + } + + if (m_PropList.has_property()) { + CreationFlags |= CL_QUEUE_PROFILING_ENABLE; + } + + cl_int Error = CL_SUCCESS; + cl_command_queue Queue; + cl_context ClContext = detail::getSyclObjImpl(m_Context)->getHandleRef(); +#ifdef CL_VERSION_2_0 + cl_queue_properties CreationFlagProperties[] = { + CL_QUEUE_PROPERTIES, CreationFlags, 0}; + Queue = clCreateCommandQueueWithProperties( + ClContext, m_Device.get(), CreationFlagProperties, + &Error); +#else + Queue = clCreateCommandQueue(ClContext, m_Device.get(), + CreationFlags, &Error); +#endif + CHECK_OCL_CODE(Error); + // TODO catch an exception and put it to list of asynchronous exceptions + + return Queue; + } + + // Warning. Returned reference will be invalid if queue_impl was destroyed. + cl_command_queue &getHandleRef() { + if (!m_Device.is_accelerator()) { + return m_CommandQueue; + } + + // To achive parallelism for FPGA with in order execution model with + // possibility of two kernels to share data with each other we shall + // create a queue for every kernel enqueued. + if (m_Queues.empty()) { + m_Queues.push_back(m_CommandQueue); + return m_CommandQueue; + } else if (m_Queues.size() < MaxNumQueues) { + m_Queues.push_back(createQueue()); + return m_Queues.back(); + } + + // If the limit of OpenCL queues is going to be exceeded - take the earliest + // used queue, wait until it finished and then reuse it. + m_QueueNumber %= MaxNumQueues; + size_t FreeQueueNum = m_QueueNumber++; + + CHECK_OCL_CODE(clFinish(m_Queues[FreeQueueNum])); + return m_Queues[FreeQueueNum]; + } template bool has_property() const { return m_PropList.has_property(); @@ -161,6 +204,12 @@ class queue_impl { property_list m_PropList; cl_command_queue m_CommandQueue = nullptr; + + // List of OpenCL queues created for FPGA device from a single SYCL queue. + vector_class m_Queues; + // Iterator through m_Queues. + size_t m_QueueNumber = 0; + bool m_OpenCLInterop = false; bool m_HostQueue = false; }; diff --git a/sycl/include/CL/sycl/detail/scheduler/commands.cpp b/sycl/include/CL/sycl/detail/scheduler/commands.cpp index 68cacd0e0f079..7d1ae84af1277 100644 --- a/sycl/include/CL/sycl/detail/scheduler/commands.cpp +++ b/sycl/include/CL/sycl/detail/scheduler/commands.cpp @@ -1,9 +1,8 @@ //==----------- commands.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -40,10 +39,10 @@ void ExecuteKernelCommand< runOnHost(); return; } - + context Context = m_Queue->get_context(); if (!m_ClKernel) { m_ClKernel = detail::ProgramManager::getInstance().getOrCreateKernel( - m_Queue->get_context(), m_KernelName.c_str()); + Context, m_KernelName.c_str()); } if (m_KernelArgs != nullptr) { @@ -99,8 +98,8 @@ void ExecuteKernelCommand< } } - std::vector CLEvents = - detail::getOrWaitEvents(std::move(DepEvents), m_Queue->get_context()); + std::vector CLEvents = detail::getOrWaitEvents( + std::move(DepEvents), detail::getSyclObjImpl(Context)); cl_event &CLEvent = Event->getHandleRef(); CLEvent = runEnqueueNDRangeKernel(m_Queue->getHandleRef(), m_ClKernel, std::move(CLEvents)); diff --git a/sycl/include/CL/sycl/detail/scheduler/commands.h b/sycl/include/CL/sycl/detail/scheduler/commands.h index ab038e6e37dd3..78870c34f5bdd 100644 --- a/sycl/include/CL/sycl/detail/scheduler/commands.h +++ b/sycl/include/CL/sycl/detail/scheduler/commands.h @@ -1,9 +1,8 @@ //==----------- commands.h -------------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/scheduler/printers.cpp b/sycl/include/CL/sycl/detail/scheduler/printers.cpp index 660f470a83654..7d09ea6144912 100644 --- a/sycl/include/CL/sycl/detail/scheduler/printers.cpp +++ b/sycl/include/CL/sycl/detail/scheduler/printers.cpp @@ -1,9 +1,8 @@ //==----------- printers.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/scheduler/requirements.h b/sycl/include/CL/sycl/detail/scheduler/requirements.h index 7e36271b82203..47968905018a5 100644 --- a/sycl/include/CL/sycl/detail/scheduler/requirements.h +++ b/sycl/include/CL/sycl/detail/scheduler/requirements.h @@ -1,9 +1,8 @@ //==----------- requirements.h ---------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp b/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp index dd12594823d1a..0509196da4bd0 100644 --- a/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp +++ b/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp @@ -1,9 +1,8 @@ //==----------- scheduler.cpp ----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/detail/scheduler/scheduler.h b/sycl/include/CL/sycl/detail/scheduler/scheduler.h index 66b61c3c868c7..cabbacff4bb11 100644 --- a/sycl/include/CL/sycl/detail/scheduler/scheduler.h +++ b/sycl/include/CL/sycl/detail/scheduler/scheduler.h @@ -1,9 +1,8 @@ //==----------- scheduler.h ------------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -22,6 +21,7 @@ #include #include #include +#include #include namespace cl { @@ -157,6 +157,10 @@ class Scheduler { // Waits for the event passed. void waitForEvent(EventImplPtr Event); + // Calls asynchronous handler for the passed event Event + // and for those other events that Event depends on. + void throwForEventRecursive(EventImplPtr Event); + // Adds new node to graph, creating an Alloca and MemMove commands if // needed. cl::sycl::event addNode(Node NewNode); @@ -190,14 +194,17 @@ class Scheduler { // void parallelReadOpt(); - static Scheduler &getInstance() { - static Scheduler instance; - return instance; - } + static Scheduler &getInstance(); enum DumpOptions { Text = 0, WholeGraph = 1, RunGraph = 2 }; bool getDumpFlagValue(DumpOptions DumpOption); + // Recursively walks through the dependencies and initializes + // the given EventsSet with the events that the Event + // waits for. The unordered_set is used to collect unuque events, + // and the unordered_set is convenient as it does not need operator<(). + void getDepEventsRecursive(std::unordered_set &EventsSet, + EventImplPtr Event); protected: // TODO: Add releasing of OpenCL buffers. @@ -211,7 +218,6 @@ class Scheduler { // Recursively generates dot records for the command passed and all that the // command depends on. void printGraphForCommand(CommandPtr Cmd, std::ostream &Stream) const; - private: Scheduler(); ~Scheduler(); @@ -226,6 +232,15 @@ class Scheduler { Scheduler(Scheduler const &) = delete; Scheduler &operator=(Scheduler const &) = delete; + + // Returns the pointer to the command associated with the given event, + // or nullptr if none is found. + CommandPtr getCmdForEvent(EventImplPtr Event); + + // Basically it is the helper method for throwForEventRecursive() now. + // It calls async handler for the command Cmd and those other + // commands that Cmd depends on. + void throwForCmdRecursive(std::shared_ptr Cmd); }; } // namespace simple_scheduler diff --git a/sycl/include/CL/sycl/device.hpp b/sycl/include/CL/sycl/device.hpp index 9b660d357f93b..753cd56591bc4 100644 --- a/sycl/include/CL/sycl/device.hpp +++ b/sycl/include/CL/sycl/device.hpp @@ -1,9 +1,8 @@ //==------------------- device.hpp - SYCL device ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/device_event.hpp b/sycl/include/CL/sycl/device_event.hpp index 9a057d39e999d..66217d89005dd 100644 --- a/sycl/include/CL/sycl/device_event.hpp +++ b/sycl/include/CL/sycl/device_event.hpp @@ -1,9 +1,8 @@ //==---------- device_event.hpp --- SYCL device event ---------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/device_selector.hpp b/sycl/include/CL/sycl/device_selector.hpp index 1e70a55d44b4c..c05b7e2ac496b 100644 --- a/sycl/include/CL/sycl/device_selector.hpp +++ b/sycl/include/CL/sycl/device_selector.hpp @@ -1,9 +1,8 @@ //==------ device_selector.hpp - SYCL device selector ---------*- C++ --*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/event.hpp b/sycl/include/CL/sycl/event.hpp index 3ac11194bbf5d..fa1ecb8f87dd2 100644 --- a/sycl/include/CL/sycl/event.hpp +++ b/sycl/include/CL/sycl/event.hpp @@ -1,9 +1,8 @@ //==---------------- event.hpp --- SYCL event ------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -41,15 +40,15 @@ class event { bool is_host() const; - void wait() const; + vector_class get_wait_list(); - // vector_class get_wait_list(); + void wait(); - // static void wait(const vector_class &eventList); + static void wait(const vector_class &EventList); - // void wait_and_throw(); + void wait_and_throw(); - // static void wait_and_throw(const vector_class &eventList); + static void wait_and_throw(const vector_class &EventList); template typename info::param_traits::return_type get_info() const; diff --git a/sycl/include/CL/sycl/exception.hpp b/sycl/include/CL/sycl/exception.hpp index 67dd1d3242e34..cce7327722576 100644 --- a/sycl/include/CL/sycl/exception.hpp +++ b/sycl/include/CL/sycl/exception.hpp @@ -1,9 +1,8 @@ //==---------------- exception.hpp - SYCL exception ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/group.hpp b/sycl/include/CL/sycl/group.hpp index 969efaa104b58..21a0bde0fc15c 100644 --- a/sycl/include/CL/sycl/group.hpp +++ b/sycl/include/CL/sycl/group.hpp @@ -1,9 +1,8 @@ //==-------------- group.hpp --- SYCL work group ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index cfd9092796198..d58de20b0c426 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -1,17 +1,18 @@ //==-------- handler.hpp --- SYCL command group handler --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #pragma once +#include #include #include #include +#include #include #include #include @@ -20,21 +21,10 @@ #include #include -#include - #include #include #include -#ifdef __SYCL_DEVICE_ONLY__ -size_t get_global_size(uint dimindx); -size_t get_local_size(uint dimindx); -size_t get_global_id(uint dimindx); -size_t get_local_id(uint dimindx); -size_t get_global_offset(uint dimindx); -size_t get_group_id(uint dimindx); -#endif - template class buffer; namespace detail { + +#ifdef __SYCL_DEVICE_ONLY__ + +#define DEFINE_INIT_SIZES(POSTFIX) \ + \ + template struct InitSizesST##POSTFIX; \ + \ + template struct InitSizesST##POSTFIX<1, DstT> { \ + static void initSize(DstT &Dst) { \ + Dst[0] = cl::__spirv::get##POSTFIX<0>(); \ + } \ + }; \ + \ + template struct InitSizesST##POSTFIX<2, DstT> { \ + static void initSize(DstT &Dst) { \ + Dst[1] = cl::__spirv::get##POSTFIX<1>(); \ + InitSizesST##POSTFIX<1, DstT>::initSize(Dst); \ + } \ + }; \ + \ + template struct InitSizesST##POSTFIX<3, DstT> { \ + static void initSize(DstT &Dst) { \ + Dst[2] = cl::__spirv::get##POSTFIX<2>(); \ + InitSizesST##POSTFIX<2, DstT>::initSize(Dst); \ + } \ + }; \ + \ + template static void init##POSTFIX(DstT &Dst) { \ + InitSizesST##POSTFIX::initSize(Dst); \ + } + +DEFINE_INIT_SIZES(GlobalSize); +DEFINE_INIT_SIZES(GlobalInvocationId) +DEFINE_INIT_SIZES(WorkgroupSize) +DEFINE_INIT_SIZES(LocalInvocationId) +DEFINE_INIT_SIZES(WorkgroupId) +DEFINE_INIT_SIZES(GlobalOffset) + +#undef DEFINE_INIT_SIZES + +#endif //__SYCL_DEVICE_ONLY__ + class queue_impl; template 0 && dimensions < 4), KernelType>::type kernelFunc) { id global_id; - for (int i = 0; i < dimensions; ++i) { - global_id[i] = get_global_id(i); - } + + detail::initGlobalInvocationId(global_id); + kernelFunc(global_id); } @@ -277,10 +309,10 @@ class handler { KernelType>::type kernelFunc) { id global_id; range global_size; - for (int i = 0; i < dimensions; ++i) { - global_id[i] = get_global_id(i); - global_size[i] = get_global_size(i); - } + + detail::initGlobalInvocationId(global_id); + detail::initGlobalSize(global_size); + item Item = detail::Builder::createItem(global_size, global_id); kernelFunc(Item); @@ -299,14 +331,12 @@ class handler { id local_id; id global_offset; - for (int i = 0; i < dimensions; ++i) { - global_size[i] = get_global_size(i); - local_size[i] = get_local_size(i); - group_id[i] = get_group_id(i); - global_id[i] = get_global_id(i); - local_id[i] = get_local_id(i); - global_offset[i] = get_global_offset(i); - } + detail::initGlobalSize(global_size); + detail::initWorkgroupSize(local_size); + detail::initWorkgroupId(group_id); + detail::initGlobalInvocationId(global_id); + detail::initLocalInvocationId(local_id); + detail::initGlobalOffset(global_offset); group Group = detail::Builder::createGroup( global_size, local_size, group_id); diff --git a/sycl/include/CL/sycl/id.hpp b/sycl/include/CL/sycl/id.hpp index d1e8ea125a721..54c551c8f2d5f 100644 --- a/sycl/include/CL/sycl/id.hpp +++ b/sycl/include/CL/sycl/id.hpp @@ -1,9 +1,8 @@ //==----------- id.hpp --- SYCL iteration id -------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -19,7 +18,7 @@ template class range; template struct id : public detail::array { public: using base = detail::array; - INLINE_IF_DEVICE id() = default; + id() = default; /* The following constructor is only available in the id struct * specialization where: dimensions==1 */ @@ -149,7 +148,7 @@ template struct id : public detail::array { }; namespace detail { -template INLINE_IF_DEVICE +template size_t getOffsetForId(range Range, id Id, id Offset) { size_t offset = 0; diff --git a/sycl/include/CL/sycl/image.hpp b/sycl/include/CL/sycl/image.hpp index fdbcdd1723a05..c3105bd0e18d3 100644 --- a/sycl/include/CL/sycl/image.hpp +++ b/sycl/include/CL/sycl/image.hpp @@ -1,9 +1,8 @@ //==------------ image.hpp -------------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/info/info_desc.hpp b/sycl/include/CL/sycl/info/info_desc.hpp index 264f3f4340ee0..48e0992376442 100644 --- a/sycl/include/CL/sycl/info/info_desc.hpp +++ b/sycl/include/CL/sycl/info/info_desc.hpp @@ -1,9 +1,8 @@ //==------- info_desc.hpp - SYCL information descriptors -------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp index 7897b0749f6a4..6701a1d64e92b 100644 --- a/sycl/include/CL/sycl/intel/sub_group.hpp +++ b/sycl/include/CL/sycl/intel/sub_group.hpp @@ -1,419 +1,344 @@ //==----------- sub_group.hpp --- SYCL sub-group ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #pragma once +#include #include #include #include #include - +#include #ifdef __SYCL_DEVICE_ONLY__ -#define __NOEXCEPT noexcept -namespace cl { -namespace __spirv { -extern size_t BuiltInSubgroupLocalInvocationId() __NOEXCEPT; -extern size_t BuiltInSubgroupSize() __NOEXCEPT; -extern size_t BuiltInSubgroupMaxSize() __NOEXCEPT; -extern size_t BuiltInSubgroupId() __NOEXCEPT; -extern size_t BuiltInNumSubgroups() __NOEXCEPT; -extern size_t BuiltInNumEnqueuedSubgroups() __NOEXCEPT; -} // namespace __spirv -} // namespace cl - -// TODO: rework to use SPIRV -typedef uint uint2 __attribute__((ext_vector_type(2))); -typedef uint uint3 __attribute__((ext_vector_type(3))); -typedef uint uint4 __attribute__((ext_vector_type(4))); -typedef uint uint8 __attribute__((ext_vector_type(8))); -typedef ushort ushort2 __attribute__((ext_vector_type(2))); -typedef ushort ushort3 __attribute__((ext_vector_type(3))); -typedef ushort ushort4 __attribute__((ext_vector_type(4))); -typedef ushort ushort8 __attribute__((ext_vector_type(8))); -size_t get_sub_group_local_id(); // BuiltInSubgroupLocalInvocationId -size_t get_sub_group_size(); // BuiltInSubgroupSize -size_t get_max_sub_group_size(); // BuiltInSubgroupMaxSize -size_t get_sub_group_id(); // BuiltInSubgroupId -size_t get_num_sub_groups(); // BuiltInNumSubgroups -size_t get_enqueued_num_sub_groups(); // BuiltInNumEnqueuedSubgroups -int sub_group_any(int); -int sub_group_all(int); -int sub_group_broadcast(int x, uint sub_grou_local_id); -int sub_group_reduce_min(int x); -int sub_group_reduce_max(int x); -int sub_group_reduce_add(int x); -int sub_group_scan_exclusive_add(int x); -int sub_group_scan_exclusive_max(int x); -int sub_group_scan_exclusive_min(int x); -int sub_group_scan_inclusive_add(int x); -int sub_group_scan_inclusive_max(int x); -int sub_group_scan_inclusive_min(int x); -int intel_sub_group_shuffle(int data, uint c); -int intel_sub_group_shuffle_up(int prev, int cur, uint c); -int intel_sub_group_shuffle_down(int cur, int next, uint c); -int intel_sub_group_shuffle_xor(int data, uint c); -uint intel_sub_group_block_read(const __global uint *p); -uint2 intel_sub_group_block_read2(const __global uint *p); -uint4 intel_sub_group_block_read4(const __global uint *p); -uint8 intel_sub_group_block_read8(const __global uint *p); -void intel_sub_group_block_write(__global uint *p, uint data); -void intel_sub_group_block_write2(__global uint *p, uint2 data); -void intel_sub_group_block_write4(__global uint *p, uint4 data); -void intel_sub_group_block_write8(__global uint *p, uint8 data); - -ushort intel_sub_group_block_read_us(const __global ushort *p); -ushort2 intel_sub_group_block_read_us2(const __global ushort *p); -ushort4 intel_sub_group_block_read_us4(const __global ushort *p); -ushort8 intel_sub_group_block_read_us8(const __global ushort *p); -void intel_sub_group_block_write_us(__global ushort *p, ushort data); -void intel_sub_group_block_write_us2(__global ushort *p, ushort2 data); -void intel_sub_group_block_write_us4(__global ushort *p, ushort4 data); -void intel_sub_group_block_write_us8(__global ushort *p, ushort8 data); -void sub_group_barrier(cl::sycl::detail::cl_mem_fence_flags flags); namespace cl { namespace sycl { template class multi_ptr; namespace intel { +template -enum class Operation { exclusive_scan, inclusive_scan, reduce }; +struct is_vec : std::false_type {}; +template +struct is_vec> : std::true_type {}; struct minimum { - Operation o; - minimum(Operation op) : o(op) {} - template T operator()(T x) { - switch (o) { - case Operation::exclusive_scan: { - return sub_group_scan_exclusive_min(x); - } - case Operation::inclusive_scan: { - return sub_group_scan_inclusive_min(x); - } - case Operation::reduce: { - return sub_group_reduce_min(x); - } - } + template + static typename std::enable_if< + !std::is_floating_point::value && std::is_signed::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupSMin(cl::__spirv::Scope::Subgroup, O, x); + } + + template + static typename std::enable_if< + !std::is_floating_point::value && std::is_unsigned::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupUMin(cl::__spirv::Scope::Subgroup, O, x); + } + + template + static typename std::enable_if::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupFMin(cl::__spirv::Scope::Subgroup, O, x); } }; struct maximum { - Operation o; - maximum(Operation op) : o(op) {} - template T operator()(T x) { - switch (o) { - case Operation::exclusive_scan: { - return sub_group_scan_exclusive_max(x); - } - case Operation::inclusive_scan: { - return sub_group_scan_inclusive_max(x); - } - case Operation::reduce: { - return sub_group_reduce_max(x); - } - } + template + static typename std::enable_if< + !std::is_floating_point::value && std::is_signed::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupSMax(cl::__spirv::Scope::Subgroup, O, x); + } + + template + static typename std::enable_if< + !std::is_floating_point::value && std::is_unsigned::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupUMax(cl::__spirv::Scope::Subgroup, O, x); + } + + template + static typename std::enable_if::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupFMax(cl::__spirv::Scope::Subgroup, O, x); } }; struct plus { - Operation o; - plus(Operation op) : o(op) {} - template T operator()(T x) { - switch (o) { - case Operation::exclusive_scan: { - return sub_group_scan_exclusive_add(x); - } - case Operation::inclusive_scan: { - return sub_group_scan_inclusive_add(x); - } - case Operation::reduce: { - return sub_group_reduce_add(x); - } - } + template + static typename std::enable_if< + !std::is_floating_point::value && std::is_integral::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupIAdd(cl::__spirv::Scope::Subgroup, O, x); + } + template + static typename std::enable_if::value, T>::type + calc(T x) { + return cl::__spirv::OpGroupFAdd(cl::__spirv::Scope::Subgroup, O, x); } }; struct sub_group { /* --- common interface members --- */ id<1> get_local_id() const { - return get_sub_group_local_id(); //*cl::__spirv::BuiltInSubgroupLocalInvocationId(); - } - range<1> get_local_range() const { - return get_sub_group_size(); // cl::__spirv::BuiltInSubgroupSize(); + return cl::__spirv::VarSubgroupLocalInvocationId; } + range<1> get_local_range() const { return cl::__spirv::VarSubgroupSize; } range<1> get_max_local_range() const { - return get_max_sub_group_size(); // cl::__spirv::BuiltInSubgroupMaxSize(); + return cl::__spirv::VarSubgroupMaxSize; } - id<1> get_group_id() const { - return get_sub_group_id(); // cl::__spirv::BuiltInSubgroupId(); - } + id<1> get_group_id() const { return cl::__spirv::VarSubgroupId; } - size_t get_group_range() const { - return get_num_sub_groups(); // cl::__spirv::BuiltInNumSubgroups(); - } + unsigned int get_group_range() const { return cl::__spirv::VarNumSubgroups; } - size_t get_uniform_group_range() const { - return get_enqueued_num_sub_groups(); // cl::__spirv::BuiltInNumEnqueuedSubgroups(); + unsigned int get_uniform_group_range() const { + return cl::__spirv::VarNumEnqueuedSubgroups; } /* --- vote / ballot functions --- */ - bool any(bool predicate) { return sub_group_any(predicate); } + bool any(bool predicate) { + return cl::__spirv::OpGroupAny(cl::__spirv::Scope::Subgroup, predicate); + } - bool all(bool predicate) { return sub_group_all(predicate); } + bool all(bool predicate) { + return cl::__spirv::OpGroupAll(cl::__spirv::Scope::Subgroup, predicate); + } /* --- collectives --- */ - template T broadcast(T x, id<1> local_id) { - return sub_group_broadcast(x, local_id.get(0)); + template + T broadcast(typename std::enable_if::value, T>::type x, + id<1> local_id) { + return cl::__spirv::OpGroupBroadcast(cl::__spirv::Scope::Subgroup, x, + local_id.get(0)); } - template T reduce(T x) { - BinaryOperation o(Operation::reduce); - return o(x); + template + T reduce(typename std::enable_if::value, T>::type x) { + return BinaryOperation::template calc(x); } - template T exclusive_scan(T x) { - BinaryOperation o(Operation::exclusive_scan); - return o(x); + template + T exclusive_scan( + typename std::enable_if::value, T>::type x) { + return BinaryOperation::template calc(x); } - template T inclusive_scan(T x) { - BinaryOperation o(Operation::inclusive_scan); - return o(x); + template + T inclusive_scan( + typename std::enable_if::value, T>::type x) { + return BinaryOperation::template calc(x); } /* --- one - input shuffles --- */ /* indices in [0 , sub - group size ) */ - template T shuffle(T x, id<1> local_id) { - return intel_sub_group_shuffle(x, local_id.get(0)); + template + typename std::enable_if::value, T>::type + shuffle(T x, id<1> local_id) { + return cl::__spirv::OpSubgroupShuffleINTEL(x, local_id.get(0)); } - template T shuffle_down(T x, uint32_t delta) { - return intel_sub_group_shuffle_down(x, x, delta); + template + typename std::enable_if::value, T>::type shuffle(T x, + id<1> local_id) { + return cl::__spirv::OpSubgroupShuffleINTEL((typename T::vector_t)x, + local_id.get(0)); } - template T shuffle_up(T x, uint32_t delta) { - return intel_sub_group_shuffle_up(x, x, delta); + template + typename std::enable_if::value, T>::type + shuffle_down(T x, uint32_t delta) { + return shuffle_down(x, x, delta); } - template T shuffle_xor(T x, id<1> value) { - return intel_sub_group_shuffle_xor(x, value.get(0)); + template + typename std::enable_if::value, T>::type + shuffle_down(T x, uint32_t delta) { + return shuffle_down(x, x, delta); } - /* --- two - input shuffles --- */ - /* indices in [0 , 2* sub - group size ) */ - template T shuffle(T x, T y, id<1> local_id) { - return intel_sub_group_shuffle_down( - x, y, local_id.get(0) - get_local_id().get(0)); + template + typename std::enable_if::value, T>::type + shuffle_up(T x, uint32_t delta) { + return shuffle_up(x, x, delta); } - template T shuffle_down(T current, T next, uint32_t delta) { - return intel_sub_group_shuffle_down(current, next, delta); - } - template T shuffle_up(T previous, T current, uint32_t delta) { - return intel_sub_group_shuffle_up(previous, current, delta); + template + typename std::enable_if::value, T>::type + shuffle_up(T x, uint32_t delta) { + return shuffle_up(x, x, delta); } - /* --- sub - group load / stores --- */ - /* these can map to SIMD or block read / write hardware where available */ - - template - typename std::enable_if::type - load(const multi_ptr src) { - uint t = intel_sub_group_block_read((const __global uint *)src.get()); - return *((T *)&t); + template + typename std::enable_if::value, T>::type + shuffle_xor(T x, id<1> value) { + return cl::__spirv::OpSubgroupShuffleXorINTEL(x, (uint32_t)value.get(0)); } - template - typename std::enable_if::type - load(const multi_ptr src) { - ushort t = - intel_sub_group_block_read_us((const __global ushort *)src.get()); - return *((T *)&t); + template + typename std::enable_if::value, T>::type shuffle_xor(T x, + id<1> value) { + return cl::__spirv::OpSubgroupShuffleXorINTEL((typename T::vector_t)x, + (uint32_t)value.get(0)); } - template - typename std::enable_if::type - load(const multi_ptr src) { - uint t = intel_sub_group_block_read((const __global uint *)src.get()); - return *((T *)&t); + /* --- two - input shuffles --- */ + /* indices in [0 , 2* sub - group size ) */ + template + typename std::enable_if::value, T>::type + shuffle(T x, T y, id<1> local_id) { + return cl::__spirv::OpSubgroupShuffleDownINTEL( + x, y, local_id.get(0) - get_local_id().get(0)); } - template - typename std::enable_if::type - load(const multi_ptr src) { - uint t = intel_sub_group_block_read_us((const __global ushort *)src.get()); - return *((T *)&t); + template + typename std::enable_if::value, T>::type shuffle(T x, T y, + id<1> local_id) { + return cl::__spirv::OpSubgroupShuffleDownINTEL( + (typename T::vector_t)x, (typename T::vector_t)y, + local_id.get(0) - get_local_id().get(0)); } - template - vec::type, N> - load(const multi_ptr src) { - uint2 t = intel_sub_group_block_read2((const __global uint *)src.get()); - return *((typename vec::vector_t *)(&t)); + template + typename std::enable_if::value, T>::type + shuffle_down(T current, T next, uint32_t delta) { + return cl::__spirv::OpSubgroupShuffleDownINTEL(current, next, delta); } - template - vec::type, - N> - load(const multi_ptr src) { - ushort2 t = - intel_sub_group_block_read_us2((const __global ushort *)src.get()); - return *((typename vec::vector_t *)(&t)); + template + typename std::enable_if::value, T>::type + shuffle_down(T current, T next, uint32_t delta) { + return cl::__spirv::OpSubgroupShuffleDownINTEL( + (typename T::vector_t)current, (typename T::vector_t)next, delta); } - template - vec::type, N> - load(const multi_ptr src) { - uint4 t = intel_sub_group_block_read4((const __global uint *)src.get()); - return *((typename vec::vector_t *)(&t)); + template + typename std::enable_if::value, T>::type + shuffle_up(T previous, T current, uint32_t delta) { + return cl::__spirv::OpSubgroupShuffleUpINTEL(previous, current, delta); } - template - vec::type, - N> - load(const multi_ptr src) { - ushort4 t = - intel_sub_group_block_read_us4((const __global ushort *)src.get()); - return *((typename vec::vector_t *)(&t)); + template + typename std::enable_if::value, T>::type + shuffle_up(T previous, T current, uint32_t delta) { + return cl::__spirv::OpSubgroupShuffleUpINTEL( + (typename T::vector_t)previous, (typename T::vector_t)current, delta); } - template - vec::type, N> + /* --- sub - group load / stores --- */ + /* these can map to SIMD or block read / write hardware where available */ + + template + typename std::enable_if<(sizeof(T) == sizeof(uint32_t) || + sizeof(T) == sizeof(uint16_t)) && + Space == access::address_space::global_space, + T>::type load(const multi_ptr src) { - uint8 t = intel_sub_group_block_read8((const __global uint *)src.get()); - return *((typename vec::vector_t *)(&t)); + if (sizeof(T) == sizeof(uint32_t)) { + uint32_t t = cl::__spirv::OpSubgroupBlockReadINTEL( + (const __global uint32_t *)src.get()); + return *((T *)(&t)); + } + uint16_t t = cl::__spirv::OpSubgroupBlockReadINTEL( + (const __global uint16_t *)src.get()); + return *((T *)(&t)); } template - vec::type, + vec::type, N> load(const multi_ptr src) { - ushort8 t = - intel_sub_group_block_read_us8((const __global ushort *)src.get()); + if (N == 1) { + return load(src); + } + if (sizeof(T) == sizeof(uint32_t)) { + typedef uint32_t ocl_t __attribute__((ext_vector_type(N))); + + ocl_t t = cl::__spirv::OpSubgroupBlockReadINTEL( + (const __global uint32_t *)src.get()); + return *((typename vec::vector_t *)(&t)); + } + typedef uint16_t ocl_t __attribute__((ext_vector_type(N))); + + ocl_t t = cl::__spirv::OpSubgroupBlockReadINTEL( + (const __global uint16_t *)src.get()); return *((typename vec::vector_t *)(&t)); } template void store(multi_ptr dst, - const typename std::enable_if::type &x) { - intel_sub_group_block_write((__global uint *)dst.get(), *((uint *)&x)); - } - - template - void store( - multi_ptr dst, - const typename std::enable_if::type &x) { - intel_sub_group_block_write_us((__global ushort *)dst.get(), - *((ushort *)&x)); + const typename std::enable_if< + (sizeof(T) == sizeof(uint32_t) || sizeof(T) == sizeof(uint16_t)) && + Space == access::address_space::global_space, + T>::type &x) { + if (sizeof(T) == sizeof(uint32_t)) { + cl::__spirv::OpSubgroupBlockWriteINTEL( + (__global uint32_t *)dst.get(), *((uint32_t *)&x)); + } else { + cl::__spirv::OpSubgroupBlockWriteINTEL( + (__global uint16_t *)dst.get(), *((uint16_t *)&x)); + } } template void store(multi_ptr dst, - const typename std::enable_if::type &x) { - intel_sub_group_block_write((__global uint *)dst.get(), *((uint *)&x)); - } - - template - void - store(multi_ptr dst, - const typename std::enable_if::type &x) { - intel_sub_group_block_write_us((__global ushort *)dst.get(), - *((ushort *)&x)); + const vec::type, N> &x) { + store(dst, x); } template void store( multi_ptr dst, - const vec< - typename std::enable_if::type, - N> &x) { - typename vec::vector_t t = x; - intel_sub_group_block_write2((__global uint *)dst.get(), *((uint2 *)&t)); - } - template - void - store(multi_ptr dst, - const vec::type, - N> &x) { - typename vec::vector_t t = x; - intel_sub_group_block_write_us2((__global ushort *)dst.get(), - *((ushort2 *)&t)); - } - - template - void store( - multi_ptr dst, - const vec< - typename std::enable_if::type, - N> &x) { - typename vec::vector_t t = x; - intel_sub_group_block_write4((__global uint *)dst.get(), *((uint4 *)&t)); - } - - template - void - store(multi_ptr dst, - const vec::type, - N> &x) { - typename vec::vector_t t = x; - intel_sub_group_block_write_us4((__global ushort *)dst.get(), - *((ushort4 *)&t)); - } - - template - void store( - multi_ptr dst, - const vec< - typename std::enable_if::type, - N> &x) { - typename vec::vector_t t = x; - intel_sub_group_block_write8((__global uint *)dst.get(), *((uint8 *)&t)); - } - - template - void - store(multi_ptr dst, - const vec::type, - N> &x) { - typename vec::vector_t t = x; - intel_sub_group_block_write_us8((__global ushort *)dst.get(), - *((ushort8 *)&t)); + const vec::type, + N> &x) { + if (sizeof(T) == sizeof(uint32_t)) { + typedef uint32_t ocl_t __attribute__((ext_vector_type(N))); + cl::__spirv::OpSubgroupBlockWriteINTEL((__global uint32_t *)dst.get(), + *((ocl_t *)&x)); + } else { + typedef uint16_t ocl_t __attribute__((ext_vector_type(N))); + cl::__spirv::OpSubgroupBlockWriteINTEL((__global uint16_t *)dst.get(), + *((ocl_t *)&x)); + } } /* --- synchronization functions --- */ void barrier(access::fence_space accessSpace = access::fence_space::global_and_local) const { - cl::sycl::detail::cl_mem_fence_flags flags; + uint32_t flags = cl::__spirv::MemorySemantics::SequentiallyConsistent; switch (accessSpace) { - case access::fence_space::local_space: - flags = cl::sycl::detail::CLK_LOCAL_MEM_FENCE; - break; case access::fence_space::global_space: - flags = cl::sycl::detail::CLK_GLOBAL_MEM_FENCE; + flags |= cl::__spirv::MemorySemantics::CrossWorkgroupMemory; + break; + case access::fence_space::local_space: + flags |= cl::__spirv::MemorySemantics::SubgroupMemory; break; case access::fence_space::global_and_local: default: - flags = cl::sycl::detail::CLK_LOCAL_MEM_FENCE | - cl::sycl::detail::CLK_GLOBAL_MEM_FENCE; + flags |= cl::__spirv::MemorySemantics::CrossWorkgroupMemory | + cl::__spirv::MemorySemantics::SubgroupMemory; break; } - ::sub_group_barrier(flags); + cl::__spirv::OpControlBarrier(cl::__spirv::Scope::Subgroup, + cl::__spirv::Scope::Workgroup, flags); } protected: diff --git a/sycl/include/CL/sycl/intel/sub_group_host.hpp b/sycl/include/CL/sycl/intel/sub_group_host.hpp index f9c43e1cbf795..ef0208f46d520 100644 --- a/sycl/include/CL/sycl/intel/sub_group_host.hpp +++ b/sycl/include/CL/sycl/intel/sub_group_host.hpp @@ -1,9 +1,8 @@ //==- sub_group_host.hpp --- SYCL sub-group for host device ---------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/item.hpp b/sycl/include/CL/sycl/item.hpp index a6c860d73d5da..30ff70b90d716 100644 --- a/sycl/include/CL/sycl/item.hpp +++ b/sycl/include/CL/sycl/item.hpp @@ -1,9 +1,8 @@ //==------------ item.hpp --- SYCL iteration item --------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp index ca3d7aae19c39..424e0abe5852c 100644 --- a/sycl/include/CL/sycl/kernel.hpp +++ b/sycl/include/CL/sycl/kernel.hpp @@ -1,9 +1,8 @@ //==--------------- kernel.hpp --- SYCL kernel -----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/macro.hpp b/sycl/include/CL/sycl/macro.hpp index 526bf234c9a0e..88f316c596b18 100644 --- a/sycl/include/CL/sycl/macro.hpp +++ b/sycl/include/CL/sycl/macro.hpp @@ -1,9 +1,8 @@ //==-------------- macro.hpp - SYCL macro header ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/math.hpp b/sycl/include/CL/sycl/math.hpp index 0b88b3de55ff1..822c38a4aacc8 100644 --- a/sycl/include/CL/sycl/math.hpp +++ b/sycl/include/CL/sycl/math.hpp @@ -1,9 +1,8 @@ //==----------- math.hpp - SYCL math functions ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -175,11 +174,21 @@ GEN_FUNC_OF_THREE_ARG(mad, double, double, double, double) // genfloatf exp (genfloatf x) GEN_FUNC_OF_ONE_ARG(native_exp, float, float) -// genfloatf fabs (genfloatf x) +// genfloat fabs (genfloatf x) GEN_FUNC_OF_ONE_ARG(fabs, float, float) GEN_FUNC_OF_ONE_ARG(fabs, double, double) // GEN_FUNC_OF_ONE_ARG(fabs, half, half) +// genfloat floor (genfloat x) +GEN_FUNC_OF_ONE_ARG(floor, float, float) +GEN_FUNC_OF_ONE_ARG(floor, double, double) +// GEN_FUNC_OF_ONE_ARG(floor, half, half) + +// genfloat ceil (genfloat x) +GEN_FUNC_OF_ONE_ARG(ceil, float, float) +GEN_FUNC_OF_ONE_ARG(ceil, double, double) +// GEN_FUNC_OF_ONE_ARG(ceil, half, half) + /* --------------- 4.13.4 Integer functions. Device version -----------------*/ // geninteger max (geninteger x, geninteger y) GEN_FUNC_OF_TWO_ARG(max, char, char, char) @@ -293,6 +302,12 @@ template T sqrt(T x) { template T fabs(T x) { return __sycl_std::fabs(x); } +template T floor(T x) { + return __sycl_std::floor(x); +} +template T ceil(T x) { + return __sycl_std::ceil(x); +} namespace native { template T exp(T x) { #ifdef __SYCL_DEVICE_ONLY__ diff --git a/sycl/include/CL/sycl/multi_ptr.hpp b/sycl/include/CL/sycl/multi_ptr.hpp index 357aa1cd26f57..185ace0a2e754 100644 --- a/sycl/include/CL/sycl/multi_ptr.hpp +++ b/sycl/include/CL/sycl/multi_ptr.hpp @@ -1,9 +1,8 @@ //==------------ multi_ptr.hpp - SYCL multi_ptr class ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/nd_item.hpp b/sycl/include/CL/sycl/nd_item.hpp index 5d34652dbcae6..cebaa1596e3f3 100644 --- a/sycl/include/CL/sycl/nd_item.hpp +++ b/sycl/include/CL/sycl/nd_item.hpp @@ -1,9 +1,8 @@ //==--------- nd_item.hpp --- SYCL iteration nd_item -----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/nd_range.hpp b/sycl/include/CL/sycl/nd_range.hpp index 689bc256dc4b5..a502328c91bb1 100644 --- a/sycl/include/CL/sycl/nd_range.hpp +++ b/sycl/include/CL/sycl/nd_range.hpp @@ -1,9 +1,8 @@ //==-------- nd_range.hpp --- SYCL iteration nd_range ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/platform.hpp b/sycl/include/CL/sycl/platform.hpp index d42ffbf5b8824..f41ca4b1b45b5 100644 --- a/sycl/include/CL/sycl/platform.hpp +++ b/sycl/include/CL/sycl/platform.hpp @@ -1,9 +1,8 @@ //==---------------- platform.hpp - SYCL platform --------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/pointers.hpp b/sycl/include/CL/sycl/pointers.hpp index f8077cadd6a24..b9df17c6a6b72 100644 --- a/sycl/include/CL/sycl/pointers.hpp +++ b/sycl/include/CL/sycl/pointers.hpp @@ -1,9 +1,8 @@ //==------------ pointers.hpp - SYCL pointers classes ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/program.hpp b/sycl/include/CL/sycl/program.hpp index 5db53194f22e6..1519e2c2412ac 100644 --- a/sycl/include/CL/sycl/program.hpp +++ b/sycl/include/CL/sycl/program.hpp @@ -1,9 +1,8 @@ //==--------------- program.hpp --- SYCL program ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/property_list.hpp b/sycl/include/CL/sycl/property_list.hpp index 706be7ea2a04e..7b562934ef88f 100644 --- a/sycl/include/CL/sycl/property_list.hpp +++ b/sycl/include/CL/sycl/property_list.hpp @@ -1,9 +1,8 @@ //==--------- property_list.hpp --- SYCL property list ---------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 824230607cf36..c14f1e75f9f63 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -1,9 +1,8 @@ //==-------------------- queue.hpp - SYCL queue ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -11,6 +10,7 @@ #include #include +#include #include #include diff --git a/sycl/include/CL/sycl/range.hpp b/sycl/include/CL/sycl/range.hpp index 80e601196ff08..0359b648f66ea 100644 --- a/sycl/include/CL/sycl/range.hpp +++ b/sycl/include/CL/sycl/range.hpp @@ -1,9 +1,8 @@ //==----------- range.hpp --- SYCL iteration range -------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/stl.hpp b/sycl/include/CL/sycl/stl.hpp index 5322fee535501..d946c48fdfbe7 100644 --- a/sycl/include/CL/sycl/stl.hpp +++ b/sycl/include/CL/sycl/stl.hpp @@ -1,9 +1,8 @@ //==----------- stl.hpp - basic STL implementation -------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/swizzles.def b/sycl/include/CL/sycl/swizzles.def index 4538fd0203a7b..48909b2245d09 100644 --- a/sycl/include/CL/sycl/swizzles.def +++ b/sycl/include/CL/sycl/swizzles.def @@ -1,9 +1,8 @@ //==---------------- swizzles.def --- SYCL types ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp index 608f9537f578f..d3a33e057b04c 100644 --- a/sycl/include/CL/sycl/types.hpp +++ b/sycl/include/CL/sycl/types.hpp @@ -1,20 +1,51 @@ //==---------------- types.hpp --- SYCL types ------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// +// Implements vec and __swizzled_vec__ classes. + #pragma once +// Define __NO_EXT_VECTOR_TYPE_ON_HOST__ to avoid using ext_vector_type +// extension even if the host compiler supports it. The same can be +// accomplished by -D__NO_EXT_VECTOR_TYPE_ON_HOST__ command line option. +#ifndef __NO_EXT_VECTOR_TYPE_ON_HOST__ +// #define __NO_EXT_VECTOR_TYPE_ON_HOST__ +#endif + +// Check if Clang's ext_vector_type attribute is available. Host compiler +// may not be Clang, and Clang may not be built with the extension. +#ifdef __clang__ +#ifndef __has_extension +#define __has_extension(x) 0 +#endif +#ifdef __HAS_EXT_VECTOR_TYPE__ +#error "Undefine __HAS_EXT_VECTOR_TYPE__ macro" +#endif +#if __has_extension(attribute_ext_vector_type) +#define __HAS_EXT_VECTOR_TYPE__ +#endif +#endif // __clang__ + +#ifdef __SYCL_USE_EXT_VECTOR_TYPE__ +#error "Undefine __SYCL_USE_EXT_VECTOR_TYPE__ macro" +#endif +#ifdef __HAS_EXT_VECTOR_TYPE__ +#if defined(__SYCL_DEVICE_ONLY__) || !defined(__NO_EXT_VECTOR_TYPE_ON_HOST__) +#define __SYCL_USE_EXT_VECTOR_TYPE__ +#endif +#elif defined(__SYCL_DEVICE_ONLY__) +// This is a soft error. We expect the device compiler to have ext_vector_type +// support, but that should not be a hard requirement. +#error "SYCL device compiler is built without ext_vector_type support" +#endif // __HAS_EXT_VECTOR_TYPE__ + #include -#ifndef __SYCL_DEVICE_ONLY__ -#include -#include -#endif // __SYCL_DEVICE_ONLY__ // 4.10.1: Scalar data types // 4.10.2: SYCL vector types @@ -105,8 +136,9 @@ using rel_t = typename std::conditional< // Special type indicating that SwizzleOp should just read value from vector - // not trying to perform any operations. Should not be called. -template class GetOp { +template class GetOp { public: + using DataT = T; DataT getValue(size_t Index) const; DataT operator()(DataT LHS, DataT Rhs); }; @@ -114,8 +146,9 @@ template class GetOp { // Special type for working SwizzleOp with scalars, stores a scalar and gives // the scalar at any index. Provides interface is compatible with SwizzleOp // operations -template class GetScalarOp { +template class GetScalarOp { public: + using DataT = T; GetScalarOp(DataT Data) : m_Data(Data) {} DataT getValue(size_t Index) const { return m_Data; } @@ -183,9 +216,24 @@ template struct LShift { } }; +template +T convertHelper(const T &Opnd) { + if (roundingMode == rounding_mode::automatic || + roundingMode == rounding_mode::rtz) { + return static_cast(Opnd); + } + if (roundingMode == rounding_mode::rtp) { + return static_cast(ceil(Opnd)); + } + // roundingMode == rounding_mode::rtn + return static_cast(floor(Opnd)); +} + } // namespace detail -template class vec { +template class vec { + using DataT = Type; + // This represent type of underlying value. There should be only one field // in the class, so vec should be equal to float16 in memory. using DataType = @@ -338,7 +386,7 @@ template class vec { void dump() { #ifndef __SYCL_DEVICE_ONLY__ for (int I = 0; I < NumElements; ++I) { - std::cout << " " << I << ": " << m_Data.s[I] << std::endl; + std::cout << " " << I << ": " << getValue(I) << std::endl; } std::cout << std::endl; #endif // __SYCL_DEVICE_ONLY__ @@ -361,12 +409,20 @@ template class vec { size_t get_count() const { return NumElements; } size_t get_size() const { return sizeof(m_Data); } - // TODO: convert() for FP types. Also, check whether rounding mode handling + // TODO: convert() for FP to FP. Also, check whether rounding mode handling // is needed for integers to FP convert. - // template - // vec convert() const; + // + // Convert to same type is no-op. + template + typename std::enable_if::value, + vec>::type + convert() const { + return *this; + } + // From Integer to Integer or FP template - typename std::enable_if::value, + typename std::enable_if::value && + std::is_integral::value, vec>::type convert() const { vec Result; @@ -375,6 +431,20 @@ template class vec { } return Result; } + // From FP to Integer + template + typename std::enable_if::value && + std::is_integral::value && + std::is_floating_point::value, + vec>::type + convert() const { + vec Result; + for (size_t I = 0; I < NumElements; ++I) { + Result.setValue( + I, detail::convertHelper(getValue(I))); + } + return Result; + } template typename std::enable_if::type @@ -415,12 +485,24 @@ template class vec { #endif #define __SYCL_LOADSTORE(Space) \ void load(size_t Offset, multi_ptr Ptr) { \ - m_Data = *multi_ptr(static_cast( \ - static_cast(Ptr + Offset * NumElements))); \ + if (NumElements != 3) { \ + m_Data = *multi_ptr(static_cast( \ + static_cast(Ptr + Offset * NumElements))); \ + return; \ + } \ + for (int I = 0; I < NumElements; I++) { \ + setValue(I, *multi_ptr(Ptr + Offset * NumElements + I)); \ + } \ } \ void store(size_t Offset, multi_ptr Ptr) const { \ - *multi_ptr(static_cast( \ - static_cast(Ptr + Offset * NumElements))) = m_Data; \ + if (NumElements != 3) { \ + *multi_ptr(static_cast( \ + static_cast(Ptr + Offset * NumElements))) = m_Data; \ + return; \ + } \ + for (int I = 0; I < NumElements; I++) { \ + *multi_ptr(Ptr + Offset * NumElements + I) = getValue(I); \ + } \ } __SYCL_LOADSTORE(access::address_space::global_space) @@ -433,7 +515,7 @@ template class vec { #error "Undefine __SYCL_BINOP macro" #endif -#ifdef __SYCL_DEVICE_ONLY__ +#ifdef __SYCL_USE_EXT_VECTOR_TYPE__ #define __SYCL_BINOP(BINOP, OPASSIGN) \ vec operator BINOP(const vec &Rhs) const { \ vec Ret; \ @@ -457,7 +539,7 @@ template class vec { *this = *this BINOP vec(Rhs); \ return *this; \ } -#else // __SYCL_DEVICE_ONLY__ +#else // __SYCL_USE_EXT_VECTOR_TYPE__ #define __SYCL_BINOP(BINOP, OPASSIGN) \ vec operator BINOP(const vec &Rhs) const { \ vec Ret; \ @@ -483,7 +565,7 @@ template class vec { *this = *this BINOP vec(Rhs); \ return *this; \ } -#endif // __SYCL_DEVICE_ONLY__ +#endif // __SYCL_USE_EXT_VECTOR_TYPE__ __SYCL_BINOP(+, +=) __SYCL_BINOP(-, -=) @@ -588,21 +670,21 @@ template class vec { vec operatorHelper(const vec &Rhs) const { vec Result; -#ifdef __SYCL_DEVICE_ONLY__ +#ifdef __SYCL_USE_EXT_VECTOR_TYPE__ Operation Op; Result.m_Data = Op(m_Data, Rhs.m_Data); -#else // __SYCL_DEVICE_ONLY__ +#else // __SYCL_USE_EXT_VECTOR_TYPE__ Operation Op; for (size_t I = 0; I < NumElements; ++I) { Result.setValue(I, Op(Rhs.getValue(I), getValue(I))); } -#endif // __SYCL_DEVICE_ONLY__ +#endif // __SYCL_USE_EXT_VECTOR_TYPE__ return Result; } // setValue and getValue should be able to operate on different underlying // types: enum cl_float#N , builtin vector float#N, builtin type float. -#ifdef __SYCL_DEVICE_ONLY__ +#ifdef __SYCL_USE_EXT_VECTOR_TYPE__ template ::type> void setValue(int Index, const DataT &Value, int) { @@ -614,7 +696,7 @@ template class vec { DataT getValue(int Index, int) const { return m_Data[Index]; } -#else +#else // __SYCL_USE_EXT_VECTOR_TYPE__ template ::type> void setValue(int Index, const DataT &Value, int) { @@ -626,7 +708,7 @@ template class vec { DataT getValue(int Index, int) const { return m_Data.s[Index]; } -#endif +#endif // __SYCL_USE_EXT_VECTOR_TYPE__ template ::type> @@ -727,6 +809,9 @@ template class OperationCurrentT, int... Indexes> class SwizzleOp { using DataT = typename VecT::element_type; + using CommonDataT = + typename std::common_type::type; using rel_t = detail::rel_t; static constexpr int getNumElements() { return sizeof...(Indexes); } @@ -751,15 +836,13 @@ class SwizzleOp { OperationCurrentT, Indexes...>, OperationCurrentT_, Idx_...>; - template - using EnableIfOneIndex = - typename std::enable_if<1 == IdxNum && - SwizzleOp::getNumElements() == IdxNum>::type; + template + using EnableIfOneIndex = typename std::enable_if< + 1 == IdxNum && SwizzleOp::getNumElements() == IdxNum, T>::type; - template - using EnableIfMultipleIndexes = - typename std::enable_if<1 != IdxNum && - SwizzleOp::getNumElements() == IdxNum>::type; + template + using EnableIfMultipleIndexes = typename std::enable_if< + 1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>::type; template using EnableIfScalarType = @@ -1158,10 +1241,16 @@ class SwizzleOp { // Begin hi/lo, even/odd, xyzw, and rgba swizzles. private: - // Indexer used in the swizzles.def. C++14 + // Indexer used in the swizzles.def. C++11 way, a bit more verbose + // than C++14 way. + struct IndexerHelper { + static const constexpr int IDXs[] = {Indexes...}; + static constexpr int get(int index) { + return IDXs[index >= getNumElements() ? 0 : index]; + } + }; static constexpr int Indexer(int index) { - const constexpr int IDXs[] = {Indexes...}; - return IDXs[index >= getNumElements() ? 0 : index]; + return IndexerHelper::get(index); } public: @@ -1229,8 +1318,22 @@ class SwizzleOp { m_RightOperation(std::move(Rhs.m_RightOperation)) {} // Either performing CurrentOperation on results of left and right operands - // or reading values from actual vector. - DataT getValue(size_t Index) const { + // or reading values from actual vector. Perform implicit type conversion when + // the number of elements == 1 + + template + CommonDataT getValue(EnableIfOneIndex Index) const { + if (std::is_same, GetOp>::value) { + std::array Idxs{Indexes...}; + return m_Vector->getValue(Idxs[Index]); + } + auto Op = OperationCurrentT(); + return Op(m_LeftOperation.getValue(Index), + m_RightOperation.getValue(Index)); + } + + template + DataT getValue(EnableIfMultipleIndexes Index) const { if (std::is_same, GetOp>::value) { std::array Idxs{Indexes...}; return m_Vector->getValue(Idxs[Index]); @@ -1275,7 +1378,7 @@ class SwizzleOp { template \ typename std::enable_if::value, vec>::type \ operator BINOP(const T &Lhs, const vec &Rhs) { \ - return vec(static_cast(Lhs)) BINOP Rhs; \ + return vec(Lhs) BINOP Rhs; \ } \ template class OperationCurrentT, int... Indexes, \ @@ -1367,7 +1470,7 @@ __SYCL_RELLOGOP(||) } // namespace sycl } // namespace cl -#ifdef __SYCL_DEVICE_ONLY__ +#ifdef __SYCL_USE_EXT_VECTOR_TYPE__ typedef char __char_t; typedef char __char2_vec_t __attribute__((ext_vector_type(2))); typedef char __char3_vec_t __attribute__((ext_vector_type(3))); @@ -1461,7 +1564,7 @@ typedef double __double16_vec_t __attribute__((ext_vector_type(16))); #define GET_CL_TYPE(target, num) __##target##num##_vec_t #define GET_SCALAR_CL_TYPE(target) target -#else // __SYCL_DEVICE_ONLY__ +#else // __SYCL_USE_EXT_VECTOR_TYPE__ // For signed char. OpenCL doesn't have any type about `signed char`, therefore // we use type alias of cl_char instead. using cl_schar = cl_char; @@ -1473,7 +1576,7 @@ using cl_schar16 = cl_char16; #define GET_CL_TYPE(target, num) cl_##target##num #define GET_SCALAR_CL_TYPE(target) cl_##target -#endif // __SYCL_DEVICE_ONLY__ +#endif // __SYCL_USE_EXT_VECTOR_TYPE__ namespace cl { namespace sycl { @@ -1484,6 +1587,12 @@ namespace sycl { using DataType = GET_CL_TYPE(base, num); \ }; +#define DECLARE_LONGLONG_CONVERTER(base, num) \ + template <> class BaseCLTypeConverter { \ + public: \ + using DataType = ::GET_CL_TYPE(base, num); \ + }; + #define DECLARE_VECTOR_CONVERTERS(base) \ namespace detail { \ DECLARE_CONVERTER(base, 2) \ @@ -1497,6 +1606,19 @@ namespace sycl { }; \ } // namespace detail +#define DECLARE_VECTOR_LONGLONG_CONVERTERS(base) \ + namespace detail { \ + DECLARE_LONGLONG_CONVERTER(base, 2) \ + DECLARE_LONGLONG_CONVERTER(base, 3) \ + DECLARE_LONGLONG_CONVERTER(base, 4) \ + DECLARE_LONGLONG_CONVERTER(base, 8) \ + DECLARE_LONGLONG_CONVERTER(base, 16) \ + template <> class BaseCLTypeConverter { \ + public: \ + using DataType = GET_SCALAR_CL_TYPE(base); \ + }; \ + } // namespace detail + #define DECLARE_SYCL_VEC_WO_CONVERTERS(base) \ using cl_##base##16 = vec; \ using cl_##base##8 = vec; \ @@ -1510,11 +1632,40 @@ namespace sycl { using base##3 = cl_##base##3; \ using base##2 = cl_##base##2; +#define DECLARE_SYCL_VEC_CHAR_WO_CONVERTERS \ + using cl_char16 = vec; \ + using cl_char8 = vec; \ + using cl_char4 = vec; \ + using cl_char3 = vec; \ + using cl_char2 = vec; \ + using cl_char = signed char; \ + using char16 = vec; \ + using char8 = vec; \ + using char4 = vec; \ + using char3 = vec; \ + using char2 = vec; + +// cl_longlong/cl_ulonglong are not supported in SYCL +#define DECLARE_SYCL_VEC_LONGLONG_WO_CONVERTERS(base) \ + using base##long16 = vec; \ + using base##long8 = vec; \ + using base##long4 = vec; \ + using base##long3 = vec; \ + using base##long2 = vec; + #define DECLARE_SYCL_VEC(base) \ DECLARE_VECTOR_CONVERTERS(base) \ DECLARE_SYCL_VEC_WO_CONVERTERS(base) -DECLARE_SYCL_VEC(char) +#define DECLARE_SYCL_VEC_CHAR \ + DECLARE_VECTOR_CONVERTERS(char) \ + DECLARE_SYCL_VEC_CHAR_WO_CONVERTERS + +#define DECLARE_SYCL_VEC_LONGLONG(base) \ + DECLARE_VECTOR_LONGLONG_CONVERTERS(base) \ + DECLARE_SYCL_VEC_LONGLONG_WO_CONVERTERS(base) + +DECLARE_SYCL_VEC_CHAR DECLARE_SYCL_VEC(schar) DECLARE_SYCL_VEC(uchar) DECLARE_SYCL_VEC(short) @@ -1523,9 +1674,8 @@ DECLARE_SYCL_VEC(int) DECLARE_SYCL_VEC(uint) DECLARE_SYCL_VEC(long) DECLARE_SYCL_VEC(ulong) -// TODO: Fix long long and unsigned long long. -// DECLARE_SYCL_VEC(longlong) -// DECLARE_SYCL_VEC(ulonglong) +DECLARE_SYCL_VEC_LONGLONG(long) +DECLARE_SYCL_VEC_LONGLONG(ulong) DECLARE_SYCL_VEC(float) DECLARE_SYCL_VEC(double) // TODO: Fix half. diff --git a/sycl/include/CL/sycl/version.hpp.in b/sycl/include/CL/sycl/version.hpp.in new file mode 100644 index 0000000000000..d4cf654f3afc9 --- /dev/null +++ b/sycl/include/CL/sycl/version.hpp.in @@ -0,0 +1,9 @@ +//==------ version.hpp --- SYCL compiler version macro ---------*- C++ -*---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#cmakedefine __SYCL_COMPILER_VERSION ${__SYCL_COMPILER_VERSION} diff --git a/sycl/source/context.cpp b/sycl/source/context.cpp index 816c491a4e2c9..973f166443d3c 100644 --- a/sycl/source/context.cpp +++ b/sycl/source/context.cpp @@ -1,31 +1,81 @@ //==---------------- context.cpp - SYCL context ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +// 4.6.2 Context class namespace cl { namespace sycl { - context::context(const vector_class &deviceList, - async_handler asyncHandler) { - if (deviceList.empty()) - throw invalid_parameter_error("First argument deviceList is empty."); - - if (deviceList[0].is_host()) { - impl = - std::make_shared(deviceList[0], asyncHandler); - } else { - impl = std::make_shared(deviceList, asyncHandler); - } - } +context::context(const async_handler &asyncHandler) + : context(default_selector().select_device(), asyncHandler) {} + +context::context(const device &dev, async_handler asyncHandler) + : context(vector_class(1, dev), asyncHandler) {} - context::context(cl_context clContext, async_handler asyncHandler) { - impl = std::make_shared(clContext, asyncHandler); +context::context(const platform &plt, async_handler asyncHandler) + : context(plt.get_devices(), asyncHandler) {} + +context::context(const vector_class &deviceList, + async_handler asyncHandler) { + if (deviceList.empty()) { + throw invalid_parameter_error("First argument deviceList is empty."); + } + if (deviceList[0].is_host()) { + impl = std::make_shared(deviceList[0], asyncHandler); + } else { + // TODO also check that devices belongs to the same platform + impl = std::make_shared(deviceList, asyncHandler); } +} + +context::context(cl_context clContext, async_handler asyncHandler) { + impl = std::make_shared(clContext, asyncHandler); +} + +template <> cl_uint context::get_info() const { + return impl->get_info(); +} + +template <> +cl::sycl::platform context::get_info() const { + return impl->get_info(); +} + +template <> +vector_class +context::get_info() const { + return impl->get_info(); +} + +bool context::operator==(const context &rhs) const { return impl == rhs.impl; } + +bool context::operator!=(const context &rhs) const { return !(*this == rhs); } + +cl_context context::get() const { return impl->get(); } + +bool context::is_host() const { return impl->is_host(); } + +platform context::get_platform() const { return impl->get_platform(); } + +vector_class context::get_devices() const { + return impl->get_devices(); +} + } // namespace sycl } // namespace cl diff --git a/sycl/source/detail/common.cpp b/sycl/source/detail/common.cpp index 8570f7203ef38..d9841251aa195 100644 --- a/sycl/source/detail/common.cpp +++ b/sycl/source/detail/common.cpp @@ -1,9 +1,8 @@ //==----------- common.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp new file mode 100644 index 0000000000000..a242069ff16e9 --- /dev/null +++ b/sycl/source/detail/context_impl.cpp @@ -0,0 +1,110 @@ +//==---------------- context_impl.cpp - SYCL context -----------*- C++ -*---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// ===--------------------------------------------------------------------=== // + +#include +#include +#include +#include +#include +#include +#include +#include + +namespace cl { +namespace sycl { +namespace detail { + +context_impl::context_impl(const device &Device, async_handler AsyncHandler) + : m_AsyncHandler(AsyncHandler), m_Devices(1, Device), m_ClContext(nullptr), + m_Platform(), m_OpenCLInterop(false), m_HostContext(true) {} + +context_impl::context_impl(const vector_class Devices, + async_handler AsyncHandler) + : m_AsyncHandler(AsyncHandler), m_Devices(Devices), m_ClContext(nullptr), + m_Platform(), m_OpenCLInterop(true), m_HostContext(false) { + m_Platform = m_Devices[0].get_platform(); + vector_class DeviceIds; + for (const auto &D : m_Devices) { + DeviceIds.push_back(D.get()); + } + cl_int Err; + m_ClContext = + clCreateContext(0, DeviceIds.size(), DeviceIds.data(), 0, 0, &Err); + // TODO catch an exception and put it to list of asynchronous exceptions + CHECK_OCL_CODE(Err); +} + +context_impl::context_impl(cl_context ClContext, async_handler AsyncHandler) + : m_AsyncHandler(AsyncHandler), m_Devices(), m_ClContext(ClContext), + m_Platform(), m_OpenCLInterop(true), m_HostContext(false) { + vector_class DeviceIds; + size_t DevicesBuffer = 0; + // TODO catch an exception and put it to list of asynchronous exceptions + CHECK_OCL_CODE(clGetContextInfo(m_ClContext, CL_CONTEXT_DEVICES, 0, nullptr, + &DevicesBuffer)); + DeviceIds.resize(DevicesBuffer / sizeof(cl_device_id)); + // TODO catch an exception and put it to list of asynchronous exceptions + CHECK_OCL_CODE(clGetContextInfo(m_ClContext, CL_CONTEXT_DEVICES, + DevicesBuffer, &DeviceIds[0], nullptr)); + + for (auto Dev : DeviceIds) { + m_Devices.emplace_back(Dev); + } + // TODO What if m_Devices if empty? m_Devices[0].get_platform() + m_Platform = platform(m_Devices[0].get_platform()); + // TODO catch an exception and put it to list of asynchronous exceptions + CHECK_OCL_CODE(clRetainContext(m_ClContext)); +} + +cl_context context_impl::get() const { + if (m_OpenCLInterop) { + // TODO catch an exception and put it to list of asynchronous exceptions + CHECK_OCL_CODE(clRetainContext(m_ClContext)); + return m_ClContext; + } + throw invalid_object_error( + "This instance of event doesn't support OpenCL interoperability."); +} + +bool context_impl::is_host() const { return m_HostContext || !m_OpenCLInterop; } +platform context_impl::get_platform() const { return m_Platform; } +vector_class context_impl::get_devices() const { return m_Devices; } + +context_impl::~context_impl() { + if (m_OpenCLInterop) { + // TODO replace CHECK_OCL_CODE_NO_EXC to CHECK_OCL_CODE and + // catch an exception and put it to list of asynchronous exceptions + CHECK_OCL_CODE_NO_EXC(clReleaseContext(m_ClContext)); + } +} + +const async_handler &context_impl::get_async_handler() const { + return m_AsyncHandler; +} + +template <> +cl_uint context_impl::get_info() const { + if (is_host()) { + return 0; + } + return get_context_info_cl::_(this->get()); +} +template <> platform context_impl::get_info() const { + return get_platform(); +} +template <> +vector_class +context_impl::get_info() const { + return get_devices(); +} + +cl_context &context_impl::getHandleRef() { return m_ClContext; } + +} // namespace detail +} // namespace sycl +} // namespace cl diff --git a/sycl/source/detail/device_info.cpp b/sycl/source/detail/device_info.cpp index 0125a054c0eb0..c578a255be2a2 100644 --- a/sycl/source/detail/device_info.cpp +++ b/sycl/source/detail/device_info.cpp @@ -1,9 +1,8 @@ //==----------- device_info.cpp --------------------------------*- C ++-*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/event_impl.cpp b/sycl/source/detail/event_impl.cpp index 6d9792da00cfd..5d00da3d2fd61 100644 --- a/sycl/source/detail/event_impl.cpp +++ b/sycl/source/detail/event_impl.cpp @@ -1,9 +1,8 @@ //==---------------- event_impl.cpp - SYCL event ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -72,6 +71,14 @@ void event_impl::wait( simple_scheduler::Scheduler::getInstance().waitForEvent(Self); } + +void event_impl::wait_and_throw( + std::shared_ptr Self) { + wait(Self); + cl::sycl::simple_scheduler::Scheduler::getInstance().throwForEventRecursive( + Self); +} + template <> cl_ulong event_impl::get_profiling_info() const { diff --git a/sycl/source/detail/force_device.cpp b/sycl/source/detail/force_device.cpp index ebe1b3ffe5551..e138ad02febf9 100644 --- a/sycl/source/detail/force_device.cpp +++ b/sycl/source/detail/force_device.cpp @@ -1,9 +1,8 @@ //==---------- force_device.cpp - Forcing SYCL device ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/force_device.hpp b/sycl/source/detail/force_device.hpp index d0a9da0a85ed4..855bd310e3c87 100644 --- a/sycl/source/detail/force_device.hpp +++ b/sycl/source/detail/force_device.hpp @@ -1,9 +1,8 @@ //==---------- force_device.hpp - Forcing SYCL device ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/helpers.cpp b/sycl/source/detail/helpers.cpp index a45ab182fd5e6..2c6d9abd4303a 100644 --- a/sycl/source/detail/helpers.cpp +++ b/sycl/source/detail/helpers.cpp @@ -1,23 +1,23 @@ //==---------------- helpers.cpp - SYCL helpers ---------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// -#include +#include #include #include +#include + namespace cl { namespace sycl { +using ContextImplPtr = std::shared_ptr; namespace detail { - -std::vector -getOrWaitEvents(std::vector DepEvents, - cl::sycl::context Context) { +std::vector getOrWaitEvents(std::vector DepEvents, + ContextImplPtr Context) { std::vector CLEvents; for (auto SyclEvent : DepEvents) { auto SyclEventImplPtr = detail::getSyclObjImpl(SyclEvent); diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp index 621d2bf708d2c..842ed158e6d89 100644 --- a/sycl/source/detail/kernel_impl.cpp +++ b/sycl/source/detail/kernel_impl.cpp @@ -1,9 +1,8 @@ //==------- kernel_impl.cpp --- SYCL kernel implementation -----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/kernel_info.cpp b/sycl/source/detail/kernel_info.cpp index 619e0f3a25408..32dde9e87b7be 100644 --- a/sycl/source/detail/kernel_info.cpp +++ b/sycl/source/detail/kernel_info.cpp @@ -1,9 +1,8 @@ //==-------- kernel_info.cpp - SYCL kernel info methods --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/platform_host.cpp b/sycl/source/detail/platform_host.cpp index fc29d8fd5adb8..63257a96df4ab 100644 --- a/sycl/source/detail/platform_host.cpp +++ b/sycl/source/detail/platform_host.cpp @@ -1,9 +1,8 @@ //==----------- platform_host.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/platform_info.cpp b/sycl/source/detail/platform_info.cpp index d1f2dcaeac34b..27669f36ae352 100644 --- a/sycl/source/detail/platform_info.cpp +++ b/sycl/source/detail/platform_info.cpp @@ -1,9 +1,8 @@ //==----------- platform_info.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/platform_opencl.cpp b/sycl/source/detail/platform_opencl.cpp index 9b66cb13b831e..f6cb39cdf52da 100644 --- a/sycl/source/detail/platform_opencl.cpp +++ b/sycl/source/detail/platform_opencl.cpp @@ -1,9 +1,8 @@ //==----------- platform_opencl.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp index 0d8e8e85f9ec2..6d16027ed0419 100644 --- a/sycl/source/detail/program_impl.cpp +++ b/sycl/source/detail/program_impl.cpp @@ -1,9 +1,8 @@ //==----- program_impl.cpp --- SYCL program implementation -----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index d5ea66c11c21d..b368c15daf87d 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1,15 +1,15 @@ //==------ program_manager.cpp --- SYCL program manager---------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include #include #include +#include #include #include @@ -32,7 +32,7 @@ static cl_device_id getFirstDevice(cl_context Context) { cl_uint NumDevices = 0; cl_int Err = clGetContextInfo(Context, CL_CONTEXT_NUM_DEVICES, sizeof(NumDevices), &NumDevices, - /*param_value_size_ret=*/ nullptr); + /*param_value_size_ret=*/nullptr); CHECK_OCL_CODE(Err); assert(NumDevices > 0 && "Context without devices?"); @@ -55,7 +55,7 @@ static cl_program createBinaryProgram(const cl_context Context, cl_uint NumDevices = 0; CHECK_OCL_CODE(clGetContextInfo(Context, CL_CONTEXT_NUM_DEVICES, sizeof(NumDevices), &NumDevices, - /*param_value_size_ret=*/ nullptr)); + /*param_value_size_ret=*/nullptr)); assert(NumDevices > 0 && "Only a single device is supported for AOT compilation"); #endif @@ -64,10 +64,9 @@ static cl_program createBinaryProgram(const cl_context Context, cl_int Err = CL_SUCCESS; cl_int BinaryStatus = CL_SUCCESS; size_t BinarySize = BinProg.size(); - const unsigned char *Binary = (const unsigned char *) &BinProg[0]; - cl_program Program = - clCreateProgramWithBinary(Context, 1, &Device, &BinarySize, &Binary, - &BinaryStatus, &Err); + const unsigned char *Binary = (const unsigned char *)&BinProg[0]; + cl_program Program = clCreateProgramWithBinary( + Context, 1, &Device, &BinarySize, &Binary, &BinaryStatus, &Err); CHECK_OCL_CODE(Err); return Program; @@ -112,16 +111,20 @@ static cl_program createProgram(const platform &Platform, return Program; } +cl_program ProgramManager::createOpenCLProgram(const context &Context) { + vector_class DeviceProg = getSpirvSource(); + cl_context ClContext = detail::getSyclObjImpl(Context)->getHandleRef(); + const platform &Platform = Context.get_platform(); + cl_program ClProgram = createProgram(Platform, ClContext, DeviceProg); + return ClProgram; +} + cl_program ProgramManager::getBuiltOpenCLProgram(const context &Context) { cl_program &ClProgram = m_CachedSpirvPrograms[Context]; if (!ClProgram) { vector_class DeviceProg = getSpirvSource(); - cl_context ClContext = Context.get(); - const platform &Platform = Context.get_platform(); - ClProgram = createProgram(Platform, ClContext, DeviceProg); - clReleaseContext(ClContext); - + ClProgram = createOpenCLProgram(Context); build(ClProgram); } return ClProgram; diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index 3c3d926560121..71b65a21c6f79 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -1,9 +1,8 @@ //==------------------ queue_impl.cpp - SYCL queue -------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -16,8 +15,10 @@ namespace sycl { namespace detail { template <> cl_uint queue_impl::get_info() const { cl_uint result = 0; - CHECK_OCL_CODE(clGetCommandQueueInfo(m_CommandQueue, CL_QUEUE_REFERENCE_COUNT, - sizeof(result), &result, nullptr)); + if (!is_host()) + CHECK_OCL_CODE(clGetCommandQueueInfo(m_CommandQueue, + CL_QUEUE_REFERENCE_COUNT, + sizeof(result), &result, nullptr)); return result; } diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 581a2c83ef0e2..403f15ae56f8c 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -1,9 +1,8 @@ //==----------- commands.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/source/detail/scheduler/printers.cpp b/sycl/source/detail/scheduler/printers.cpp index bae08bed4204c..56b55f2a6e2c0 100644 --- a/sycl/source/detail/scheduler/printers.cpp +++ b/sycl/source/detail/scheduler/printers.cpp @@ -1,9 +1,8 @@ //==----------- printers.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/detail/scheduler/scheduler.cpp b/sycl/source/detail/scheduler/scheduler.cpp index 08321704c0cea..982b301dfcf74 100644 --- a/sycl/source/detail/scheduler/scheduler.cpp +++ b/sycl/source/detail/scheduler/scheduler.cpp @@ -1,9 +1,8 @@ //==----------- scheduler.cpp ----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -17,6 +16,7 @@ #include #include #include +#include #include namespace cl { @@ -40,16 +40,31 @@ void Node::addInteropArg(shared_ptr_class Ptr, size_t Size, m_InteropArgs.emplace_back(Ptr, Size, ArgIndex, BufReq); } -// Waits for the event passed. -void Scheduler::waitForEvent(EventImplPtr Event) { +CommandPtr Scheduler::getCmdForEvent(EventImplPtr Event) { + // TODO: Currently, this method searches for the command in + // m_BuffersEvolution, which seems expensive, especially + // taking into account that this operation may be called + // from another loop. Need to optimize this method, for example, + // by adding a direct link from 'event' to the 'command' it + // is associated with. for (auto &BufEvolution : m_BuffersEvolution) { for (auto &Cmd : BufEvolution.second) { if (detail::getSyclObjImpl(Cmd->getEvent()) == Event) { - enqueueAndWaitForCommand(Cmd); - return; + return Cmd; } } } + return nullptr; +} + +// Waits for the event passed. +void Scheduler::waitForEvent(EventImplPtr Event) { + auto Cmd = getCmdForEvent(Event); + if (Cmd) { + enqueueAndWaitForCommand(Cmd); + return; + } + for (auto &Evnt : m_EventsWithoutRequirements) { if (Evnt == Event) { Evnt->waitInternal(); @@ -58,6 +73,51 @@ void Scheduler::waitForEvent(EventImplPtr Event) { } } +// Calls async handler for the given command Cmd and those other +// commands that Cmd depends on. +void Scheduler::throwForCmdRecursive(std::shared_ptr Cmd) { + if (Cmd == nullptr) { + return; + } + + auto QImpl = Cmd->getQueue(); + QImpl->throw_asynchronous(); + + std::vector, BufferReqPtr>> Deps = + Cmd->getDependencies(); + for (auto D : Deps) { + throwForCmdRecursive(D.first); + } +} + +// Calls async handler for the given event Event and those other +// events that Event depends on. +void Scheduler::throwForEventRecursive(EventImplPtr Event) { + auto Cmd = getCmdForEvent(Event); + if (Cmd) { + throwForCmdRecursive(Cmd); + } +} + +void Scheduler::getDepEventsRecursive( + std::unordered_set &EventsSet, + EventImplPtr Event) { + auto Cmd = getCmdForEvent(Event); + if (Cmd == nullptr) { + return; + } + + std::vector, BufferReqPtr>> Deps = + Cmd->getDependencies(); + for (auto D : Deps) { + auto DepEvent = D.first->getEvent(); + EventsSet.insert(DepEvent); + + auto DepEventImpl = cl::sycl::detail::getSyclObjImpl(DepEvent); + getDepEventsRecursive(EventsSet, DepEventImpl); + } +} + void Scheduler::print(std::ostream &Stream) const { Stream << "======================================" << std::endl; Stream << "Graph dump" << std::endl; @@ -183,6 +243,11 @@ void Scheduler::printGraphForCommand(CommandPtr Cmd, Cmd->printDot(Stream); } +Scheduler &Scheduler::getInstance() { + static Scheduler Instance; + return Instance; +} + } // namespace simple_scheduler } // namespace sycl } // namespace cl diff --git a/sycl/source/device.cpp b/sycl/source/device.cpp index 564a3fbfa2613..e847616f930c7 100644 --- a/sycl/source/device.cpp +++ b/sycl/source/device.cpp @@ -1,9 +1,8 @@ //==------------------- device.cpp -----------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/device_selector.cpp b/sycl/source/device_selector.cpp index 9df21450d95e2..2dbfc1a1a3f9a 100644 --- a/sycl/source/device_selector.cpp +++ b/sycl/source/device_selector.cpp @@ -1,9 +1,8 @@ //==------ device_selector.cpp - SYCL device selector ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/event.cpp b/sycl/source/event.cpp index 83789ed350416..60e9531e2a04a 100644 --- a/sycl/source/event.cpp +++ b/sycl/source/event.cpp @@ -1,9 +1,8 @@ //==---------------- event.cpp --- SYCL event ------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -11,9 +10,12 @@ #include #include #include +#include + #include #include +#include namespace cl { namespace sycl { @@ -31,7 +33,30 @@ cl_event event::get() { return impl->get(); } bool event::is_host() const { return impl->is_host(); } -void event::wait() const { impl->wait(impl); } +void event::wait() { impl->wait(impl); } + +void event::wait(const vector_class &EventList) { + for (auto E : EventList) { + E.wait(); + } +} + +void event::wait_and_throw() { impl->wait_and_throw(impl); } + +void event::wait_and_throw(const vector_class &EventList) { + for (auto E : EventList) { + E.wait_and_throw(); + } +} + +vector_class event::get_wait_list() { + std::unordered_set DepEventsSet; + cl::sycl::simple_scheduler::Scheduler::getInstance().getDepEventsRecursive( + DepEventsSet, impl); + + vector_class DepEventsVec(DepEventsSet.begin(), DepEventsSet.end()); + return DepEventsVec; +} event::event(std::shared_ptr event_impl) : impl(event_impl) {} diff --git a/sycl/source/exception.cpp b/sycl/source/exception.cpp index 7827a9d0add0c..9cfa5207e1f7b 100644 --- a/sycl/source/exception.cpp +++ b/sycl/source/exception.cpp @@ -1,9 +1,8 @@ //==---------------- exception.cpp - SYCL exception ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp index 0bf30b1a584fe..6a9efd9e662fc 100644 --- a/sycl/source/kernel.cpp +++ b/sycl/source/kernel.cpp @@ -1,9 +1,8 @@ //==--------------- kernel.cpp --- SYCL kernel -----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/platform.cpp b/sycl/source/platform.cpp index 3d45c56d05123..34815cfedc8e4 100644 --- a/sycl/source/platform.cpp +++ b/sycl/source/platform.cpp @@ -1,9 +1,8 @@ //==----------- platform.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index 6f9fa01bcda7d..a6df2836e3da2 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -1,9 +1,8 @@ //==-------------- queue.cpp -----------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -20,8 +19,9 @@ queue::queue(const context &syclContext, const device_selector &deviceSelector, return deviceSelector(d1) < deviceSelector(d2); }; - *this = queue(*std::max_element(Devs.begin(), Devs.end(), Comp), asyncHandler, - propList); + const device &syclDevice = *std::max_element(Devs.begin(), Devs.end(), Comp); + impl = std::make_shared(syclDevice, syclContext, + asyncHandler, propList); } queue::queue(const device &syclDevice, const async_handler &asyncHandler, diff --git a/sycl/source/spirv_ops.cpp b/sycl/source/spirv_ops.cpp index 640e85139144e..905dc6f51fc96 100644 --- a/sycl/source/spirv_ops.cpp +++ b/sycl/source/spirv_ops.cpp @@ -1,9 +1,8 @@ //===------------- spirv_ops.cpp - SPIRV operations -----------------------===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/aot/with-llvm-bc.cpp b/sycl/test/aot/with-llvm-bc.cpp index 730876d99de56..6f796f386f732 100644 --- a/sycl/test/aot/with-llvm-bc.cpp +++ b/sycl/test/aot/with-llvm-bc.cpp @@ -5,14 +5,12 @@ // // Only CPU supports LLVM IR bitcode as a binary // RUN: %CPU_RUN_PLACEHOLDER %t.out -// XFAIL: * //==----- with-llvm-bc.cpp - SYCL kernel with LLVM IR bitcode as binary ----==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/basic_tests/accessor/accessor.cpp b/sycl/test/basic_tests/accessor/accessor.cpp index 269b5c20f2f4a..3af6e1f919778 100644 --- a/sycl/test/basic_tests/accessor/accessor.cpp +++ b/sycl/test/basic_tests/accessor/accessor.cpp @@ -5,10 +5,9 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out //==----------------accessor.cpp - SYCL accessor basic test ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/accessor_static_check.cpp b/sycl/test/basic_tests/accessor_static_check.cpp index 75780e3c4f048..d4f38d080faef 100644 --- a/sycl/test/basic_tests/accessor_static_check.cpp +++ b/sycl/test/basic_tests/accessor_static_check.cpp @@ -4,10 +4,9 @@ // RUN: %clang --sycl -fsyntax-only %s //==--- accessor_static_check.cpp - Static checks for SYCL accessors -------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/accessor_syntax_only.cpp b/sycl/test/basic_tests/accessor_syntax_only.cpp index 913879fbe1c63..b37c56baeb6e3 100644 --- a/sycl/test/basic_tests/accessor_syntax_only.cpp +++ b/sycl/test/basic_tests/accessor_syntax_only.cpp @@ -1,9 +1,8 @@ //==--- accessor_syntax_only.cpp - Syntax checks for SYCL accessors --------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// // This test is supposed to check that interface of sycl::accessor diff --git a/sycl/test/basic_tests/buffer/buffer.cpp b/sycl/test/basic_tests/buffer/buffer.cpp index d1c0b85cb2f21..02450b91abbfc 100644 --- a/sycl/test/basic_tests/buffer/buffer.cpp +++ b/sycl/test/basic_tests/buffer/buffer.cpp @@ -5,10 +5,9 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out //==------------------- buffer.cpp - SYCL buffer basic test ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/basic_tests/buffer/buffer_interop.cpp b/sycl/test/basic_tests/buffer/buffer_interop.cpp index 8d83fafe9f4a0..18f4fd003d9ac 100644 --- a/sycl/test/basic_tests/buffer/buffer_interop.cpp +++ b/sycl/test/basic_tests/buffer/buffer_interop.cpp @@ -4,10 +4,9 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out //==------------------- buffer.cpp - SYCL buffer basic test ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/context.cpp b/sycl/test/basic_tests/context.cpp index 899104d76edc3..2e71f9a45e299 100644 --- a/sycl/test/basic_tests/context.cpp +++ b/sycl/test/basic_tests/context.cpp @@ -3,10 +3,9 @@ //==--------------- context.cpp - SYCL context test ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/basic_tests/device.cpp b/sycl/test/basic_tests/device.cpp index 73759a595618a..802d305f46111 100644 --- a/sycl/test/basic_tests/device.cpp +++ b/sycl/test/basic_tests/device.cpp @@ -3,10 +3,9 @@ //==--------------- device.cpp - SYCL device test --------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/basic_tests/device_event.cpp b/sycl/test/basic_tests/device_event.cpp index 338f97bcdb0d9..cb82884935aa5 100644 --- a/sycl/test/basic_tests/device_event.cpp +++ b/sycl/test/basic_tests/device_event.cpp @@ -6,10 +6,9 @@ //==--------device_event.cpp - SYCL class device_event test ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/event.cpp b/sycl/test/basic_tests/event.cpp index 7ed59aca0891a..e8e7752a2573f 100644 --- a/sycl/test/basic_tests/event.cpp +++ b/sycl/test/basic_tests/event.cpp @@ -2,10 +2,9 @@ // RUN: env SYCL_DEVICE_TYPE=HOST %t.out //==--------------- event.cpp - SYCL event test ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/group.cpp b/sycl/test/basic_tests/group.cpp index bb90e84c47bed..9fdc254f10a4d 100644 --- a/sycl/test/basic_tests/group.cpp +++ b/sycl/test/basic_tests/group.cpp @@ -4,10 +4,9 @@ //==--------------- group.cpp - SYCL group test ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/handler/handler_mem_op.cpp b/sycl/test/basic_tests/handler/handler_mem_op.cpp index 3732ce0daad64..3ed7421d85219 100644 --- a/sycl/test/basic_tests/handler/handler_mem_op.cpp +++ b/sycl/test/basic_tests/handler/handler_mem_op.cpp @@ -4,10 +4,9 @@ // RUN: %GPU_RUN_PLACEHOLDER %t.out //==- handler.cpp - SYCL handler explicit memory operations test -*- C++-*--==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/basic_tests/id.cpp b/sycl/test/basic_tests/id.cpp index fe858cf2a0517..28ea042607e04 100644 --- a/sycl/test/basic_tests/id.cpp +++ b/sycl/test/basic_tests/id.cpp @@ -4,10 +4,9 @@ //==--------------- id.cpp - SYCL id test ----------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/info.cpp b/sycl/test/basic_tests/info.cpp index 7afd9fc4cc14a..823ac768943cf 100644 --- a/sycl/test/basic_tests/info.cpp +++ b/sycl/test/basic_tests/info.cpp @@ -5,10 +5,9 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out //==----------------info.cpp - SYCL objects get_info() test ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/item.cpp b/sycl/test/basic_tests/item.cpp index 82e9b21b0b771..205d7b484b07b 100644 --- a/sycl/test/basic_tests/item.cpp +++ b/sycl/test/basic_tests/item.cpp @@ -3,10 +3,9 @@ // RUN: %t.out //==--------------- item.cpp - SYCL item test ------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/macros.cpp b/sycl/test/basic_tests/macros.cpp new file mode 100644 index 0000000000000..db584857f0ce6 --- /dev/null +++ b/sycl/test/basic_tests/macros.cpp @@ -0,0 +1,19 @@ +// RUN: %clang -std=c++11 -g %s -o %t.out -lstdc++ -lOpenCL -lsycl +//==------------------- macros.cpp - SYCL buffer basic test ----------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +int main() { + std::cout << "SYCL language version: " << CL_SYCL_LANGUAGE_VERSION + << std::endl; + std::cout << "SYCL compiler version: " << __SYCL_COMPILER_VERSION + << std::endl; + return 0; +} diff --git a/sycl/test/basic_tests/nd_item.cpp b/sycl/test/basic_tests/nd_item.cpp index 836f8a272fae2..200524c0147d4 100644 --- a/sycl/test/basic_tests/nd_item.cpp +++ b/sycl/test/basic_tests/nd_item.cpp @@ -3,10 +3,9 @@ // RUN: %t.out //==--------------- nd_item.cpp - SYCL nd_item test ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/nd_range.cpp b/sycl/test/basic_tests/nd_range.cpp index 7b0563155e75d..1d9d75be096c7 100644 --- a/sycl/test/basic_tests/nd_range.cpp +++ b/sycl/test/basic_tests/nd_range.cpp @@ -3,10 +3,9 @@ // RUN: %t.out //==--------------- nd_range.cpp - SYCL nd_range test ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/platform.cpp b/sycl/test/basic_tests/platform.cpp index cda8b9625b2ae..e59c99f28878d 100644 --- a/sycl/test/basic_tests/platform.cpp +++ b/sycl/test/basic_tests/platform.cpp @@ -3,10 +3,9 @@ // RUN: %t.out //==--------------- platform.cpp - SYCL platform test ----------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/property_list.cpp b/sycl/test/basic_tests/property_list.cpp index c06023c12ae3f..d1faeb950ba60 100644 --- a/sycl/test/basic_tests/property_list.cpp +++ b/sycl/test/basic_tests/property_list.cpp @@ -4,10 +4,9 @@ // CHECK: PASSED //==--------------- property_list.cpp - SYCL property list test ------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/queue.cpp b/sycl/test/basic_tests/queue.cpp index 1be9bb49a36c0..55af303f862b3 100644 --- a/sycl/test/basic_tests/queue.cpp +++ b/sycl/test/basic_tests/queue.cpp @@ -3,10 +3,9 @@ // RUN: env SYCL_DEVICE_TYPE=HOST %t.out //==--------------- queue.cpp - SYCL queue test ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include @@ -88,10 +87,18 @@ int main() { queue Queue(pl); try { Queue.throw_asynchronous(); - } - catch (const std::bad_function_call& e) { - std::cout << "Default asynchronous handler call failed: " << e.what() << std::endl; + } catch (const std::bad_function_call &e) { + std::cout << "Default asynchronous handler call failed: " << e.what() + << std::endl; throw; } } + + { + default_selector Selector; + device Device = Selector.select_device(); + context Context(Device); + queue Queue(Context, Selector); + assert(Context == Queue.get_context()); + } } diff --git a/sycl/test/basic_tests/range.cpp b/sycl/test/basic_tests/range.cpp index 00793deefe3d0..d80d60bc32c3b 100644 --- a/sycl/test/basic_tests/range.cpp +++ b/sycl/test/basic_tests/range.cpp @@ -3,10 +3,9 @@ // RUN: %t.out //==--------------- range.cpp - SYCL range test ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/range_error.cpp b/sycl/test/basic_tests/range_error.cpp index 5e7cd271c5ae1..ec657ace220c0 100644 --- a/sycl/test/basic_tests/range_error.cpp +++ b/sycl/test/basic_tests/range_error.cpp @@ -1,10 +1,9 @@ // RUN: %clang -std=c++11 -Xclang -verify %s -Xclang -verify-ignore-unexpected=note,warning -fsyntax-only //==--------------- range_error.cpp - SYCL range error test ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/subdevice.cpp b/sycl/test/basic_tests/subdevice.cpp index 5a55ac9b9527a..a0a5a03eaf5b2 100644 --- a/sycl/test/basic_tests/subdevice.cpp +++ b/sycl/test/basic_tests/subdevice.cpp @@ -5,10 +5,9 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out //==------------ subdevice.cpp - SYCL subdevice basic test -----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/basic_tests/swizzle_op.cpp b/sycl/test/basic_tests/swizzle_op.cpp index 5af41e58a1730..1e92462272208 100644 --- a/sycl/test/basic_tests/swizzle_op.cpp +++ b/sycl/test/basic_tests/swizzle_op.cpp @@ -5,10 +5,9 @@ // RUNx: %ACC_RUN_PLACEHOLDER %t.out //==------------ swizzle_op.cpp - SYCL SwizzleOp basic test ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #define SYCL_SIMPLE_SWIZZLES @@ -228,4 +227,33 @@ int main() { assert(results[2] == 3); assert(results[3] == 4); } + + { + cl::sycl::cl_uint results[4] = {0}; + { + buffer b(results, range<1>(4)); + queue myQueue; + myQueue.submit([&](handler &cgh) { + auto B = b.get_access(cgh); + cgh.single_task([=]() { + cl::sycl::uchar4 vec; + cl::sycl::uint add = 254; + cl::sycl::uchar factor = 2; + vec.x() = 2; + vec.y() = 4; + vec.z() = 6; + vec.w() = 8; + + B[0] = add + vec.x() / factor; + B[1] = add + vec.y() / factor; + B[2] = add + vec.z() / factor; + B[3] = add + vec.w() / factor; + }); + }); + } + assert(results[0] == 255); + assert(results[1] == 256); + assert(results[2] == 257); + assert(results[3] == 258); + } } diff --git a/sycl/test/basic_tests/types.cpp b/sycl/test/basic_tests/types.cpp index cf6358f558cb6..069d737f7eea7 100644 --- a/sycl/test/basic_tests/types.cpp +++ b/sycl/test/basic_tests/types.cpp @@ -1,10 +1,9 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl //==--------------- types.cpp - SYCL types test ----------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include diff --git a/sycl/test/basic_tests/vectors.cpp b/sycl/test/basic_tests/vectors.cpp index 132c862d3afa1..46bc40cfec830 100644 --- a/sycl/test/basic_tests/vectors.cpp +++ b/sycl/test/basic_tests/vectors.cpp @@ -3,10 +3,9 @@ // RUN: %t.out //==--------------- vectors.cpp - SYCL vectors test ------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/fpga_tests/fpga_queue.cpp b/sycl/test/fpga_tests/fpga_queue.cpp new file mode 100644 index 0000000000000..5c8b86abb96bc --- /dev/null +++ b/sycl/test/fpga_tests/fpga_queue.cpp @@ -0,0 +1,154 @@ +// RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +//==------------- fpga_queue.cpp - SYCL FPGA queues test -------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +#include +#include + +using namespace cl::sycl; + +const int dataSize = 32; +const int maxNumQueues = 256; + +void GetCLQueue(event sycl_event, std::set& cl_queues) { + try { + cl_command_queue cl_queue; + cl_event cl_event = sycl_event.get(); + cl_int error = clGetEventInfo(cl_event, CL_EVENT_COMMAND_QUEUE, + sizeof(cl_queue), &cl_queue, nullptr); + assert(CL_SUCCESS == error && "Failed to obtain queue from OpenCL event"); + + cl_queues.insert(cl_queue); + } catch (invalid_object_error e) { + std::cout << "Failed to get OpenCL queue from SYCL event: " << e.what() + << std::endl; + } +} + +int main() { + int data[dataSize] = {0}; + + { + queue Queue; + std::set cl_queues; + event sycl_event; + + // Purpose of this test is to check how many OpenCL queues are being + // created from 1 SYCL queue for FPGA device. For that we submit 3 kernels + // expecting 3 OpenCL queues created as a result. + buffer bufA (data, range<1>(dataSize)); + buffer bufB (data, range<1>(dataSize)); + buffer bufC (data, range<1>(dataSize)); + + sycl_event = Queue.submit([&](handler& cgh) { + auto writeBuffer = bufA.get_access(cgh); + + // Create a range. + auto myRange = range<1>(dataSize); + + // Create a kernel. + auto myKernel = ([=](id<1> idx) { + writeBuffer[idx] = idx[0]; + }); + + cgh.parallel_for(myRange, myKernel); + }); + GetCLQueue(sycl_event, cl_queues); + + sycl_event = Queue.submit([&](handler& cgh) { + auto writeBuffer = bufB.get_access(cgh); + + // Create a range. + auto myRange = range<1>(dataSize); + + // Create a kernel. + auto myKernel = ([=](id<1> idx) { + writeBuffer[idx] = idx[0]; + }); + + cgh.parallel_for(myRange, myKernel); + }); + GetCLQueue(sycl_event, cl_queues); + + sycl_event = Queue.submit([&](handler& cgh) { + auto readBufferA = bufA.get_access(cgh); + auto readBufferB = bufB.get_access(cgh); + auto writeBuffer = bufC.get_access(cgh); + + // Create a range. + auto myRange = range<1>(dataSize); + + // Create a kernel. + auto myKernel = ([=](id<1> idx) { + writeBuffer[idx] = readBufferA[idx] + readBufferB[idx]; + }); + + cgh.parallel_for(myRange, myKernel); + }); + GetCLQueue(sycl_event, cl_queues); + + int result = cl_queues.size(); + device dev = Queue.get_device(); + int expected_result = dev.is_accelerator() ? 3 : dev.is_host() ? 0 : 1; + + if (expected_result != result) { + std::cout << "Result Num of queues = " << result << std::endl + << "Expected Num of queues = 3" << std::endl; + + return -1; + } + + auto readBufferC = bufC.get_access(); + for (size_t i = 0; i != dataSize; ++i) { + if (readBufferC[i] != 2 * i) { + std::cout << "Result mismatches " << readBufferC[i] << " Vs expected " + << 2 * i << " for index " << i << std::endl; + } + } + } + + { + queue Queue; + std::set cl_queues; + event sycl_event; + + // Check limits of OpenCL queues creation for accelerator device. + buffer buf (&data[0], range<1>(1)); + + for (size_t i = 0; i != maxNumQueues + 1; ++i) { + sycl_event = Queue.submit([&](handler& cgh) { + auto Buffer = buf.get_access(cgh); + + // Create a kernel. + auto myKernel = ([=]() { + Buffer[0] = 0; + }); + + cgh.single_task(myKernel); + }); + GetCLQueue(sycl_event, cl_queues); + } + + int result = cl_queues.size(); + device dev = Queue.get_device(); + int expected_result = dev.is_accelerator() ? maxNumQueues : + dev.is_host() ? 0 : 1; + + if (expected_result != result) { + std::cout << "Result Num of queues = " << result << std::endl + << "Expected Num of queues = " << maxNumQueues << std::endl; + + return -1; + } + } + + return 0; +} diff --git a/sycl/test/functor/kernel_functor.cpp b/sycl/test/functor/kernel_functor.cpp index b3fb07683d288..13680964326b9 100644 --- a/sycl/test/functor/kernel_functor.cpp +++ b/sycl/test/functor/kernel_functor.cpp @@ -7,10 +7,9 @@ //==--- kernel_functor.cpp - Functors as SYCL kernel test ------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/kernel-and-program/kernel-and-program.cpp b/sycl/test/kernel-and-program/kernel-and-program.cpp index 46f8cb3346208..cabb326e755d9 100644 --- a/sycl/test/kernel-and-program/kernel-and-program.cpp +++ b/sycl/test/kernel-and-program/kernel-and-program.cpp @@ -5,10 +5,9 @@ // RUNx: %ACC_RUN_PLACEHOLDER %t.out //==--- kernel-and-program.cpp - SYCL kernel/program test ------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/lit.cfg b/sycl/test/lit.cfg index 390ba1e2e32eb..6b8a469163cc9 100644 --- a/sycl/test/lit.cfg +++ b/sycl/test/lit.cfg @@ -39,6 +39,7 @@ if 'LD_LIBRARY_PATH' in os.environ: else: config.environment['LD_LIBRARY_PATH'] = config.llvm_build_libs_dir +config.substitutions.append( ('%clang_cc1', ' ' + config.clang + ' -cc1 ') ) config.substitutions.append( ('%clang', ' ' + config.clang + ' -I'+config.opencl_include ) ) config.substitutions.append( ('%llvm_build_libs_dir', config.llvm_build_libs_dir ) ) config.substitutions.append( ('%opencl_include', config.opencl_include ) ) diff --git a/sycl/test/math/math.cpp b/sycl/test/math/math.cpp index 9db4e4365b335..963fe8db62bb3 100644 --- a/sycl/test/math/math.cpp +++ b/sycl/test/math/math.cpp @@ -5,10 +5,9 @@ // RUN: %ACC_RUN_PLACEHOLDER %t.out //==--------------- math.cpp - SYCL math test ------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -51,14 +50,15 @@ int main() { resultAccessor[wiID.get(0)] += cl::sycl::native::exp(2.f); resultAccessor[wiID.get(0)] += cl::sycl::fabs(-2.f); resultAccessor[wiID.get(0)] += cl::sycl::fabs(1.0); + resultAccessor[wiID.get(0)] += cl::sycl::floor(-3.4); + resultAccessor[wiID.get(0)] += cl::sycl::ceil(2.4); }); }); } - for (size_t i = 0; i < 10; ++i) { - /* Result of addition of 2 + 1 + 7.389... + 2 + 1*/ - assert(result[i] > 13 && result[i] < 14 && - "Expected result[i] > 13 && result[i] < 14"); + /* Result of addition of 2 + 1 + 7.389... + 2 + 1 -4 + 3*/ + assert(result[i] > 12 && result[i] < 13 && + "Expected result[i] > 12 && result[i] < 13"); } return 0; diff --git a/sycl/test/multi_ptr/multi_ptr.cpp b/sycl/test/multi_ptr/multi_ptr.cpp index b1006b82db17d..e7f0e8d73f82c 100644 --- a/sycl/test/multi_ptr/multi_ptr.cpp +++ b/sycl/test/multi_ptr/multi_ptr.cpp @@ -6,10 +6,9 @@ //==--------------- multi_ptr.cpp - SYCL multi_ptr test --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/multisource/multisource.cpp b/sycl/test/multisource/multisource.cpp index 8889ceee8ba1b..4269e7064ab73 100644 --- a/sycl/test/multisource/multisource.cpp +++ b/sycl/test/multisource/multisource.cpp @@ -1,9 +1,8 @@ //==--------------- multisource.cpp ----------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/program_manager/program_manager.cpp b/sycl/test/program_manager/program_manager.cpp index 7d1877dc5e5d2..59681c0469b80 100644 --- a/sycl/test/program_manager/program_manager.cpp +++ b/sycl/test/program_manager/program_manager.cpp @@ -5,10 +5,9 @@ //==--- program_manager.cpp - SYCL program manager test --------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/regression/kernel_name_class.cpp b/sycl/test/regression/kernel_name_class.cpp index 40a2678d86384..4e4d2e835ef55 100644 --- a/sycl/test/regression/kernel_name_class.cpp +++ b/sycl/test/regression/kernel_name_class.cpp @@ -6,10 +6,9 @@ //==-- kernel_name_class.cpp - SYCL kernel naming variants test ------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/scheduler/Dump.cpp b/sycl/test/scheduler/Dump.cpp index 169f3a9b5bd22..f372aac545a32 100644 --- a/sycl/test/scheduler/Dump.cpp +++ b/sycl/test/scheduler/Dump.cpp @@ -1,9 +1,8 @@ //==--------------- Dump.cpp - Test SYCL scheduler graph dumping -----------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// // RUN: %clang -std=c++11 -g %s -o %t.out -lstdc++ -lOpenCL -lsycl diff --git a/sycl/test/scheduler/MultipleDevices.cpp b/sycl/test/scheduler/MultipleDevices.cpp index 23b9646dc01e9..d7efff9314ef1 100644 --- a/sycl/test/scheduler/MultipleDevices.cpp +++ b/sycl/test/scheduler/MultipleDevices.cpp @@ -2,10 +2,9 @@ // RUN: %t.out //===- MultipleDevices.cpp - Test checkking multi-device execution --------===// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/scheduler/parallelReadOpt.cpp b/sycl/test/scheduler/parallelReadOpt.cpp index 7f06ae0cf37da..026c54283ac3e 100644 --- a/sycl/test/scheduler/parallelReadOpt.cpp +++ b/sycl/test/scheduler/parallelReadOpt.cpp @@ -20,10 +20,9 @@ //==---- parallelReadOpt.cpp - SYCL scheduler parallel read test -----------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/separate-compile/test.cpp b/sycl/test/separate-compile/test.cpp index fb1d108e5cc91..7da291dc64599 100755 --- a/sycl/test/separate-compile/test.cpp +++ b/sycl/test/separate-compile/test.cpp @@ -1,30 +1,25 @@ // >> ---- compile src1 // >> device compilation... -// RUN: %clang -std=c++11 --sycl -Xclang -fsycl-int-header=sycl_ihdr_a.h %s -c -o a_kernel.spv +// RUN: %clang -std=c++11 --sycl -Xclang -fsycl-int-header=sycl_ihdr_a.h %s -c -o a_kernel.bc // >> host compilation... // RUN: %clang -std=c++11 -include sycl_ihdr_a.h -g -c %s -o a.o // // >> ---- compile src2 // >> device compilation... -// RUN: %clang -DB_CPP=1 -std=c++11 --sycl -Xclang -fsycl-int-header=sycl_ihdr_b.h %s -c -o b_kernel.spv +// RUN: %clang -DB_CPP=1 -std=c++11 --sycl -Xclang -fsycl-int-header=sycl_ihdr_b.h %s -c -o b_kernel.bc // >> host compilation... // RUN: %clang -DB_CPP=1 -std=c++11 -include sycl_ihdr_b.h -g -c %s -o b.o // // >> ---- bundle .o with .spv // >> run bundler -// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=a.o,a_kernel.spv -outputs=a_fat.o -// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=b.o,b_kernel.spv -outputs=b_fat.o +// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=a.o,a_kernel.bc -outputs=a_fat.o +// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -inputs=b.o,b_kernel.bc -outputs=b_fat.o // // >> ---- unbundle fat objects -// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=a.o,a_kernel.spv -inputs=a_fat.o -unbundle -// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=b.o,b_kernel.spv -inputs=b_fat.o -unbundle +// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=a.o,a_kernel.bc -inputs=a_fat.o -unbundle +// RUN: clang-offload-bundler -type=o -targets=host-x86_64,sycl-spir64-pc-linux-gnu -outputs=b.o,b_kernel.bc -inputs=b_fat.o -unbundle // // >> ---- link device code -// >> convert to bitcode -// RUN: llvm-spirv -r -o=a_kernel.bc a_kernel.spv -// RUN: llvm-spirv -r -o=b_kernel.bc b_kernel.spv -// -// >> link bitcode // RUN: llvm-link -o=app.bc a_kernel.bc b_kernel.bc // // >> convert linked .bc to spirv @@ -44,10 +39,9 @@ //==----------- test.cpp - Tests SYCL separate compilation -----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #ifdef B_CPP diff --git a/sycl/test/struct_param/struct_kernel_param.cpp b/sycl/test/struct_param/struct_kernel_param.cpp index 5438aa7b5eefe..7ff7cb005f83c 100644 --- a/sycl/test/struct_param/struct_kernel_param.cpp +++ b/sycl/test/struct_param/struct_kernel_param.cpp @@ -6,10 +6,9 @@ //==-struct_kernel_param.cpp-Checks passing structs as kernel params--------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/sub_group/barrier.cpp b/sycl/test/sub_group/barrier.cpp index dc61f9bd4c1f6..3f7d8bd855871 100644 --- a/sycl/test/sub_group/barrier.cpp +++ b/sycl/test/sub_group/barrier.cpp @@ -1,15 +1,13 @@ // RUN: %clang -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==---------- barrier.cpp - SYCL sub_group barrier test -------------------==// +//==---------- barrier.cpp - SYCL sub_group barrier test -------*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -47,8 +45,8 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { }); }); auto addacc = addbuf.template get_access(); - auto sgsizeacc = sgsizebuf.get_access(); + size_t sg_size = sgsizeacc[0]; int WGid = -1, SGid = 0; T add = 0; @@ -75,18 +73,12 @@ int main() { std::cout << "Skipping test\n"; return 0; } - /* Limit work-group size to avoid type overflow. */ - check(Queue, 120, 30); - check(Queue, 1024, 256); check(Queue); - check(Queue); + check(Queue); check(Queue); - check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp16")) { - check(Queue); - } + check(Queue); check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp64")) { + if (Queue.get_device().has_extension("cl_khr_fp64")) { check(Queue); } std::cout << "Test passed." << std::endl; diff --git a/sycl/test/sub_group/broadcast.cpp b/sycl/test/sub_group/broadcast.cpp index 70561e4c2d024..f6f8465bb25dc 100644 --- a/sycl/test/sub_group/broadcast.cpp +++ b/sycl/test/sub_group/broadcast.cpp @@ -1,15 +1,13 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==--------- broadcast.cpp - SYCL sub_group broadcast test ----------------==// +//==--------- broadcast.cpp - SYCL sub_group broadcast test ----*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -28,16 +26,18 @@ template void check(queue &Queue) { auto sgsizeacc = sgsizebuf.get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { intel::sub_group SG = NdItem.get_sub_group(); - if (NdItem.get_global_id(0) == 0) - sgsizeacc[0] = SG.get_max_local_range()[0]; /*Broadcast GID of element with SGLID == SGID */ syclacc[NdItem.get_global_id()] = SG.broadcast(NdItem.get_global_id(0), SG.get_group_id()); + if (NdItem.get_global_id(0) == 0) + sgsizeacc[0] = SG.get_max_local_range()[0]; }); }); auto syclacc = syclbuf.template get_access(); auto sgsizeacc = sgsizebuf.get_access(); size_t sg_size = sgsizeacc[0]; + if (sg_size == 0) + sg_size = L; int WGid = -1, SGid = 0; for (int j = 0; j < G; j++) { if (j % L % sg_size == 0) { @@ -61,17 +61,12 @@ int main() { std::cout << "Skipping test\n"; return 0; } - check(Queue); - check(Queue); check(Queue); - check(Queue); + check(Queue); check(Queue); - check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp16")) { - check(Queue); - } + check(Queue); check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp64")) { + if (Queue.get_device().has_extension("cl_khr_fp64")) { check(Queue); } std::cout << "Test passed." << std::endl; diff --git a/sycl/test/sub_group/common.cpp b/sycl/test/sub_group/common.cpp index f63795b7879e2..2dbc08a6fcd16 100644 --- a/sycl/test/sub_group/common.cpp +++ b/sycl/test/sub_group/common.cpp @@ -1,31 +1,29 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==-------------- common.cpp - SYCL sub_group common test -----------------==// +//==-------------- common.cpp - SYCL sub_group common test -----*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// + #include "helper.hpp" #include - using namespace cl::sycl; struct Data { - size_t local_id; - size_t local_range; - size_t max_local_range; - size_t group_id; - size_t group_range; - size_t uniform_group_range; + unsigned int local_id; + unsigned int local_range; + unsigned int max_local_range; + unsigned int group_id; + unsigned int group_range; + unsigned int uniform_group_range; }; -void check(queue &Queue, const int G, const int L) { +void check(queue &Queue, unsigned int G, unsigned int L) { try { nd_range<1> NdRange(G, L); @@ -47,16 +45,24 @@ void check(queue &Queue, const int G, const int L) { }); }); auto syclacc = syclbuf.get_access(); - size_t max_sg = - Queue.get_device().get_info(); - size_t num_sg = L / max_sg + (L % max_sg ? 1 : 0); + unsigned int max_sg = get_sg_size(Queue.get_device()); + unsigned int num_sg = L / max_sg + (L % max_sg ? 1 : 0); for (int j = 0; j < G; j++) { - size_t group_id = j % L / max_sg; - size_t local_range = + unsigned int group_id = j % L / max_sg; + unsigned int local_range = (group_id + 1 == num_sg) ? (L - group_id * max_sg) : max_sg; exit_if_not_equal(syclacc[j].local_id, j % L % max_sg, "local_id"); exit_if_not_equal(syclacc[j].local_range, local_range, "local_range"); - exit_if_not_equal(syclacc[j].max_local_range, max_sg, "max_local_range"); + // TODO: Currently workgroup size affects this paramater on CPU and does + // not on GPU. Remove if when it is aligned. + if (Queue.get_device().get_info() == + info::device_type::cpu) { + exit_if_not_equal(syclacc[j].max_local_range, std::min(max_sg, L), + "max_local_range"); + } else { + exit_if_not_equal(syclacc[j].max_local_range, max_sg, + "max_local_range"); + } exit_if_not_equal(syclacc[j].group_id, group_id, "group_id"); exit_if_not_equal(syclacc[j].group_range, num_sg, "group_range"); exit_if_not_equal(syclacc[j].uniform_group_range, num_sg, diff --git a/sycl/test/sub_group/common_ocl.cpp b/sycl/test/sub_group/common_ocl.cpp index eba8a48d3a7ce..194a9f3bb3e48 100644 --- a/sycl/test/sub_group/common_ocl.cpp +++ b/sycl/test/sub_group/common_ocl.cpp @@ -1,51 +1,54 @@ +// RUN: %clang_cc1 -x cl -cl-std=CL2.0 %S/sg.cl -triple spir64-unknown-unknown -emit-spirv -o %T/kernel_ocl.spv -include opencl-c.h // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out -// RUN: %ACC_RUN_PLACEHOLDER %t.out -//==--- common_ocl.cpp - basic SG methods in SYCL vs OpenCL ---------------==// +// RUN: %CPU_RUN_PLACEHOLDER %t.out %T/kernel_ocl.spv +// RUN: %GPU_RUN_PLACEHOLDER %t.out %T/kernel_ocl.spv +// RUN: %ACC_RUN_PLACEHOLDER %t.out %T/kernel_ocl.spv +//==--- common_ocl.cpp - basic SG methods in SYCL vs OpenCL ---*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include "helper.hpp" #include - +#include +#include using namespace cl::sycl; struct Data { - size_t local_id; - size_t local_range; - size_t max_local_range; - size_t group_id; - size_t group_range; - size_t uniform_group_range; + unsigned int local_id; + unsigned int local_range; + unsigned int max_local_range; + unsigned int group_id; + unsigned int group_range; + unsigned int uniform_group_range; }; -void check(queue &Queue, const int G, const int L) { - +void check(queue &Queue, const int G, const int L, const char *SpvFile) { try { nd_range<1> NdRange(G, L); buffer oclbuf(G); buffer syclbuf(G); - program Prog(Queue.get_context()); - Prog.build_with_source( - "struct Data { size_t local_id; size_t local_range; size_t " - "max_local_range; size_t group_id; size_t group_range; \n" - "size_t uniform_group_range; };\n" - "kernel void ocl_subgr(global struct Data* a) {\n" - "size_t id = get_global_id(0);" - "a[id].local_id = get_sub_group_local_id();\n" - "a[id].local_range = get_sub_group_size();\n" - "a[id].max_local_range = get_max_sub_group_size();\n" - "a[id].group_id = get_sub_group_id();\n" - "a[id].group_range = get_num_sub_groups();\n" - "a[id].uniform_group_range = get_num_sub_groups(); }"); + std::ifstream File(SpvFile, std::ios::binary); + if (!File.is_open()) { + std::cerr << std::strerror(errno); + throw compile_program_error("Cannot open SPIRV file\n"); + } + File.seekg(0, std::ios::end); + vector_class Spv(File.tellg()); + File.seekg(0); + File.read(Spv.data(), Spv.size()); + File.close(); + int Err; + cl_program ClProgram = clCreateProgramWithIL(Queue.get_context().get(), + Spv.data(), Spv.size(), &Err); + CHECK_OCL_CODE(Err); + CHECK_OCL_CODE( + clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr)); + program Prog(Queue.get_context(), ClProgram); Queue.submit([&](handler &cgh) { auto oclacc = oclbuf.get_access(cgh); cgh.set_args(oclacc); @@ -90,17 +93,17 @@ void check(queue &Queue, const int G, const int L) { exit(1); } } -int main() { +int main(int argc, char **argv) { queue Queue; - if (!core_sg_supported(Queue.get_device())) { + if (!core_sg_supported(Queue.get_device()) || argc != 2) { std::cout << "Skipping test\n"; return 0; } - check(Queue, 240, 80); - check(Queue, 8, 4); - check(Queue, 24, 12); - check(Queue, 1024, 256); + check(Queue, 240, 80, argv[1]); + check(Queue, 8, 4, argv[1]); + check(Queue, 24, 12, argv[1]); + check(Queue, 1024, 256, argv[1]); std::cout << "Test passed." << std::endl; return 0; } diff --git a/sycl/test/sub_group/helper.hpp b/sycl/test/sub_group/helper.hpp index 1a786c5b639e9..d7ab271208b0b 100644 --- a/sycl/test/sub_group/helper.hpp +++ b/sycl/test/sub_group/helper.hpp @@ -1,23 +1,128 @@ //==---------- helper.hpp - SYCL sub_group helper functions ----------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include #include using namespace cl::sycl; + +template struct utils { + static T1 add_vec(const vec &v); + static bool cmp_vec(const vec &v, const vec &r); + static std::string stringify_vec(const vec &v); +}; +template struct utils { + static T2 add_vec(const vec &v) { return v.s0(); } + static bool cmp_vec(const vec &v, const vec &r) { + return v.s0() == r.s0(); + } + static std::string stringify_vec(const vec &v) { + return std::to_string((T2)v.s0()); + } +}; +template struct utils { + static T2 add_vec(const vec &v) { return v.s0() + v.s1(); } + static bool cmp_vec(const vec &v, const vec &r) { + return v.s0() == r.s0() && v.s1() == r.s1(); + } + static std::string stringify_vec(const vec &v) { + return std::string("(") + std::to_string((T2)v.s0()) + ", " + + std::to_string((T2)v.s1()) + " )"; + } +}; +template struct utils { + static T2 add_vec(const vec &v) { + return v.s0() + v.s1() + v.s2() + v.s3(); + } + static bool cmp_vec(const vec &v, const vec &r) { + return v.s0() == r.s0() && v.s1() == r.s1() && v.s2() == r.s2() && + v.s3() == r.s3(); + } + static std::string stringify_vec(const vec &v) { + return std::string("(") + std::to_string((T2)v.s0()) + ", " + + std::to_string((T2)v.s1()) + std::to_string((T2)v.s2()) + ", " + + std::to_string((T2)v.s3()) + " )"; + } +}; +template struct utils { + static T2 add_vec(const vec &v) { + return v.s0() + v.s1() + v.s2() + v.s3() + v.s4() + v.s5() + v.s6() + + v.s7(); + } + static bool cmp_vec(const vec &v, const vec &r) { + return v.s0() == r.s0() && v.s1() == r.s1() && v.s2() == r.s2() && + v.s3() == r.s3() && v.s4() == r.s4() && v.s5() == r.s5() && + v.s6() == r.s6() && v.s7() == r.s7(); + } + static std::string stringify_vec(const vec &v) { + return std::string("(") + std::to_string((T2)v.s0()) + ", " + + std::to_string((T2)v.s1()) + std::to_string((T2)v.s2()) + ", " + + std::to_string((T2)v.s3()) + std::to_string((T2)v.s4()) + ", " + + std::to_string((T2)v.s5()) + std::to_string((T2)v.s6()) + ", " + + std::to_string((T2)v.s7()) + " )"; + } +}; + +template struct utils { + static T2 add_vec(const vec &v) { + return v.s0() + v.s1() + v.s2() + v.s3() + v.s4() + v.s5() + v.s6() + + v.s7() + v.s8() + v.s9() + v.sA() + v.sB() + v.sC() + v.sD() + + v.sE() + v.sF(); + } + static bool cmp_vec(const vec &v, const vec &r) { + return v.s0() == r.s0() && v.s1() == r.s1() && v.s2() == r.s2() && + v.s3() == r.s3() && v.s4() == r.s4() && v.s5() == r.s5() && + v.s6() == r.s6() && v.s7() == r.s7() && v.s8() == r.s8() && + v.s9() == r.s9() && v.sA() == r.sA() && v.sB() == r.sB() && + v.sC() == r.sC() && v.sD() == r.sD() && v.sE() == r.sE() && + v.sF() == r.sF(); + } + static std::string stringify_vec(const vec &v) { + return std::string("(") + std::to_string((T2)v.s0()) + ", " + + std::to_string((T2)v.s1()) + std::to_string((T2)v.s2()) + ", " + + std::to_string((T2)v.s3()) + std::to_string((T2)v.s4()) + ", " + + std::to_string((T2)v.s5()) + std::to_string((T2)v.s6()) + ", " + + std::to_string((T2)v.s7()) + std::to_string((T2)v.s8()) + ", " + + std::to_string((T2)v.s9()) + std::to_string((T2)v.sA()) + ", " + + std::to_string((T2)v.sB()) + std::to_string((T2)v.sC()) + ", " + + std::to_string((T2)v.sE()) + std::to_string((T2)v.sD()) + ", " + + std::to_string((T2)v.sF()) + " )"; + } +}; + template void exit_if_not_equal(T val, T ref, const char *name) { - if (std::fabs(val - ref) > 0.01) { - std::cout << "Unexpected result for " << name << ": " << val - << " expected value: " << ref << std::endl; - exit(1); + if (std::is_floating_point::value) { + if (std::fabs(val - ref) > 0.01) { + std::cout << "Unexpected result for " << name << ": " << (double)val + << " expected value: " << (double)ref << std::endl; + exit(1); + } + } else { + if (std::abs(val - ref) != 0) { + std::cout << "Unexpected result for " << name << ": " << (long)val + << " expected value: " << (long)ref << std::endl; + exit(1); + } } } -/* CPU returns max number of SG, GPU returns mux SG size for + +template +void exit_if_not_equal_vec(vec val, vec ref, const char *name) { + if (!utils::cmp_vec(ref, val)) { + std::cout << "Unexpected result for " << name << ": " + << utils::stringify_vec(val) + << " expected value: " << utils::stringify_vec(ref) + << std::endl; + + exit(0); + } +} + +/* CPU returns max number of SG, GPU returns max SG size for * CL_DEVICE_MAX_NUM_SUB_GROUPS device parameter. This function aligns the * value. * */ diff --git a/sycl/test/sub_group/info.cpp b/sycl/test/sub_group/info.cpp index ccbb97f07ea13..e0526c42d980e 100644 --- a/sycl/test/sub_group/info.cpp +++ b/sycl/test/sub_group/info.cpp @@ -3,12 +3,11 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==------------- info.cpp - SYCL sub_group parameters test ----------------==// +//==------------- info.cpp - SYCL sub_group parameters test ----*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/test/sub_group/load_store.cpp b/sycl/test/sub_group/load_store.cpp index 621cc71185b7d..19cb8ef10183c 100644 --- a/sycl/test/sub_group/load_store.cpp +++ b/sycl/test/sub_group/load_store.cpp @@ -1,15 +1,13 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out //==----------- load_store.cpp - SYCL sub_group load/store test ------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -18,29 +16,9 @@ template class sycl_subgr; using namespace cl::sycl; -template struct utils { - static T1 add_vec(const vec &v); -}; -template struct utils { - static T2 add_vec(const vec &v) { return v.s0(); } -}; -template struct utils { - static T2 add_vec(const vec &v) { return v.s0() + v.s1(); } -}; -template struct utils { - static T2 add_vec(const vec &v) { - return v.s0() + v.s1() + v.s2() + v.s3(); - } -}; -template struct utils { - static T2 add_vec(const vec &v) { - return v.s0() + v.s1() + v.s2() + v.s3() + v.s4() + v.s5() + v.s6() + - v.s7(); - } -}; template void check(queue &Queue) { - const int G = 128, L = 128; + const int G = 1024, L = 64; try { nd_range<1> NdRange(G, L); buffer syclbuf(G); @@ -58,9 +36,6 @@ template void check(queue &Queue) { cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { intel::sub_group SG = NdItem.get_sub_group(); if (SG.get_group_id().get(0) % N == 0) { - if (NdItem.get_global_id(0) == 0) - sgsizeacc[0] = SG.get_max_local_range()[0]; - size_t WGSGoffset = NdItem.get_group(0) * L + SG.get_group_id().get(0) * SG.get_max_local_range().get(0); @@ -70,13 +45,15 @@ template void check(queue &Queue) { vec v(utils::add_vec(SG.load(mp))); SG.store(mp, v); } + if (NdItem.get_global_id(0) == 0) + sgsizeacc[0] = SG.get_max_local_range()[0]; }); }); auto acc = syclbuf.template get_access(); auto sgsizeacc = sgsizebuf.get_access(); size_t sg_size = sgsizeacc[0]; int WGid = -1, SGid = 0; - for (int j = 0; j < G; j++) { + for (int j = 0; j < (G - (sg_size * N)); j++) { if (j % L % sg_size == 0) { SGid++; } @@ -92,11 +69,14 @@ template void check(queue &Queue) { ref += (T)(j + i * sg_size) + 0.1; } } - std::string s("Vector<"); - s += std::string(typeid(ref).name()) + std::string(",") + - std::to_string(N) + std::string(">[") + std::to_string(j) + - std::string("]"); - exit_if_not_equal(acc[j], ref, s.c_str()); + /* There is no defined out-of-range behavior for these functions. */ + if ((SGid + N) * sg_size < L) { + std::string s("Vector<"); + s += std::string(typeid(ref).name()) + std::string(",") + + std::to_string(N) + std::string(">[") + std::to_string(j) + + std::string("]"); + exit_if_not_equal(acc[j], ref, s.c_str()); + } } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); @@ -171,7 +151,7 @@ int main() { check(Queue); check(Queue); check(Queue); - typedef uint aligned_uint __attribute__((aligned(16))); + typedef unsigned int aligned_uint __attribute__((aligned(16))); check(Queue); check(Queue); check(Queue); @@ -184,20 +164,21 @@ int main() { check(Queue); check(Queue); } - if (Queue.get_device().has_extension("cl_khr_fp16") && - Queue.get_device().has_extension("cl_intel_subgroups_short")) { + if (Queue.get_device().has_extension("cl_intel_subgroups_short")) { typedef short aligned_short __attribute__((aligned(16))); check(Queue); check(Queue); check(Queue); check(Queue); check(Queue); - typedef half aligned_half __attribute__((aligned(16))); - check(Queue); - check(Queue); - check(Queue); - check(Queue); - check(Queue); + if (Queue.get_device().has_extension("cl_khr_fp16")) { + typedef half aligned_half __attribute__((aligned(16))); + check(Queue); + check(Queue); + check(Queue); + check(Queue); + check(Queue); + } } std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test/sub_group/reduce.cpp b/sycl/test/sub_group/reduce.cpp index d6b3e31781e71..5d78586910d99 100644 --- a/sycl/test/sub_group/reduce.cpp +++ b/sycl/test/sub_group/reduce.cpp @@ -1,15 +1,13 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==--------------- reduce.cpp - SYCL sub_group reduce test ----------------==// +//==--------------- reduce.cpp - SYCL sub_group reduce test ----*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -23,17 +21,12 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { buffer minbuf(G); buffer maxbuf(G); buffer addbuf(G); - buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { auto minacc = minbuf.template get_access(cgh); auto maxacc = maxbuf.template get_access(cgh); auto addacc = addbuf.template get_access(cgh); - auto sgsizeacc = - sgsizebuf.template get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { intel::sub_group sg = NdItem.get_sub_group(); - if (NdItem.get_global_id(0) == 0) - sgsizeacc[0] = sg.get_max_local_range()[0]; minacc[NdItem.get_global_id()] = sg.reduce(NdItem.get_global_id(0)); maxacc[NdItem.get_global_id()] = @@ -45,8 +38,7 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { auto minacc = minbuf.template get_access(); auto maxacc = maxbuf.template get_access(); auto addacc = addbuf.template get_access(); - auto sgsizeacc = sgsizebuf.template get_access(); - size_t sg_size = sgsizeacc[0]; + size_t sg_size = get_sg_size(Queue.get_device()); int WGid = -1, SGid = 0; T max = 0, add = 0; for (int j = 0; j < G; j++) { @@ -78,18 +70,12 @@ int main() { std::cout << "Skipping test\n"; return 0; } - /* Limit work-group size to avoid type overflow. */ - check(Queue, 120, 30); - check(Queue); check(Queue); - check(Queue); + check(Queue); check(Queue); - check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp16")) { - check(Queue); - } + check(Queue); check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp64")) { + if (Queue.get_device().has_extension("cl_khr_fp64")) { check(Queue); } std::cout << "Test passed." << std::endl; diff --git a/sycl/test/sub_group/scan.cpp b/sycl/test/sub_group/scan.cpp index a9eb942a67f17..683733d7ab271 100644 --- a/sycl/test/sub_group/scan.cpp +++ b/sycl/test/sub_group/scan.cpp @@ -1,15 +1,13 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==--------------- scan.cpp - SYCL sub_group scan test --------------------==// +//==--------------- scan.cpp - SYCL sub_group scan test --------*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// @@ -27,7 +25,6 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { buffer mininbuf(G); buffer maxinbuf(G); buffer addinbuf(G); - buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { auto minexacc = minexbuf.template get_access(cgh); @@ -41,11 +38,8 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { maxinbuf.template get_access(cgh); auto addinacc = addinbuf.template get_access(cgh); - auto sgsizeacc = sgsizebuf.get_access(cgh); cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { intel::sub_group SG = NdItem.get_sub_group(); - if (NdItem.get_global_id(0) == 0) - sgsizeacc[0] = SG.get_max_local_range()[0]; minexacc[NdItem.get_global_id()] = SG.exclusive_scan(NdItem.get_global_id(0)); maxexacc[NdItem.get_global_id()] = @@ -67,8 +61,7 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { auto maxinacc = maxinbuf.template get_access(); auto addinacc = addinbuf.template get_access(); - auto sgsizeacc = sgsizebuf.get_access(); - size_t sg_size = sgsizeacc[0]; + size_t sg_size = get_sg_size(Queue.get_device()); int WGid = -1, SGid = 0; T add = 0; for (int j = 0; j < G; j++) { @@ -104,18 +97,12 @@ int main() { std::cout << "Skipping test\n"; return 0; } - /* Limit work-group size to avoid type overflow. */ - check(Queue, 120, 30); - check(Queue); check(Queue); - check(Queue); + check(Queue); check(Queue); - check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp16")) { - check(Queue); - } + check(Queue); check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp64")) { + if (Queue.get_device().has_extension("cl_khr_fp64")) { check(Queue); } std::cout << "Test passed." << std::endl; diff --git a/sycl/test/sub_group/sg.cl b/sycl/test/sub_group/sg.cl new file mode 100644 index 0000000000000..0dcee4129807e --- /dev/null +++ b/sycl/test/sub_group/sg.cl @@ -0,0 +1,25 @@ +//==-------------- sg.cl - OpenCL reference kernel file --------*- C++ -*---==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// ===--------------------------------------------------------------------=== // +#pragma OPENCL EXTENSION cl_khr_subgroups : enable +struct Data { + uint local_id; + uint local_range; + uint max_local_range; + uint group_id; + uint group_range; + uint uniform_group_range; +}; +__kernel void ocl_subgr(__global struct Data *a) { + uint id = get_global_id(0); + a[id].local_id = get_sub_group_local_id(); + a[id].local_range = get_sub_group_size(); + a[id].max_local_range = get_max_sub_group_size(); + a[id].group_id = get_sub_group_id(); + a[id].group_range = get_num_sub_groups(); + a[id].uniform_group_range = get_num_sub_groups(); +} diff --git a/sycl/test/sub_group/shuffle.cpp b/sycl/test/sub_group/shuffle.cpp index 0b0db7ca073aa..7f246e2161d8c 100644 --- a/sycl/test/sub_group/shuffle.cpp +++ b/sycl/test/sub_group/shuffle.cpp @@ -1,22 +1,128 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out -// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUNx: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==------------ shuffle.cpp - SYCL sub_group shuffle test ----------------==// +//==------------ shuffle.cpp - SYCL sub_group shuffle test -----*- C++ -*---==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #include "helper.hpp" #include -template class sycl_subgr; +template class sycl_subgr; + using namespace cl::sycl; +template +void check(queue &Queue, size_t G = 240, size_t L = 60) { + try { + nd_range<1> NdRange(G, L); + buffer> buf2(G); + buffer> buf2_up(G); + buffer> buf2_down(G); + buffer> buf(G); + buffer> buf_up(G); + buffer> buf_down(G); + buffer> buf_xor(G); + buffer sgsizebuf(1); + Queue.submit([&](handler &cgh) { + auto acc2 = buf2.template get_access(cgh); + auto acc2_up = buf2_up.template get_access(cgh); + auto acc2_down = + buf2_down.template get_access(cgh); + + auto acc = buf.template get_access(cgh); + auto acc_up = buf_up.template get_access(cgh); + auto acc_down = + buf_down.template get_access(cgh); + auto acc_xor = buf_xor.template get_access(cgh); + auto sgsizeacc = sgsizebuf.get_access(cgh); + + cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { + intel::sub_group SG = NdItem.get_sub_group(); + uint32_t wggid = NdItem.get_global_id(0); + uint32_t sgid = SG.get_group_id().get(0); + vec vwggid(wggid), vsgid(sgid); + if (wggid == 0) + sgsizeacc[0] = SG.get_max_local_range()[0]; + /* 1 for odd subgroups and 2 for even*/ + acc2[NdItem.get_global_id()] = + SG.shuffle(vec(1), vec(2), + (sgid % 2) ? 1 : SG.get_max_local_range()[0]); + /* GID-SGID */ + acc2_up[NdItem.get_global_id()] = SG.shuffle_up(vwggid, vwggid, sgid); + /* GID-SGID or SGLID if GID+SGID > SGsize*/ + acc2_down[NdItem.get_global_id()] = + SG.shuffle_down(vwggid, vec(SG.get_local_id().get(0)), sgid); + + /*GID of middle element in every subgroup*/ + acc[NdItem.get_global_id()] = + SG.shuffle(vwggid, SG.get_max_local_range()[0] / 2); + /* Save GID-SGID */ + acc_up[NdItem.get_global_id()] = SG.shuffle_up(vwggid, sgid); + /* Save GID+SGID */ + acc_down[NdItem.get_global_id()] = SG.shuffle_down(vwggid, sgid); + /* Save GID XOR SGID */ + acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(vwggid, sgid); + }); + }); + auto acc = buf.template get_access(); + auto acc_up = buf_up.template get_access(); + auto acc_down = buf_down.template get_access(); + auto acc2 = buf2.template get_access(); + auto acc2_up = buf2_up.template get_access(); + auto acc2_down = buf2_down.template get_access(); + auto acc_xor = buf_xor.template get_access(); + auto sgsizeacc = sgsizebuf.get_access(); + + size_t sg_size = sgsizeacc[0]; + int SGid = 0; + for (int j = 0; j < G; j++) { + if (j % L % sg_size == 0) { + SGid++; + } + if (j % L == 0) { + SGid = 0; + } + /*GID of middle element in every subgroup*/ + exit_if_not_equal_vec( + acc[j], vec(j / L * L + SGid * sg_size + sg_size / 2), + "shuffle"); + /* 1 for odd subgroups and 2 for even*/ + exit_if_not_equal_vec(acc2[j], vec((SGid % 2) ? 1 : 2), + "shuffle2"); + /* Value GID+SGID for all element except last SGID in SG*/ + if (j % L % sg_size + SGid < sg_size && j % L + SGid < L) { + exit_if_not_equal_vec(acc_down[j], vec(j + SGid), "shuffle_down"); + exit_if_not_equal_vec(acc2_down[j], vec(j + SGid), + "shuffle2_down"); + } else { /* SGLID for GID+SGid */ + if (j % L + SGid < L) /* Do not go out LG*/ + exit_if_not_equal_vec(acc2_down[j], + vec((j + SGid) % L % sg_size), + "shuffle2_down"); + } + /* Value GID-SGID for all element except first SGID in SG*/ + if (j % L % sg_size >= SGid) { + exit_if_not_equal_vec(acc_up[j], vec(j - SGid), "shuffle_up"); + exit_if_not_equal_vec(acc2_up[j], vec(j - SGid), "shuffle2_up"); + } else { /* SGLID for GID-SGid */ + if (j % L - SGid + sg_size < L) /* Do not go out LG*/ + exit_if_not_equal_vec(acc2_up[j], vec(j - SGid + sg_size), + "shuffle2_up"); + } + /* GID XOR SGID */ + exit_if_not_equal_vec(acc_xor[j], vec(j ^ SGid), "shuffle_xor"); + } + } catch (exception e) { + std::cout << "SYCL exception caught: " << e.what(); + exit(1); + } +} + template void check(queue &Queue, size_t G = 240, size_t L = 60) { try { nd_range<1> NdRange(G, L); @@ -39,37 +145,32 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { auto acc_down = buf_down.template get_access(cgh); auto acc_xor = buf_xor.template get_access(cgh); - auto sgsizeacc = sgsizebuf.get_access(cgh); - cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { + + cgh.parallel_for>(NdRange, [=](nd_item<1> NdItem) { intel::sub_group SG = NdItem.get_sub_group(); - if (NdItem.get_global_id(0) == 0) + uint32_t wggid = NdItem.get_global_id(0); + uint32_t sgid = SG.get_group_id().get(0); + if (wggid == 0) sgsizeacc[0] = SG.get_max_local_range()[0]; /* 1 for odd subgroups and 2 for even*/ - acc2[NdItem.get_global_id()] = SG.shuffle( - 1, 2, - (SG.get_group_id().get(0) % 2) ? 1 : SG.get_max_local_range()[0]); + acc2[NdItem.get_global_id()] = + SG.shuffle(1, 2, (sgid % 2) ? 1 : SG.get_max_local_range()[0]); /* GID-SGID */ - acc2_up[NdItem.get_global_id()] = - SG.shuffle_up(NdItem.get_global_id(0), NdItem.get_global_id(0), - SG.get_group_id().get(0)); + acc2_up[NdItem.get_global_id()] = SG.shuffle_up(wggid, wggid, sgid); /* GID-SGID or SGLID if GID+SGID > SGsize*/ - acc2_down[NdItem.get_global_id()] = SG.shuffle_down( - NdItem.get_global_id(0), SG.get_local_id().get(0), - SG.get_group_id().get(0)); + acc2_down[NdItem.get_global_id()] = + SG.shuffle_down(wggid, SG.get_local_id().get(0), sgid); /*GID of middle element in every subgroup*/ - acc[NdItem.get_global_id()] = SG.shuffle( - NdItem.get_global_id(0), SG.get_max_local_range()[0] / 2); + acc[NdItem.get_global_id()] = + SG.shuffle(wggid, SG.get_max_local_range()[0] / 2); /* Save GID-SGID */ - acc_up[NdItem.get_global_id()] = - SG.shuffle_up(NdItem.get_global_id(0), SG.get_group_id().get(0)); + acc_up[NdItem.get_global_id()] = SG.shuffle_up(wggid, sgid); /* Save GID+SGID */ - acc_down[NdItem.get_global_id()] = SG.shuffle_down( - NdItem.get_global_id(0), SG.get_group_id().get(0)); + acc_down[NdItem.get_global_id()] = SG.shuffle_down(wggid, sgid); /* Save GID XOR SGID */ - acc_xor[NdItem.get_global_id()] = - SG.shuffle_xor(NdItem.get_global_id(0), SG.get_group_id()); + acc_xor[NdItem.get_global_id()] = SG.shuffle_xor(wggid, sgid); }); }); auto acc = buf.template get_access(); @@ -79,8 +180,8 @@ template void check(queue &Queue, size_t G = 240, size_t L = 60) { auto acc2_up = buf2_up.template get_access(); auto acc2_down = buf2_down.template get_access(); auto acc_xor = buf_xor.template get_access(); - auto sgsizeacc = sgsizebuf.get_access(); + size_t sg_size = sgsizeacc[0]; int SGid = 0; for (int j = 0; j < G; j++) { @@ -126,17 +227,28 @@ int main() { std::cout << "Skipping test\n"; return 0; } - check(Queue); - check(Queue); + + if (Queue.get_device().has_extension("cl_intel_subgroups_short")) { + check(Queue); + check(Queue); + } check(Queue); - check(Queue); + check(Queue); + check(Queue); + check(Queue); + check(Queue); + check(Queue); + check(Queue); + check(Queue); + check(Queue); + check(Queue); check(Queue); - check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp16")) { + check(Queue); + if (Queue.get_device().has_extension("cl_khr_fp16")) { check(Queue); } check(Queue); - if (!Queue.get_device().has_extension("cl_khr_fp64")) { + if (Queue.get_device().has_extension("cl_khr_fp64")) { check(Queue); } std::cout << "Test passed." << std::endl; diff --git a/sycl/test/sub_group/vote.cpp b/sycl/test/sub_group/vote.cpp index a11343a857dbb..ffc440549c836 100644 --- a/sycl/test/sub_group/vote.cpp +++ b/sycl/test/sub_group/vote.cpp @@ -1,15 +1,13 @@ // RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL -lsycl // RUN: env SYCL_DEVICE_TYPE=HOST %t.out -// TODO: Enable when use SPIRV operations instead direct built-ins calls. -// RUNx: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -//==--------------- vote.cpp - SYCL sub_group vote test --------------------==// +//==--------------- vote.cpp - SYCL sub_group vote test --*- C++ -*---------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/tools/get_device_count_by_type.cpp b/sycl/tools/get_device_count_by_type.cpp index b803ddba0b0ef..074c29d5bf81f 100644 --- a/sycl/tools/get_device_count_by_type.cpp +++ b/sycl/tools/get_device_count_by_type.cpp @@ -1,9 +1,8 @@ //==-- get_device_count_by_type.cpp - Get device count by type -------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// diff --git a/sycl/tools/sycl-check.cpp b/sycl/tools/sycl-check.cpp index d55e8170830b8..82de5f52cb03f 100644 --- a/sycl/tools/sycl-check.cpp +++ b/sycl/tools/sycl-check.cpp @@ -1,9 +1,8 @@ //==----------- sycl-check.cpp ---------------------------------------------==// // -// The LLVM Compiler Infrastructure -// -// This file is distributed under the University of Illinois Open Source -// License. See LICENSE.TXT for details. +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===//