Merge remote-tracking branch 'origin/sycl' into private/asachkov/fix-integer-vec-conversions

This commit also fixes conversions to/from `std::byte`

Author: AlexeySachkov

.github/CODEOWNERS

@@ -59,6 +59,8 @@ xptifw/ @intel/llvm-reviewers-runtime
 
 # DPC++ tools
 llvm/ @intel/dpcpp-tools-reviewers
+clang/lib/CodeGen/BackendUtil.cpp @intel/dpcpp-cfe-reviewers @intel/dpcpp-tools-reviewers
+clang/include/clang/CodeGen/BackendUtil.h @intel/dpcpp-cfe-reviewers @intel/dpcpp-tools-reviewers
 
 # OpenCL CPU RT
 opencl-cpu/ @intel/ocl-cpu-rt-write

README.md

@@ -81,7 +81,7 @@ guidelines.
 
 ## Late-outline OpenMP\* and OpenMP\* Offload
 
-See [openmp](/tree/openmp) branch.
+See [openmp](/openmp) branch.
 
 # License
 

clang/include/clang/Basic/DiagnosticCommonKinds.td

@@ -436,4 +436,10 @@ def warn_try_not_valid_on_target : Warning<
   "target '%0' does not support exception handling;"
   " 'catch' block is ignored">,
   InGroup<OpenMPTargetException>;
+
+// Launch bound
+def warn_launch_bounds_missing_attr: Warning<
+  "%0 attribute ignored, as it requires: maximum work group size"
+  "%select{| and minimum work groups per compute unit}1 to be also specified">,
+  InGroup<IgnoredAttributes>;
 }

clang/lib/CodeGen/Targets/NVPTX.cpp

@@ -245,29 +245,47 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes(
       // And kernel functions are not subject to inlining
       F->addFnAttr(llvm::Attribute::NoInline);
     }
+    bool HasMaxWorkGroupSize = false;
+    bool HasMinWorkGroupPerCU = false;
     if (const auto *MWGS = FD->getAttr<SYCLIntelMaxWorkGroupSizeAttr>()) {
       auto MaxThreads = (*MWGS->getZDimVal()).getExtValue() *
                         (*MWGS->getYDimVal()).getExtValue() *
                         (*MWGS->getXDimVal()).getExtValue();
-      if (MaxThreads > 0)
+      if (MaxThreads > 0) {
         addNVVMMetadata(F, "maxntidx", MaxThreads);
+        HasMaxWorkGroupSize = true;
+      }
+    }
 
-      auto attrValue = [&](Expr *E) {
-        const auto *CE = cast<ConstantExpr>(E);
-        std::optional<llvm::APInt> Val = CE->getResultAsAPSInt();
-        return Val->getZExtValue();
-      };
-
-      if (const auto *MWGPCU =
-              FD->getAttr<SYCLIntelMinWorkGroupsPerComputeUnitAttr>()) {
+    auto attrValue = [&](Expr *E) {
+      const auto *CE = cast<ConstantExpr>(E);
+      std::optional<llvm::APInt> Val = CE->getResultAsAPSInt();
+      return Val->getZExtValue();
+    };
+
+    if (const auto *MWGPCU =
+            FD->getAttr<SYCLIntelMinWorkGroupsPerComputeUnitAttr>()) {
+      if (!HasMaxWorkGroupSize && FD->hasAttr<OpenCLKernelAttr>()) {
+        M.getDiags().Report(D->getLocation(),
+                            diag::warn_launch_bounds_missing_attr)
+            << MWGPCU << 0;
+      } else {
         // The value is guaranteed to be > 0, pass it to the metadata.
         addNVVMMetadata(F, "minnctapersm", attrValue(MWGPCU->getValue()));
+        HasMinWorkGroupPerCU = true;
+      }
+    }
 
-        if (const auto *MWGPMP =
-                FD->getAttr<SYCLIntelMaxWorkGroupsPerMultiprocessorAttr>()) {
-          // The value is guaranteed to be > 0, pass it to the metadata.
-          addNVVMMetadata(F, "maxclusterrank", attrValue(MWGPMP->getValue()));
-        }
+    if (const auto *MWGPMP =
+            FD->getAttr<SYCLIntelMaxWorkGroupsPerMultiprocessorAttr>()) {
+      if ((!HasMaxWorkGroupSize || !HasMinWorkGroupPerCU) &&
+          FD->hasAttr<OpenCLKernelAttr>()) {
+        M.getDiags().Report(D->getLocation(),
+                            diag::warn_launch_bounds_missing_attr)
+            << MWGPMP << 1;
+      } else {
+        // The value is guaranteed to be > 0, pass it to the metadata.
+        addNVVMMetadata(F, "maxclusterrank", attrValue(MWGPMP->getValue()));
       }
     }
   }

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5398,8 +5398,16 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
         break;
       }
     }
-    if (HasFPGA) {
+    // At -O0, imply -fsycl-disable-range-rounding.
+    bool DisableRangeRounding = false;
+    if (Arg *A = Args.getLastArg(options::OPT_O_Group)) {
+      if (A->getOption().matches(options::OPT_O0))
+        DisableRangeRounding = true;
+    }
+    if (DisableRangeRounding || HasFPGA)
       CmdArgs.push_back("-fsycl-disable-range-rounding");
+
+    if (HasFPGA) {
       // Pass -fintelfpga to both the host and device SYCL compilations if set.
       CmdArgs.push_back("-fintelfpga");
     }
@@ -6106,10 +6114,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
     A->render(Args, CmdArgs);
   }
 
-  if (Arg *A = Args.getLastArg(options::OPT_faltmathlib_EQ)) {
-    StringRef Name = A->getValue();
+  if (Arg *A = Args.getLastArg(options::OPT_faltmathlib_EQ))
     A->render(Args, CmdArgs);
-  }
 
   if (Args.hasFlag(options::OPT_fmerge_all_constants,
                    options::OPT_fno_merge_all_constants, false))

clang/test/Driver/sycl-offload.c

@@ -428,6 +428,26 @@
 // RUN:   | FileCheck -check-prefix=CHK-TOOLS-OPTS2 %s
 // CHK-TOOLS-OPTS2: clang-offload-wrapper{{.*}} "-link-opts=-DFOO1 -DFOO2"
 
+/// -fsycl-disable-range-rounding settings
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl \
+// RUN:        -fsycl-targets=spir64 -O0 %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-DISABLE-RANGE-ROUNDING %s
+// RUN: %clang_cl -### -fsycl -fsycl-targets=spir64 -Od %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-DISABLE-RANGE-ROUNDING %s
+// CHK-DISABLE-RANGE-ROUNDING: "-fsycl-disable-range-rounding"
+
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl \
+// RUN:        -fsycl-targets=spir64 -O2 %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-RANGE-ROUNDING %s
+// RUN: %clang_cl -### -fsycl -fsycl-targets=spir64 -O2 %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-RANGE-ROUNDING %s
+// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl \
+// RUN:        -fsycl-targets=spir64 %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-RANGE-ROUNDING %s
+// RUN: %clang_cl -### -fsycl -fsycl-targets=spir64 %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=CHK-RANGE-ROUNDING %s
+// CHK-RANGE-ROUNDING-NOT: "-fsycl-disable-range-rounding"
+
 /// ###########################################################################
 
 /// Verify that triple-boundarch pairs are correct with multi-targetting

clang/test/SemaSYCL/lb_sm_90.cpp

@@ -1,8 +1,9 @@
-// RUN: %clang_cc1 -internal-isystem %S/Inputs %s -triple nvptx64-nvidia-cuda -target-cpu sm_90 -fsycl-is-device -fsyntax-only -Wno-c++23-extensions -verify
-// expected-no-diagnostics
+// RUN: %clang_cc1 -internal-isystem %S/Inputs %s -triple nvptx64-nvidia-cuda -target-cpu sm_90 -fsycl-is-device -fsyntax-only -Wno-c++23-extensions -verify -S
 
 // Maximum work groups per multi-processor, mapped to maxclusterrank PTX
-// directive, is an SM_90 feature, make sure that no warnings/errors are issued.
+// directive, is an SM_90 feature. Attributes need to be used in sequence:
+// max_work_group_size, min_work_groups_per_cu, max_work_groups_per_mp, warn on
+// missing attributes in sequences.
 
 #include "sycl.hpp"
 
@@ -13,19 +14,44 @@ template <int N1, int N2, int N3> class Functor {
   operator()() const {}
 };
 
+// expected-warning@+1 {{'max_work_groups_per_mp' attribute ignored, as it requires: maximum work group size and minimum work groups per compute unit to be also specified}}
+template <int N1, int N2> class Functor_2 {
+public:
+  [[intel::max_work_group_size(1, 1, N1),
+    intel::max_work_groups_per_mp(N2)]] void
+  operator()() const {}
+};
+
 int main() {
   sycl::queue Q{};
 
   Q.submit([&](sycl::handler &cgh) {
-    cgh.single_task<class T1>( [=] [[intel::max_work_group_size(1, 1, 256),
-                                          intel::min_work_groups_per_cu(2),
-                                          intel::max_work_groups_per_mp(4)]] (
-) { volatile int A = 42; });
+    cgh.single_task<class T1>(
+        [=] [[intel::max_work_group_size(1, 1, 256),
+              intel::min_work_groups_per_cu(2),
+              intel::max_work_groups_per_mp(4)]] () { volatile int A = 42; });
+
+    // expected-warning@+2 {{'max_work_groups_per_mp' attribute ignored, as it requires: maximum work group size and minimum work groups per compute unit to be also specified}}
+    cgh.single_task<class T2>(
+        [=] [[intel::max_work_group_size(1, 1, 256),
+              intel::max_work_groups_per_mp(4)]] () { volatile int A = 42; });
+
+    // expected-warning@+2 {{'max_work_groups_per_mp' attribute ignored, as it requires: maximum work group size and minimum work groups per compute unit to be also specified}}
+    cgh.single_task<class T3>(
+        [=] [[intel::max_work_groups_per_mp(4)]] () { volatile int A = 42; });
+
+    // expected-warning@+2 {{'min_work_groups_per_cu' attribute ignored, as it requires: maximum work group size to be also specified}} cgh.single_task<class T4>(
+    cgh.single_task<class T4>(
+        [=] [[intel::min_work_groups_per_cu(4)]] () { volatile int A = 42; });
   });
 
   Q.submit([&](sycl::handler &cgh) {
     cgh.single_task<class F>(Functor<512, 8, 16>{});
   });
 
+  Q.submit([&](sycl::handler &cgh) {
+    cgh.single_task<class F2>(Functor_2<512, 8>{});
+  });
+
   return 0;
 }

llvm/lib/SYCLLowerIR/CompileTimeProperties.def

@@ -60,3 +60,8 @@ SYCL_COMPILE_TIME_PROPERTY("sycl-bi-directional-ports-true", 5885,
 // SPIR-V Spec: https://github.com/KhronosGroup/SPIRV-Registry/blob/main/extensions/INTEL/SPV_INTEL_cache_controls.asciidoc
 SYCL_COMPILE_TIME_PROPERTY("sycl-prefetch-hint", 6442, DecorValueTy::uint32)
 SYCL_COMPILE_TIME_PROPERTY("sycl-prefetch-hint-nt", 6442, DecorValueTy::uint32)
+
+// The corresponding SPIR-V OpCodes for cache control properties
+SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-hint", 6442, DecorValueTy::uint32)
+SYCL_COMPILE_TIME_PROPERTY("sycl-cache-read-assertion", 6442, DecorValueTy::uint32)
+SYCL_COMPILE_TIME_PROPERTY("sycl-cache-write-hint", 6443, DecorValueTy::uint32)