Merge from 'sycl' to 'sycl-web' (triSYCL#47)

.github/CODEOWNERS

@@ -0,0 +1,19 @@
+clang/ @erichkeane, @Fznamznon
+
+clang/**/Driver @mdtoguchi @AGindinson
+
+llvm-spirv/ @AlexeySotkin, @AlexeySachkov
+
+opencl-aot/ @dm-vodopyanov, @AlexeySachkov, @romanovvlad
+
+libdevice/ @asavonic, @vzakhari
+
+sycl/doc/extensions/ @mkinsner, @jbrodman
+
+sycl/doc/ @pvchupin, @kbobrovs
+
+sycl/ @romanovvlad, @bader
+
+xpti/ @tovinkere, @andykaylor
+
+* @bader

.github/workflows/gh_pages.yml

@@ -13,14 +13,17 @@ jobs:
         ref: sycl
         path: repo
     - name: Install deps
-      run: sudo apt-get install -y doxygen graphviz ssh ninja-build
+      run: |
+        sudo apt-get install -y doxygen graphviz ssh ninja-build
+        sudo pip3 install sphinx recommonmark sphinx_markdown_tables
     - name: Build Docs
       run: |
         mkdir -p $GITHUB_WORKSPACE/build
         cd $GITHUB_WORKSPACE/build
         python $GITHUB_WORKSPACE/repo/buildbot/configure.py -w $GITHUB_WORKSPACE \
         -s $GITHUB_WORKSPACE/repo -o $GITHUB_WORKSPACE/build -t Release --docs
         cmake --build . --target doxygen-sycl
+        cmake --build . --target docs-sycl-html
     - name: Deploy
       env:
         SSH_KEY: ${{secrets.ACTIONS_DEPLOY_KEY}}
@@ -32,7 +35,11 @@ jobs:
         ssh-add -k ~/.ssh/id_rsa
         git clone git@github.com:intel/llvm-docs.git docs
         cd $GITHUB_WORKSPACE/docs
-        yes | \cp -rf $GITHUB_WORKSPACE/build/tools/sycl/doc/doxygen/html/* .
+        git rm -rf .
+        touch .nojekyll
+        yes | \cp -rf $GITHUB_WORKSPACE/build/tools/sycl/doc/html/* .
+        mkdir doxygen
+        yes | \cp -rf $GITHUB_WORKSPACE/build/tools/sycl/doc/doxygen/html/* doxygen/
         git config --global user.name "iclsrc"
         git config --global user.email "ia.compiler.tools.git@intel.com"
         git add .

buildbot/configure.py

@@ -20,6 +20,7 @@ def do_configure(args):
     sycl_build_pi_cuda = 'OFF'
     llvm_enable_assertions = 'ON'
     llvm_enable_doxygen = 'OFF'
+    llvm_enable_sphinx = 'OFF'
     llvm_build_shared_libs = 'OFF'
 
     if platform.system() == 'Linux':
@@ -38,6 +39,7 @@ def do_configure(args):
 
     if args.docs:
         llvm_enable_doxygen = 'ON'
+        llvm_enable_sphinx = 'ON'
 
     if args.shared_libs:
         llvm_build_shared_libs = 'ON'
@@ -63,6 +65,7 @@ def do_configure(args):
         "-DCMAKE_INSTALL_PREFIX={}".format(install_dir),
         "-DSYCL_INCLUDE_TESTS=ON", # Explicitly include all kinds of SYCL tests.
         "-DLLVM_ENABLE_DOXYGEN={}".format(llvm_enable_doxygen),
+        "-DLLVM_ENABLE_SPHINX={}".format(llvm_enable_sphinx),
         "-DBUILD_SHARED_LIBS={}".format(llvm_build_shared_libs),
         "-DSYCL_ENABLE_XPTI_TRACING=ON", # Explicitly turn on XPTI tracing
         llvm_dir

clang/CMakeLists.txt

@@ -237,14 +237,7 @@ set(ENABLE_X86_RELAX_RELOCATIONS OFF CACHE BOOL
 set(ENABLE_EXPERIMENTAL_NEW_PASS_MANAGER FALSE CACHE BOOL
   "Enable the experimental new pass manager by default.")
 
-# Clang tool executes cc1 commands in the same process after b4a99a0
-# It causes increasing memory consumption for compilations where several
-# source files are passed (for C++ application) or several build steps (e.g.
-# for SYCL application we have host, device and integration header step
-# per source file). Memory is not freed for all cc1 commands until end
-# This change forces clang driver use old behavior untill memory issue
-# is fixed.
-set(CLANG_SPAWN_CC1 ON CACHE BOOL
+set(CLANG_SPAWN_CC1 OFF CACHE BOOL
     "Whether clang should use a new process for the CC1 invocation")
 
 # TODO: verify the values against LangStandards.def?

clang/include/clang/Basic/Attr.td

@@ -236,6 +236,10 @@ class DefaultIntArgument<string name, int default> : IntArgument<name, 1> {
   int Default = default;
 }
 
+class DefaultUnsignedArgument<string name, int default> : UnsignedArgument<name, 1> {
+  int Default = default;
+}
+
 // This argument is more complex, it includes the enumerator type name,
 // a list of strings to accept, and a list of enumerators to map them to.
 class EnumArgument<string name, string type, list<string> values,
@@ -1676,6 +1680,57 @@ def SYCLIntelFPGAMaxConcurrency : Attr {
   let Documentation = [SYCLIntelFPGAMaxConcurrencyAttrDocs];
 }
 
+def SYCLIntelFPGALoopCoalesce : Attr {
+  let Spellings = [CXX11<"intelfpga","loop_coalesce">];
+  let Args = [ExprArgument<"NExpr">];
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let HasCustomTypeTransform = 1;
+  let AdditionalMembers = [{
+    static const char *getName() {
+      return "loop_coalesce";
+    }
+  }];
+  let Documentation = [SYCLIntelFPGALoopCoalesceAttrDocs];
+}
+
+def SYCLIntelFPGADisableLoopPipelining : Attr {
+  let Spellings = [CXX11<"intelfpga","disable_loop_pipelining">];
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let HasCustomTypeTransform = 1;
+  let AdditionalMembers = [{
+    static const char *getName() {
+      return "disable_loop_pipelining";
+    }
+  }];
+  let Documentation = [SYCLIntelFPGADisableLoopPipeliningAttrDocs];
+}
+
+def SYCLIntelFPGAMaxInterleaving : Attr {
+  let Spellings = [CXX11<"intelfpga","max_interleaving">];
+  let Args = [ExprArgument<"NExpr">];
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let HasCustomTypeTransform = 1;
+  let AdditionalMembers = [{
+    static const char *getName() {
+      return "max_interleaving";
+    }
+  }];
+  let Documentation = [SYCLIntelFPGAMaxInterleavingAttrDocs];
+}
+
+def SYCLIntelFPGASpeculatedIterations : Attr {
+  let Spellings = [CXX11<"intelfpga","speculated_iterations">];
+  let Args = [ExprArgument<"NExpr">];
+  let LangOpts = [SYCLIsDevice, SYCLIsHost];
+  let HasCustomTypeTransform = 1;
+  let AdditionalMembers = [{
+    static const char *getName() {
+      return "speculated_iterations";
+    }
+  }];
+  let Documentation = [SYCLIntelFPGASpeculatedIterationsAttrDocs];
+}
+
 def IntelFPGALocalNonConstVar : SubsetSubject<Var,
                               [{S->hasLocalStorage() &&
                                 S->getKind() != Decl::ImplicitParam &&
@@ -2471,13 +2526,16 @@ def NoDeref : TypeAttr {
   let Documentation = [NoDerefDocs];
 }
 
+// Default arguments in ReqWorkGroupSize can be used only with
+// intel::reqd_work_group_size spelling.
 def ReqdWorkGroupSize : InheritableAttr {
   let Spellings = [GNU<"reqd_work_group_size">,
-                   CXX11<"cl","reqd_work_group_size">];
-  let Args = [UnsignedArgument<"XDim">, UnsignedArgument<"YDim">,
-              UnsignedArgument<"ZDim">];
+                   CXX11<"intel","reqd_work_group_size">, 
+                   CXX11<"cl","reqd_work_group_size">]; 
+  let Args = [UnsignedArgument<"XDim">, DefaultUnsignedArgument<"YDim", 1>,
+              DefaultUnsignedArgument<"ZDim", 1>];
   let Subjects = SubjectList<[Function], ErrorDiag>;
-  let Documentation = [Undocumented];
+  let Documentation = [ReqdWorkGroupSizeAttrDocs];
 }
 
 def WorkGroupSizeHint :  InheritableAttr {

clang/include/clang/Basic/AttrDocs.td

@@ -2003,6 +2003,36 @@ device kernel, the attribute is ignored and it is not propagated to a kernel.
   }];
 }
 
+def ReqdWorkGroupSizeAttrDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "reqd_work_group_size";
+  let Content = [{
+This attribute is documented by both OpenCL and SYCL standards
+and allows to specify exact *local_work_size* which must be used as
+argument to **clEnqueueNDRangeKernel** (in OpenCL) or to
+**parallel_for** in SYCL. This allows the compiler to optimize the
+generated code appropriately for the kernel to which attribute is applied.
+
+While semantic of this attribute is the same between OpenCL and SYCL,
+spelling is a bit different: 
+
+SYCL 1.2.1 describes ``[[cl::reqd_work_group_size(X, Y, Z)]]`` spelling: this
+attribute is legal on device functions and is propagated down to any caller of 
+those device functions, such that the kernel attributes are the sum of all 
+attributes of all device functions called in this kernel.
+See section 6.7 Attributes for more details.
+
+As Intel extension, ``[[intel::reqd_work_group_size(X, Y, Z)]]`` spelling is allowed
+which features optional arguments `Y` and `Z`, those simplifies its usage if
+only 1- or 2-dimensional ND-range is assumed by a programmer. These arguments
+defaults to ``1``.
+
+In OpenCL C, this attribute is available in GNU spelling
+(``__attribute__((reqd_work_group_size(X, Y, Z)))``), see section
+6.7.2 Optional Attribute Qualifiers of OpenCL 1.2 specification for details.
+  }];
+}
+
 def SYCLIntelMaxWorkGroupSizeAttrDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "max_work_group_size (IntelFPGA)";
@@ -2110,6 +2140,53 @@ be applied multiple times to the same loop.
   }];
 }
 
+def SYCLIntelFPGALoopCoalesceAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "loop_coalesce";
+  let Content = [{
+This attribute applies to a loop. Indicates that the loop nest should be
+coalesced into a single loop without affecting functionality. Parameter N is
+optional. If specified, it shall be a positive integer, and indicates how many
+of the nested loop levels should be coalesced.
+  }];
+}
+
+def SYCLIntelFPGADisableLoopPipeliningAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "disable_loop_pipelining";
+  let Content = [{
+This attribute applies to a loop. Disables pipelining of the loop data path,
+causing the loop to be executed serially. Cannot be used on the same loop in
+conjunction with max_interleaving, speculated_iterations, max_concurrency, ii
+or ivdep.
+  }];
+}
+
+def SYCLIntelFPGAMaxInterleavingAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "max_interleaving";
+  let Content = [{
+This attribute applies to a loop. Places a maximum limit N on the number of
+interleaved invocations of an inner loop by an outer loop (note, this does not
+mean that this attribute can only be applied to inner loops in user code - outer
+loops in user code may still be contained in an implicit loop due to NDRange).
+Parameter N is mandatory, and shall be non-negative integer. Cannot be
+used on the same loop in conjunction with disable_loop_pipelining.
+  }];
+}
+
+def SYCLIntelFPGASpeculatedIterationsAttrDocs : Documentation {
+  let Category = DocCatVariable;
+  let Heading = "speculated_iterations";
+  let Content = [{
+This attribute applies to a loop. Specifies the number of concurrent speculated
+iterations that will be in flight for a loop invocation (i.e. the exit
+condition for these iterations will not have been evaluated yet).
+Parameter N is mandatory, and may either be 0, or a positive integer. Cannot be
+used on the same loop in conjunction with disable_loop_pipelining.
+  }];
+}
+
 def SYCLDeviceIndirectlyCallableDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "intel::device_indirectly_callable";

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10666,9 +10666,9 @@ def err_builtin_launder_invalid_arg : Error<
 // SYCL-specific diagnostics
 def err_sycl_attribute_address_space_invalid : Error<
   "address space is outside the valid range of values">;
-def err_sycl_kernel_name_class_not_top_level : Error<
- "kernel name class and its template argument classes' declarations can only "
- "nest in a namespace: %0">;
+def err_sycl_kernel_incorrectly_named : Error<
+  "kernel %select{name is missing"
+  "|needs to have a globally-visible name}0">;
 def err_sycl_restrict : Error<
   "SYCL kernel cannot "
 "%select{use a non-const global variable"

clang/include/clang/Driver/Action.h

@@ -17,6 +17,7 @@
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/iterator_range.h"
+#include <initializer_list>
 #include <string>
 
 namespace llvm {
@@ -78,9 +79,10 @@ class Action {
     SYCLPostLinkJobClass,
     PartialLinkJobClass,
     BackendCompileJobClass,
+    FileTableTformJobClass,
 
     JobClassFirst = PreprocessJobClass,
-    JobClassLast = BackendCompileJobClass
+    JobClassLast = FileTableTformJobClass
   };
 
   // The offloading kind determines if this action is binded to a particular
@@ -679,6 +681,13 @@ class SYCLPostLinkJobAction : public JobAction {
   static bool classof(const Action *A) {
     return A->getKind() == SYCLPostLinkJobClass;
   }
+
+  void setRTSetsSpecConstants(bool Val) { RTSetsSpecConsts = Val; }
+
+  bool getRTSetsSpecConstants() const { return RTSetsSpecConsts; }
+
+private:
+  bool RTSetsSpecConsts = true;
 };
 
 class PartialLinkJobAction : public JobAction {
@@ -705,6 +714,47 @@ class BackendCompileJobAction : public JobAction {
   }
 };
 
+// Represents a file table transformation action. The order of inputs to a
+// FileTableTformJobAction at construction time must accord with the tforms
+// added later - some tforms "consume" inputs. For example, "replace column"
+// needs another file to read the replacement column from.
+class FileTableTformJobAction : public JobAction {
+  void anchor() override;
+
+public:
+  struct Tform {
+    enum Kind { EXTRACT, EXTRACT_DROP_TITLE, REPLACE };
+
+    Tform() = default;
+    Tform(Kind K, std::initializer_list<StringRef> Args) : TheKind(K) {
+      for (auto A : Args)
+        TheArgs.emplace_back(A.str());
+    }
+
+    Kind TheKind;
+    SmallVector<std::string, 2> TheArgs;
+  };
+
+  FileTableTformJobAction(Action *Input, types::ID OutputType);
+  FileTableTformJobAction(ActionList &Inputs, types::ID OutputType);
+
+  // Deletes all columns except the one with given name.
+  void addExtractColumnTform(StringRef ColumnName, bool WithColTitle = true);
+
+  // Replaces a column with title <From> in this table with a column with title
+  // <To> from another file table passed as input to this action.
+  void addReplaceColumnTform(StringRef From, StringRef To);
+
+  static bool classof(const Action *A) {
+    return A->getKind() == FileTableTformJobClass;
+  }
+
+  const ArrayRef<Tform> getTforms() const { return Tforms; }
+
+private:
+  SmallVector<Tform, 2> Tforms; // transformation actions requested
+};
+
 } // namespace driver
 } // namespace clang
 

clang/include/clang/Driver/CC1Options.td

@@ -919,6 +919,8 @@ def fsycl_std_layout_kernel_params: Flag<["-"], "fsycl-std-layout-kernel-params"
 def fsycl_allow_func_ptr : Flag<["-"], "fsycl-allow-func-ptr">,
   HelpText<"Allow function pointers in SYCL device.">;
 def fno_sycl_allow_func_ptr : Flag<["-"], "fno-sycl-allow-func-ptr">;
+def fsycl_enable_optimizations: Flag<["-"], "fsycl-enable-optimizations">,
+  HelpText<"Experimental flag enabling standard optimization in the front-end.">;
 
 } // let Flags = [CC1Option]
 

clang/include/clang/Driver/ToolChain.h

@@ -147,6 +147,7 @@ class ToolChain {
   mutable std::unique_ptr<Tool> SYCLPostLink;
   mutable std::unique_ptr<Tool> PartialLink;
   mutable std::unique_ptr<Tool> BackendCompiler;
+  mutable std::unique_ptr<Tool> FileTableTform;
 
   Tool *getClang() const;
   Tool *getFlang() const;
@@ -161,6 +162,7 @@ class ToolChain {
   Tool *getSYCLPostLink() const;
   Tool *getPartialLink() const;
   Tool *getBackendCompiler() const;
+  Tool *getTableTform() const;
 
   mutable std::unique_ptr<SanitizerArgs> SanitizerArguments;
   mutable std::unique_ptr<XRayArgs> XRayArguments;