Merge branch 'sycl' into pulldown-ww29

Author: dwoodwor-intel

clang/test/SemaSYCL/Inputs/CL/sycl/detail/kernel_desc.hpp

@@ -35,6 +35,17 @@ __SYCL_INLINE_NAMESPACE(cl) {
     int offset;
   };
 
+  template <bool Cond, typename TrueT, typename FalseT>
+  struct conditional {
+    using type = TrueT;
+  };
+  template <typename TrueT, typename FalseT>
+  struct conditional<false, TrueT, FalseT> {
+    using type = FalseT;
+  };
+
+  using int64_t = conditional<sizeof(long) == 8, long, long long>::type;
+
   template <class KernelNameType> struct KernelInfo {
     static constexpr unsigned getNumParams() { return 0; }
     static const kernel_param_desc_t &getParamDesc(int) {
@@ -43,6 +54,7 @@ __SYCL_INLINE_NAMESPACE(cl) {
     }
     static constexpr const char *getName() { return ""; }
     static constexpr bool isESIMD() { return 0; }
+    static constexpr int64_t getKernelSize() { return 0; }
   };
   } // namespace detail
   } // namespace sycl

libclc/ptx-nvidiacl/libspirv/group/collectives.cl

@@ -249,13 +249,34 @@ __CLC_SUBGROUP_COLLECTIVE(FMax, __CLC_MAX, half, -HALF_MAX)
 __CLC_SUBGROUP_COLLECTIVE(FMax, __CLC_MAX, float, -FLT_MAX)
 __CLC_SUBGROUP_COLLECTIVE(FMax, __CLC_MAX, double, -DBL_MAX)
 
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseAndKHR, __CLC_AND, and, uchar, ~0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseOrKHR, __CLC_OR, or, uchar, 0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseXorKHR, __CLC_XOR, xor, uchar, 0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseAndKHR, __CLC_AND, and, char, ~0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseOrKHR, __CLC_OR, or, char, 0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseXorKHR, __CLC_XOR, xor, char, 0)
+
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseAndKHR, __CLC_AND, and, ushort, ~0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseOrKHR, __CLC_OR, or, ushort, 0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseXorKHR, __CLC_XOR, xor, ushort, 0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseAndKHR, __CLC_AND, and, short, ~0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseOrKHR, __CLC_OR, or, short, 0)
+__CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseXorKHR, __CLC_XOR, xor, short, 0)
+
 __CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseAndKHR, __CLC_AND, and, uint, ~0)
 __CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseOrKHR, __CLC_OR, or, uint, 0)
 __CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseXorKHR, __CLC_XOR, xor, uint, 0)
 __CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseAndKHR, __CLC_AND, and, int, ~0)
 __CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseOrKHR, __CLC_OR, or, int, 0)
 __CLC_SUBGROUP_COLLECTIVE_REDUX(BitwiseXorKHR, __CLC_XOR, xor, int, 0)
 
+__CLC_SUBGROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, ulong, ~0l)
+__CLC_SUBGROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, ulong, 0l)
+__CLC_SUBGROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, ulong, 0l)
+__CLC_SUBGROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, long, ~0l)
+__CLC_SUBGROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, long, 0l)
+__CLC_SUBGROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, long, 0l)
+
 #undef __CLC_SUBGROUP_COLLECTIVE_BODY
 #undef __CLC_SUBGROUP_COLLECTIVE
 #undef __CLC_SUBGROUP_COLLECTIVE_REDUX
@@ -376,13 +397,34 @@ __CLC_GROUP_COLLECTIVE(FMax, __CLC_MAX, half, -HALF_MAX)
 __CLC_GROUP_COLLECTIVE(FMax, __CLC_MAX, float, -FLT_MAX)
 __CLC_GROUP_COLLECTIVE(FMax, __CLC_MAX, double, -DBL_MAX)
 
+__CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, uchar, ~0)
+__CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, uchar, 0)
+__CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, uchar, 0)
+__CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, char, ~0)
+__CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, char, 0)
+__CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, char, 0)
+
+__CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, ushort, ~0)
+__CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, ushort, 0)
+__CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, ushort, 0)
+__CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, short, ~0)
+__CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, short, 0)
+__CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, short, 0)
+
 __CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, uint, ~0)
 __CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, uint, 0)
 __CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, uint, 0)
 __CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, int, ~0)
 __CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, int, 0)
 __CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, int, 0)
 
+__CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, ulong, ~0l)
+__CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, ulong, 0l)
+__CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, ulong, 0l)
+__CLC_GROUP_COLLECTIVE(BitwiseAndKHR, __CLC_AND, long, ~0l)
+__CLC_GROUP_COLLECTIVE(BitwiseOrKHR, __CLC_OR, long, 0l)
+__CLC_GROUP_COLLECTIVE(BitwiseXorKHR, __CLC_XOR, long, 0l)
+
 // half requires additional mangled entry points
 _CLC_DEF _CLC_CONVERGENT half _Z17__spirv_GroupFAddjjDF16_(uint scope, uint op,
                                                            half x) {

sycl/CMakeLists.txt

@@ -28,7 +28,7 @@ include(SYCLUtils)
 set(SYCL_MAJOR_VERSION 5)
 set(SYCL_MINOR_VERSION 7)
 set(SYCL_PATCH_VERSION 0)
-set(SYCL_DEV_ABI_VERSION 0)
+set(SYCL_DEV_ABI_VERSION 1)
 if (SYCL_ADD_DEV_VERSION_POSTFIX)
   set(SYCL_VERSION_POSTFIX "-${SYCL_DEV_ABI_VERSION}")
 endif()

sycl/doc/EnvironmentVariables.md

@@ -157,6 +157,7 @@ variables in production code.</span>
 | `SYCL_PI_LEVEL_ZERO_USE_COPY_ENGINE_FOR_D2D_COPY` (experimental) | Integer | Allows the use of copy engine, if available in the device, in Level Zero plugin for device to device copy operations. The default is 0. This option is experimental and will be removed once heuristics are added to make a decision about use of copy engine for device to device copy operations. |
 | `SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS` | Any(\*) | Enable support of device-scope events whose state is not visible to the host. If enabled mode is SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS=1 the Level Zero plugin would create all events having device-scope only and create proxy host-visible events for them when their status is needed (wait/query) on the host. If enabled mode is SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS=2 the Level Zero plugin would create all events having device-scope and add proxy host-visible event at the end of each command-list submission. The default is 2, meaning only the last event in a batch is host-visible. |
 | `SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS` | Integer | When set to a positive value enables use of Level Zero immediate commandlists, which means there is no batching and all commands are immediately submitted for execution. Default is 0. Note: When immediate commandlist usage is enabled it is necessary to also set SYCL_PI_LEVEL_ZERO_DEVICE_SCOPE_EVENTS to either 0 or 1.  |
+| `SYCL_PI_LEVEL_ZERO_USE_MULTIPLE_COMMANDLIST_BARRIERS` | Integer | When set to a positive value enables use of multiple Level Zero commandlists when submitting barriers. Default is 0. |
 
 ## Debugging variables for CUDA Plugin
 

sycl/include/CL/sycl/detail/kernel_desc.hpp

@@ -13,6 +13,8 @@
 #include <CL/sycl/detail/defines_elementary.hpp>
 #include <CL/sycl/detail/export.hpp>
 
+#include <cstdint>
+
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
@@ -79,6 +81,7 @@ template <class KernelNameType> struct KernelInfo {
   static constexpr const char *getFunctionName() { return ""; }
   static constexpr unsigned getLineNumber() { return 0; }
   static constexpr unsigned getColumnNumber() { return 0; }
+  static constexpr int64_t getKernelSize() { return 0; }
 };
 #else
 template <char...> struct KernelInfoData {
@@ -93,6 +96,7 @@ template <char...> struct KernelInfoData {
   static constexpr const char *getFunctionName() { return ""; }
   static constexpr unsigned getLineNumber() { return 0; }
   static constexpr unsigned getColumnNumber() { return 0; }
+  static constexpr int64_t getKernelSize() { return 0; }
 };
 
 // C++14 like index_sequence and make_index_sequence
@@ -135,6 +139,9 @@ template <class KernelNameType> struct KernelInfo {
   static constexpr const char *getFunctionName() { return ""; }
   static constexpr unsigned getLineNumber() { return 0; }
   static constexpr unsigned getColumnNumber() { return 0; }
+  static constexpr int64_t getKernelSize() {
+    return SubKernelInfo::getKernelSize();
+  }
 };
 #endif //__SYCL_UNNAMED_LAMBDA__
 

sycl/include/CL/sycl/handler.hpp

@@ -731,13 +731,25 @@ class __SYCL_EXPORT handler {
           "kernel_handler is not yet supported by host device.",
           PI_ERROR_INVALID_OPERATION);
     }
+
     KernelType *KernelPtr =
         ResetHostKernel<KernelType, LambdaArgType, Dims>(KernelFunc);
 
     using KI = sycl::detail::KernelInfo<KernelName>;
+    constexpr bool KernelHasName =
+        KI::getName() != nullptr && KI::getName()[0] != '\0';
+
+    // Some host compilers may have different captures from Clang. Currently
+    // there is no stable way of handling this when extracting the captures, so
+    // a static assert is made to fail for incompatible kernel lambdas.
+    static_assert(!KernelHasName || sizeof(KernelFunc) == KI::getKernelSize(),
+                  "Unexpected kernel lambda size. This can be caused by an "
+                  "external host compiler producing a lambda with an "
+                  "unexpected layout. This is a limitation of the compiler.");
+
     // Empty name indicates that the compilation happens without integration
     // header, so don't perform things that require it.
-    if (KI::getName() != nullptr && KI::getName()[0] != '\0') {
+    if (KernelHasName) {
       // TODO support ESIMD in no-integration-header case too.
       MArgs.clear();
       extractArgsAndReqsFromLambda(reinterpret_cast<char *>(KernelPtr),

sycl/include/sycl/ext/intel/experimental/online_compiler.hpp

@@ -219,13 +219,5 @@ online_compiler<source_language::cm>::compile(const std::string &src) {
 } // namespace experimental
 } // namespace intel
 } // namespace ext
-
-namespace ext {
-namespace __SYCL2020_DEPRECATED(
-    "use 'ext::intel::experimental' instead") intel {
-  using namespace ext::intel::experimental;
-} // namespace intel
-} // namespace ext
-
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)