Merge branch 'sycl' into empty_task_removal

Author: KseniyaTikhomirova

.github/workflows/gh_pages.yml

@@ -16,7 +16,7 @@ jobs:
     runs-on: ubuntu-latest
     if: github.repository == 'intel/llvm'
     steps:
-    - uses: actions/checkout@v2
+    - uses: actions/checkout@v3
       with:
         path: repo
     - name: Install deps

.github/workflows/sycl_containers.yaml

@@ -28,7 +28,7 @@ jobs:
     runs-on: ubuntu-latest
     steps:
       - name: Checkout
-        uses: actions/checkout@v2
+        uses: actions/checkout@v3
         with:
           fetch-depth: 2
       - name: Build and Push Container
@@ -47,7 +47,7 @@ jobs:
     runs-on: ubuntu-latest
     steps:
       - name: Checkout
-        uses: actions/checkout@v2
+        uses: actions/checkout@v3
         with:
           fetch-depth: 2
       - name: Build and Push Container
@@ -70,18 +70,17 @@ jobs:
     needs: base_image_ubuntu2004
     steps:
       - name: Checkout
-        uses: actions/checkout@v2
+        uses: actions/checkout@v3
         with:
           fetch-depth: 2
       - name: Get dependencies configuration
         id: deps
         run: |
           DEPS=`cat devops/dependencies.json`
-          DEPS="${DEPS//'%'/'%25'}"
-          DEPS="${DEPS//$'\n'/'%0A'}"
-          DEPS="${DEPS//$'\r'/'%0D'}"
+          DEPS="${DEPS//$'\r'/''}"
+          DEPS="${DEPS//$'\n'/' '}"
           echo $DEPS
-          echo "::set-output name=deps::$DEPS"
+          echo "deps=$DEPS" >>$GITHUB_OUTPUT
       - name: Build and Push Container
         uses: ./devops/actions/build_container
         with:
@@ -108,18 +107,17 @@ jobs:
     needs: base_image_ubuntu2004
     steps:
       - name: Checkout
-        uses: actions/checkout@v2
+        uses: actions/checkout@v3
         with:
           fetch-depth: 2
       - name: Get dependencies configuration
         id: deps
         run: |
           DEPS=`cat devops/dependencies.json`
-          DEPS="${DEPS//'%'/'%25'}"
-          DEPS="${DEPS//$'\n'/'%0A'}"
-          DEPS="${DEPS//$'\r'/'%0D'}"
+          DEPS="${DEPS//$'\r'/''}"
+          DEPS="${DEPS//$'\n'/' '}"
           echo $DEPS
-          echo "::set-output name=deps::$DEPS"
+          echo "deps=$DEPS" >>$GITHUB_OUTPUT
       - name: Build and Push Container
         uses: ./devops/actions/build_container
         with:

.github/workflows/sycl_linux_build_and_test.yml

@@ -161,12 +161,12 @@ jobs:
     - name: Pack LIT
       run: tar -cJf lit.tar.xz -C $GITHUB_WORKSPACE/src/llvm/utils/lit .
     - name: Upload toolchain
-      uses: actions/upload-artifact@v2
+      uses: actions/upload-artifact@v3
       with:
         name: sycl_linux_${{ inputs.build_artifact_suffix }}
         path: llvm_sycl.tar.xz
     - name: Upload LIT
-      uses: actions/upload-artifact@v2
+      uses: actions/upload-artifact@v3
       with:
         name: sycl_lit_${{ inputs.build_artifact_suffix }}
         path: lit.tar.xz
@@ -226,7 +226,7 @@ jobs:
           sudo -E /opt/install_drivers.sh --all
         fi
     # FIXME cached_checkout fails here, but works everywhere else
-    - uses: actions/checkout@v2
+    - uses: actions/checkout@v3
       with:
         persist-credentials: false
         path: llvm
@@ -277,7 +277,7 @@ jobs:
         fi
     # FIXME cached_checkout fails here, but works everywhere else
     # TODO: figure out if we remove this action
-    - uses: actions/checkout@v2
+    - uses: actions/checkout@v3
       with:
         path: llvm
     # TODO should this action be packed into container as well?

.github/workflows/sycl_macos_build_and_test.yml

@@ -66,7 +66,7 @@ jobs:
     - name: Pack toolchain
       run: tar -cJf llvm_sycl.tar.xz -C $GITHUB_WORKSPACE/build/install .
     - name: Upload toolchain
-      uses: actions/upload-artifact@v2
+      uses: actions/upload-artifact@v3
       with:
         name: sycl_macos_${{ inputs.build_artifact_suffix }}
         path: llvm_sycl.tar.xz

.github/workflows/sycl_nightly.yml

@@ -47,8 +47,8 @@ jobs:
     runs-on: ubuntu-latest
     needs: ubuntu2004_build_test
     steps:
-    - uses: actions/checkout@v2
-    - uses: actions/download-artifact@v2
+    - uses: actions/checkout@v3
+    - uses: actions/download-artifact@v3
       with:
         name: sycl_linux_default
         path: devops/

.github/workflows/sycl_post_commit.yml

@@ -44,13 +44,13 @@ jobs:
     runs-on: ubuntu-20.04
     if: github.repository == 'intel/llvm'
     steps:
-    - uses: actions/checkout@v2
+    - uses: actions/checkout@v3
       with:
         path: src
     - name: Install Ubuntu deps
       run: sudo apt install -y ninja-build ccache
     - name: Setup Cache
-      uses: actions/cache@v2
+      uses: actions/cache@v3
       id: cache
       with:
         path: ${{ github.workspace }}/cache
@@ -104,7 +104,7 @@ jobs:
     - name: Pack
       run: tar -czvf llvm_sycl.tar.gz -C $GITHUB_WORKSPACE/build/install .
     - name: Upload artifacts
-      uses: actions/upload-artifact@v1
+      uses: actions/upload-artifact@v3
       with:
         name: sycl_linux_${{ matrix.config }}
         path: llvm_sycl.tar.gz

.github/workflows/sycl_precommit.yml

@@ -29,7 +29,7 @@ jobs:
     steps:
     - name: 'PR commits + 1'
       run: echo "PR_FETCH_DEPTH=$(( ${{ github.event.pull_request.commits }} + 1 ))" >> "${GITHUB_ENV}"
-    - uses: actions/checkout@v2
+    - uses: actions/checkout@v3
       with:
         ref: ${{ github.event.pull_request.head.sha }}
         persist-credentials: false

.github/workflows/sycl_update_gpu_driver.yml

@@ -10,7 +10,7 @@ jobs:
     runs-on: ubuntu-latest
     if: github.repository == 'intel/llvm'
     steps:
-    - uses: actions/checkout@v2
+    - uses: actions/checkout@v3
     - name: Update dependencies file
       run: |
         version="$(python3 devops/scripts/update_drivers.py linux)"

.github/workflows/sycl_windows_build_and_test.yml

@@ -23,7 +23,7 @@ jobs:
         git config --global core.autocrlf false
         echo "C:\Program Files\Git\usr\bin" | Out-File -FilePath $env:GITHUB_PATH -Encoding utf8 -Append
         echo "SCCACHE_DIR=D:\github\_work\cache\${{ inputs.build_cache_suffix }}" | Out-File -FilePath $env:GITHUB_ENV -Encoding utf8 -Append
-    - uses: actions/checkout@v2
+    - uses: actions/checkout@v3
       with:
         path: src
         fetch-depth: 1
@@ -78,7 +78,7 @@ jobs:
       shell: bash
       run: cmake --build build --target deploy-sycl-toolchain
     - name: Upload toolchain
-      uses: actions/upload-artifact@v2
+      uses: actions/upload-artifact@v3
       with:
         name: sycl_windows_default
         path: install/**/*

.github/workflows/sync-main.yml

@@ -8,7 +8,7 @@ jobs:
     runs-on: ubuntu-latest
     if: github.repository == 'intel/llvm'
     steps:
-      - uses: actions/checkout@v2
+      - uses: actions/checkout@v3
         with:
           # persist-credentials: false allows us to use our own credentials for
           # pushing to the repository.  Otherwise, the default github actions token

clang/include/clang/Basic/Attr.td

@@ -1302,11 +1302,13 @@ def SYCLType: InheritableAttr {
                            ["accessor", "local_accessor", "spec_constant",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
-                            "private_memory", "aspect"],
+                            "private_memory", "aspect", "annotated_ptr", "annotated_arg",
+                            "stream", "sampler"],
                            ["accessor", "local_accessor", "spec_constant",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
-                            "private_memory", "aspect"]>];
+                            "private_memory", "aspect", "annotated_ptr", "annotated_arg",
+                            "stream", "sampler"]>];
   // Only used internally by SYCL implementation
   let Documentation = [InternalOnly];
 }
@@ -1603,6 +1605,20 @@ def SYCLAddIRAttrCommonMembers : SYCLAddIRAttrMemberCodeHolder<[{
     if (ValueQType->isIntegralOrEnumerationType() ||
         ValueQType->isFloatingType())
       return Value.getAsString(Context, ValueQType);
+    if (ValueQType->isArrayType() &&
+        (ValueQType->getArrayElementTypeNoTypeQual()->isCharType() ||
+         ValueQType->getArrayElementTypeNoTypeQual()
+                   ->isIntegralOrEnumerationType())) {
+      SmallString<10> StrBuffer;
+      for (unsigned I = 0; I < Value.getArraySize(); ++I) {
+        const APValue &ArrayElem = Value.getArrayInitializedElt(I);
+        char C = static_cast<char>(ArrayElem.getInt().getExtValue());
+        if (C == 0)
+          break;
+        StrBuffer += C;
+      }
+      return std::string(StrBuffer);
+    }
     return None;
   }
 
@@ -1628,6 +1644,33 @@ def SYCLAddIRAttrCommonMembers : SYCLAddIRAttrMemberCodeHolder<[{
                                       ValueE->getType()->isSignedIntegerType());
       return std::string(IntegerStrBuffer);
     }
+    if (const auto *InitListE = dyn_cast<InitListExpr>(ValueE)) {
+      if (InitListE->isStringLiteralInit()) {
+        const Expr *StringInitE = InitListE->getInit(0)->IgnoreParenImpCasts();
+        return getValidAttributeValueAsString(StringInitE, Context);
+      }
+
+      SmallString<10> StrBuffer;
+      for (const auto *InitE : InitListE->inits()) {
+        const Expr *InitNoImpCastE = InitE->IgnoreParenImpCasts();
+        char C = 0;
+        if (const auto *CharacterVal =
+              dyn_cast<CharacterLiteral>(InitNoImpCastE))
+          C = static_cast<char>(CharacterVal->getValue());
+        else if (const auto *IntegerVal =
+              dyn_cast<IntegerLiteral>(InitNoImpCastE))
+          C = static_cast<char>(IntegerVal->getValue().getSExtValue());
+        else
+          return None;
+
+        // Null terminator will end the string reading.
+        if (C == 0)
+          break;
+
+        StrBuffer += C;
+      }
+      return std::string(StrBuffer);
+    }
 
     const auto *ValueCE = dyn_cast<ConstantExpr>(ValueE);
     if (!ValueCE)
@@ -3427,6 +3470,7 @@ def WorkGroupSizeHint :  InheritableAttr {
     }
   }];
   let Documentation = [WorkGroupSizeHintAttrDocs];
+  let SupportsNonconformingLambdaSyntax = 1;
 }
 
 def InitPriority : InheritableAttr, TargetSpecificAttr<TargetSupportsInitPriority> {

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10390,8 +10390,8 @@ def warn_opencl_generic_address_space_arg : Warning<
   "passing non-generic address space pointer to %0"
   " may cause dynamic conversion affecting performance">,
   InGroup<Conversion>, DefaultIgnore;
-def err_bad_union_kernel_param_members : Error<
-  "%0 cannot be used inside a union kernel parameter">;
+def err_bad_kernel_param_data_members : Error<
+  "%0 cannot be a data member of a %select{union|struct}1 kernel parameter">;
 
 // OpenCL v2.0 s6.13.6 -- Builtin Pipe Functions
 def err_opencl_builtin_pipe_first_arg : Error<
@@ -11887,6 +11887,9 @@ def err_sycl_add_ir_attribute_invalid_value : Error<
 def err_sycl_add_ir_attribute_invalid_filter : Error<
   "initializer list in the first argument of %0 must contain only string "
   "literals">;
+def warn_sycl_old_and_new_kernel_attributes : Warning<
+  "kernel has both attribute %0 and kernel properties; conflicting properties "
+  "are ignored">, InGroup<IgnoredAttributes>;
 
 // errors of expect.with.probability
 def err_probability_not_constant_float : Error<

clang/include/clang/Basic/LangOptions.def

@@ -279,6 +279,7 @@ LANGOPT(IntelFPGA         , 1, 0, "Perform ahead-of-time compilation for FPGA")
 LANGOPT(SYCLAllowFuncPtr  , 1, 0, "Allow function pointers in SYCL device code")
 LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for SYCL kernel parameters")
 LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels")
+LANGOPT(SYCLForceInlineKernelLambda , 1, 0, "Force inline SYCL kernel lambdas in entry point")
 LANGOPT(SYCLESIMDForceStatelessMem, 1, 0, "Make accessors use USM memory in ESIMD kernels")
 ENUM_LANGOPT(SYCLVersion  , SYCLMajorVersion, 2, SYCL_None, "Version of the SYCL standard used")
 LANGOPT(DeclareSPIRVBuiltins, 1, 0, "Declare SPIR-V builtin functions")

clang/include/clang/Driver/Driver.h

@@ -797,6 +797,11 @@ class Driver {
   /// targets.
   mutable llvm::StringMap<StringRef> SYCLUniqueIDList;
 
+  /// Vector of Macros that need to be added to the Host compilation in a
+  /// SYCL based offloading scenario.  These macros are gathered during
+  /// construction of the device compilations.
+  mutable std::vector<std::string> SYCLTargetMacroArgs;
+
   /// Return the typical executable name for the specified driver \p Mode.
   static const char *getExecutableForDriverMode(DriverMode Mode);
 
@@ -867,6 +872,17 @@ class Driver {
   void createAppendedFooterInput(Action *&Input, Compilation &C,
                                  const llvm::opt::ArgList &Args) const;
 
+  /// addSYCLTargetMacroArg - Add the given macro to the vector of args to be
+  /// added to the host compilation step.
+  void addSYCLTargetMacroArg(const llvm::opt::ArgList &Args,
+                             StringRef Macro) const {
+    SYCLTargetMacroArgs.push_back(Args.MakeArgString(Macro));
+  }
+  /// getSYCLTargetMacroArgs - return the previously gathered macro target args.
+  llvm::ArrayRef<std::string> getSYCLTargetMacroArgs() const {
+    return SYCLTargetMacroArgs;
+  }
+
   /// setSYCLUniqueID - set the Unique ID that is used for all FE invocations
   /// when performing compilations for SYCL.
   void addSYCLUniqueID(StringRef UniqueID, StringRef FileName) const {

clang/include/clang/Driver/Options.td

@@ -2942,6 +2942,12 @@ defm sycl_unnamed_lambda
               " >= clang::LangOptions::SYCLMajorVersion::SYCL_2020")>,
           PosFlag<SetTrue, [], "Allow">, NegFlag<SetFalse, [], "Disallow">,
           BothFlags<[CC1Option, CoreOption], " unnamed SYCL lambda kernels">>;
+defm sycl_force_inline_kernel_lambda
+    : BoolFOption<
+          "sycl-force-inline-kernel-lambda", LangOpts<"SYCLForceInlineKernelLambda">,
+          DefaultTrue,
+          PosFlag<SetTrue, [], "Allow">, NegFlag<SetFalse, [], "Disallow">,
+          BothFlags<[CC1Option, CoreOption], " force inline SYCL kernels lambda in entry point">>;
 def fsycl_help_EQ : Joined<["-"], "fsycl-help=">,
   Flags<[NoXarchOption, CoreOption]>, HelpText<"Emit help information from the "
   "related offline compilation tool. Valid values: all, fpga, gen, x86_64.">,

clang/include/clang/Sema/Sema.h

@@ -10984,6 +10984,7 @@ class Sema final {
   SYCLIntelMaxWorkGroupSizeAttr *
   MergeSYCLIntelMaxWorkGroupSizeAttr(Decl *D,
                                      const SYCLIntelMaxWorkGroupSizeAttr &A);
+  void CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D);
   SYCLAddIRAttributesFunctionAttr *MergeSYCLAddIRAttributesFunctionAttr(
       Decl *D, const SYCLAddIRAttributesFunctionAttr &A);
   void AddSYCLAddIRAttributesFunctionAttr(Decl *D,

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -80,19 +80,30 @@ void CodeGenTypes::addRecordTypeName(const RecordDecl *RD,
                 OS << "i" << TTy->getIntegerBitWidth();
                 break;
               }
-            } else if (TTy->isBFloatTy())
+            } else if (TTy->isHalfTy()) {
+              OS << "half";
+            } else if (TTy->isFloatTy()) {
+              OS << "float";
+            } else if (TTy->isDoubleTy()) {
+              OS << "double";
+            } else if (TTy->isBFloatTy()) {
               OS << "bfloat16";
-            else if (TTy->isStructTy()) {
+            } else if (TTy->isStructTy()) {
               StringRef LlvmTyName = TTy->getStructName();
-              // Emit half/bfloat16 for sycl[::*]::{half,bfloat16}
+              // Emit half/bfloat16/tf32 for sycl[::*]::{half,bfloat16,tf32}
               if (LlvmTyName.startswith("class.sycl::") ||
                   LlvmTyName.startswith("class.__sycl_internal::"))
                 LlvmTyName = LlvmTyName.rsplit("::").second;
+              if (LlvmTyName != "half" && LlvmTyName != "bfloat16" &&
+                  LlvmTyName != "tf32")
+                llvm_unreachable("Wrong matrix base type!");
               OS << LlvmTyName;
-            } else
-              TTy->print(OS, false, true);
-          } else if (TemplateArg.getKind() == TemplateArgument::Integral)
+            } else {
+              llvm_unreachable("Wrong matrix base type!");
+            }
+          } else if (TemplateArg.getKind() == TemplateArgument::Integral) {
             OS << TemplateArg.getAsIntegral();
+          }
         }
         Ty->setName(OS.str());
         return;