Merge remote-tracking branch sycl into win_fpga_lit

Author: vladimirlaz

buildbot/dependency.conf

@@ -4,15 +4,15 @@ ocl_cpu_rt_ver=2020.10.7.0.15
 # https://github.com/intel/llvm/releases/download/2020-WW31/win-oclcpuexp-2020.10.7.0.15_rel.zip
 ocl_cpu_rt_ver_win=2020.10.7.0.15
 # Same GPU driver supports Level Zero and OpenCL:
-# https://github.com/intel/compute-runtime/releases/tag/20.25.17111
+# https://github.com/intel/compute-runtime/releases/tag/20.29.17408
 ocl_gpu_rt_ver=20.29.17408
 # Same GPU driver supports Level Zero and OpenCL:
 # https://downloadmirror.intel.com/29674/a08/igfx_win10_100.8336.zip
 ocl_gpu_rt_ver_win=27.20.100.8336
 intel_sycl_ver=build
-# https://github.com/oneapi-src/oneTBB/releases/download/v2020.2/tbb-2020.2-lin.tgz
+# TODO provide URL for Linux TBB driver
 tbb_ver=2021.1.8.515
-# https://github.com/oneapi-src/oneTBB/releases/download/v2020.2/tbb-2020.2-win.zip
+# TODO provide URL for Windows TBB driver
 tbb_ver_win=2021.1.8.514
 # https://github.com/intel/llvm/releases/download/2020-WW31/fpgaemu-2020.10.7.0.15_rel.tar.gz
 ocl_fpga_emu_ver=2020.10.7.0.15

llvm-spirv/lib/SPIRV/SPIRVWriter.cpp

@@ -563,10 +563,11 @@ SPIRVFunction *LLVMToSPIRV::transFunctionDecl(Function *F) {
       // Order of integer numbers in MD node follows the order of function
       // parameters on which we shall attach the appropriate decoration. Add
       // decoration only if MD value is not negative.
-      BM->addCapability(CapabilityFPGABufferLocationINTEL);
       int LocID = getMDOperandAsInt(BufferLocation, ArgNo);
-      if (LocID >= 0)
+      if (LocID >= 0) {
+        BM->addCapability(CapabilityFPGABufferLocationINTEL);
         BA->addDecorate(DecorationBufferLocationINTEL, LocID);
+      }
     }
   }
   if (Attrs.hasAttribute(AttributeList::ReturnIndex, Attribute::ZExt))

sycl/CMakeLists.txt

@@ -107,6 +107,9 @@ if( NOT OpenCL_INCLUDE_DIRS )
     INSTALL_COMMAND   ""
     STEP_TARGETS      build
     COMMENT           "Downloading OpenCL headers."
+    LOG_DOWNLOAD 1
+    LOG_UPDATE 1
+    LOG_BUILD 1
   )
   add_definitions(-DCL_TARGET_OPENCL_VERSION=220)
 else()
@@ -154,6 +157,11 @@ if( NOT OpenCL_LIBRARIES )
     STEP_TARGETS      configure,build,install
     DEPENDS           ocl-headers
     BUILD_BYPRODUCTS ${OpenCL_LIBRARIES}
+    LOG_DOWNLOAD 1
+    LOG_UPDATE 1
+    LOG_CONFIGURE 1
+    LOG_BUILD 1
+    LOG_INSTALL 1
   )
   ExternalProject_Add_Step(ocl-icd llvminstall
     COMMAND ${CMAKE_COMMAND} -E copy_directory <INSTALL_DIR>/ ${LLVM_BINARY_DIR}

sycl/doc/extensions/GroupAlgorithms/SYCL_INTEL_group_algorithms.asciidoc

@@ -51,9 +51,7 @@ John Pennycook, Intel (john 'dot' pennycook 'at' intel 'dot' com)
 
 == Dependencies
 
-This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-6 and the following extensions:
-
-- +SYCL_INTEL_extended_atomics+
+This extension is written against the SYCL 1.2.1 specification, Revision v1.2.1-6.
 
 == Overview
 
@@ -69,10 +67,6 @@ The extension introduces the following functions:
 - +reduce+
 - +exclusive_scan+
 - +inclusive_scan+
-- +barrier+
-
-The definitions and behavior of the following functions are based on equivalents in the SYCL 2020 provisional specification:
-- +barrier+
 
 === Alignment with OpenCL vs C++
 
@@ -258,21 +252,6 @@ The return types of the collective functions in {cpp}17 are not deduced from the
 |Perform an inclusive scan over the values in the range [_first_, _last_) using the operator _binary_op_, which must be one of the group algorithms library function objects.  The value written to +result + i+ is the inclusive scan of the first +i+ values in the range and an initial value specified by _init_. Returns a pointer to the end of the output range. _first_, _last_, _result_, _binary_op_ and _init_ must be the same for all work-items in the group. _binary_op(init, *first)_ must return a value of type _T_.
 |===
 
-==== Synchronization
-
-The behavior of memory fences in this section is aligned with the single happens-before relationship defined by the +SYCL_INTEL_extended_atomics+ extension.
-
-|===
-|Function|Description
-
-|+template <typename Group> void barrier(Group g);+
-|Synchronize all work-items in the group, and ensure that all memory accesses to any address space prior to the barrier are visible to all work-items in the group after the barrier.  The scope of the group memory fences implied by this barrier is the narrowest scope including all work-items in the group.
-
-|+template <typename Group> void barrier(Group g, memory_scope scope);+
-|Synchronize all work-items in the group, and ensure that all memory accesses to any address space prior to the barrier are visible to all work-items specified by _scope_ after the barrier.  The scope of the group memory fences implied by this barrier is controlled by _scope_ and must be broader than the narrowest scope including all work-items in the group.  If the specified _scope_ is narrower than the narrowest scope including all work-items in the group, the _scope_ argument is ignored.
-
-|===
-
 == Issues
 
 None.
@@ -291,7 +270,6 @@ None.
 |========================================
 |Rev|Date|Author|Changes
 |1|2020-01-30|John Pennycook|*Initial public working draft*
-|2|2020-07-28|John Pennycook|*Add group barrier*
 |========================================
 
 //************************************************************************

sycl/doc/extensions/SubGroupAlgorithms/SYCL_INTEL_sub_group_algorithms.asciidoc

@@ -70,7 +70,6 @@ This extension adds sub-group support to all of the functions from +SYCL_INTEL_g
 - +reduce+
 - +exclusive_scan+
 - +inclusive_scan+
-- +barrier+
 
 It additionally introduces a number of functions that are currently specific to sub-groups:
 
@@ -166,7 +165,6 @@ None.
 |========================================
 |Rev|Date|Author|Changes
 |1|2020-03-16|John Pennycook|*Initial public working draft*
-|2|2020-07-28|John Pennycook|*Add group barrier*
 |========================================
 
 //************************************************************************

sycl/include/CL/sycl/intel/group_algorithm.hpp

@@ -13,7 +13,6 @@
 #include <CL/sycl/detail/spirv.hpp>
 #include <CL/sycl/detail/type_traits.hpp>
 #include <CL/sycl/group.hpp>
-#include <CL/sycl/intel/atomic.hpp>
 #include <CL/sycl/intel/functional.hpp>
 #include <CL/sycl/intel/sub_group.hpp>
 
@@ -78,15 +77,6 @@ template <> inline id<3> linear_id_to_id(range<3> r, size_t linear_id) {
   return result;
 }
 
-// TODO: Replace with Group::fence_scope from SYCL 2020 provisional
-template <typename Group> struct FenceScope {
-  static constexpr intel::memory_scope value = intel::memory_scope::work_group;
-};
-
-template <> struct FenceScope<intel::sub_group> {
-  static constexpr intel::memory_scope value = intel::memory_scope::sub_group;
-};
-
 template <typename T, class BinaryOperation> struct identity {};
 
 template <typename T, typename V> struct identity<T, intel::plus<V>> {
@@ -906,42 +896,6 @@ template <typename Group> bool leader(Group g) {
 #endif
 }
 
-template <typename Group> void barrier(Group, memory_scope scope) {
-  static_assert(sycl::detail::is_generic_group<Group>::value,
-                "Group algorithms only support the sycl::group and "
-                "intel::sub_group class.");
-#ifdef __SYCL_DEVICE_ONLY__
-  // MemoryScope must be broader than Group scope for correctness
-  auto GroupScope = detail::FenceScope<Group>::value;
-  auto BroadestScope = (scope > GroupScope) ? scope : GroupScope;
-  auto MemoryScope = sycl::detail::spirv::getScope(BroadestScope);
-  auto ExecutionScope = sycl::detail::spirv::group_scope<Group>::value;
-  __spirv_ControlBarrier(ExecutionScope, MemoryScope,
-                         __spv::MemorySemanticsMask::AcquireRelease |
-                             __spv::MemorySemanticsMask::SubgroupMemory |
-                             __spv::MemorySemanticsMask::WorkgroupMemory |
-                             __spv::MemorySemanticsMask::CrossWorkgroupMemory);
-#else
-  (void)scope;
-  throw runtime_error("Group algorithms are not supported on host device.",
-                      PI_INVALID_DEVICE);
-#endif
-}
-
-template <typename Group> void barrier(Group g) {
-  static_assert(sycl::detail::is_generic_group<Group>::value,
-                "Group algorithms only support the sycl::group and "
-                "intel::sub_group class.");
-#ifdef __SYCL_DEVICE_ONLY__
-  auto MemoryScope = detail::FenceScope<Group>::value;
-  barrier(g, MemoryScope);
-#else
-  (void)g;
-  throw runtime_error("Group algorithms are not supported on host device.",
-                      PI_INVALID_DEVICE);
-#endif
-}
-
 } // namespace intel
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/plugins/level_zero/CMakeLists.txt

@@ -33,6 +33,11 @@ if (NOT DEFINED LEVEL_ZERO_LIBRARY OR NOT DEFINED LEVEL_ZERO_INCLUDE_DIR)
                -DOpenCL_INCLUDE_DIR=${OpenCL_INCLUDE_DIRS}
                -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR>
                -DCMAKE_INSTALL_LIBDIR:PATH=lib${LLVM_LIBDIR_SUFFIX}
+               LOG_DOWNLOAD 1
+               LOG_UPDATE 1
+               LOG_CONFIGURE 1
+               LOG_BUILD 1
+               LOG_INSTALL 1
                ${AUX_CMAKE_FLAGS}
     STEP_TARGETS      configure,build,install
     DEPENDS           ocl-headers

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -3123,7 +3123,8 @@ pi_result piEnqueueMemUnmap(pi_queue Queue, pi_mem MemObj, void *MappedPtr,
   // piEnqueueMemBufferMap, but can only do so after the above copy
   // is completed. Instead of waiting for It here (blocking), we shall
   // do so in piEventRelease called for the pi_event tracking the unmap.
-  (*Event)->CommandData = MemObj->MapHostPtr ? nullptr : MappedPtr;
+  if (Event)
+    (*Event)->CommandData = MemObj->MapHostPtr ? nullptr : MappedPtr;
 
   // Execute command list asynchronously, as the event will be used
   // to track down its completion.

sycl/test/group-algorithm/barrier.cpp

@@ -36,7 +36,7 @@
 config.test_exec_root = os.path.join(config.sycl_obj_root, 'test')
 
 # Propagate some variables from the host environment.
-llvm_config.with_system_environment(['PATH', 'OCL_ICD_FILENAME', 'SYCL_DEVICE_ALLOWLIST', 'SYCL_CONFIG_FILE_NAME'])
+llvm_config.with_system_environment(['PATH', 'OCL_ICD_FILENAMES', 'SYCL_DEVICE_ALLOWLIST', 'SYCL_CONFIG_FILE_NAME'])
 
 # Configure LD_LIBRARY_PATH or corresponding os-specific alternatives
 if platform.system() == "Linux":

sycl/test/lit.cfg.py

@@ -8,12 +8,10 @@
 
 #include "helper.hpp"
 #include <CL/sycl.hpp>
-template <typename T>
-class sycl_subgr;
+template <typename T> class sycl_subgr;
 using namespace cl::sycl;
-template <typename T>
-void check(queue &Queue) {
-  const int G = 240, L = 60;
+template <typename T> void check(queue &Queue) {
+  const int G = 256, L = 64;
   try {
     nd_range<1> NdRange(G, L);
     buffer<T> syclbuf(G);
@@ -23,9 +21,10 @@ void check(queue &Queue) {
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
       cgh.parallel_for<sycl_subgr<T>>(NdRange, [=](nd_item<1> NdItem) {
         intel::sub_group SG = NdItem.get_sub_group();
-        /*Broadcast GID of element with SGLID == SGID */
+        /*Broadcast GID of element with SGLID == SGID % SGMLR*/
         syclacc[NdItem.get_global_id()] =
-            broadcast(SG, T(NdItem.get_global_id(0)), SG.get_group_id());
+            broadcast(SG, T(NdItem.get_global_id(0)),
+                      SG.get_group_id() % SG.get_max_local_range()[0]);
         if (NdItem.get_global_id(0) == 0)
           sgsizeacc[0] = SG.get_max_local_range()[0];
       });
@@ -44,7 +43,8 @@ void check(queue &Queue) {
         WGid++;
         SGid = 0;
       }
-      exit_if_not_equal<T>(syclacc[j], L * WGid + SGid + SGid * sg_size,
+      exit_if_not_equal<T>(syclacc[j],
+                           L * WGid + SGid % sg_size + SGid * sg_size,
                            "broadcasted value");
     }
   } catch (exception e) {

sycl/test/sub_group/broadcast.hpp

@@ -1,6 +1,5 @@
-// UNSUPPORTED: cuda || cpu
+// UNSUPPORTED: cuda
 // CUDA compilation and runtime do not yet support sub-groups.
-// #2245 failed on OpenCL CPU (2020.10.7.0.15) with avx2 instruction set
 //
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out
@@ -24,7 +23,7 @@ template <typename T> class pointer_kernel;
 using namespace cl::sycl;
 
 template <typename T>
-void check_pointer(queue &Queue, size_t G = 240, size_t L = 60) {
+void check_pointer(queue &Queue, size_t G = 256, size_t L = 64) {
   try {
     nd_range<1> NdRange(G, L);
     buffer<T *> buf(G);
@@ -118,7 +117,7 @@ void check_pointer(queue &Queue, size_t G = 240, size_t L = 60) {
 }
 
 template <typename T, typename Generator>
-void check_struct(queue &Queue, Generator &Gen, size_t G = 240, size_t L = 60) {
+void check_struct(queue &Queue, Generator &Gen, size_t G = 256, size_t L = 64) {
 
   // Fill a vector with values that will be shuffled
   std::vector<T> values(G);

sycl/test/sub_group/generic-shuffle.cpp

@@ -1,5 +1,7 @@
-// UNSUPPORTED: cuda
+// UNSUPPORTED: cuda || cpu
 // CUDA compilation and runtime do not yet support sub-groups.
+// #2252 Disable until all variants of built-ins are available in OpenCL CPU
+// runtime for every supported ISA
 //
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out

sycl/test/sub_group/load_store.cpp

@@ -1,5 +1,7 @@
-// UNSUPPORTED: cuda
+// UNSUPPORTED: cuda || cpu
 // CUDA compilation and runtime do not yet support sub-groups.
+// #2252 Disable until all variants of built-ins are available in OpenCL CPU
+// runtime for every supported ISA
 //
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out

sycl/test/sub_group/reduce.cpp

@@ -9,14 +9,13 @@
 #include "helper.hpp"
 #include <CL/sycl.hpp>
 
-template <typename T, class BinaryOperation>
-class sycl_subgr;
+template <typename T, class BinaryOperation> class sycl_subgr;
 
 using namespace cl::sycl;
 
 template <typename T, class BinaryOperation>
 void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
-              size_t G = 240, size_t L = 60) {
+              size_t G = 256, size_t L = 64) {
   try {
     nd_range<1> NdRange(G, L);
     buffer<T> buf(G);
@@ -65,8 +64,7 @@ void check_op(queue &Queue, T init, BinaryOperation op, bool skip_init = false,
   }
 }
 
-template <typename T>
-void check(queue &Queue, size_t G = 240, size_t L = 60) {
+template <typename T> void check(queue &Queue, size_t G = 256, size_t L = 64) {
   // limit data range for half to avoid rounding issues
   if (std::is_same<T, cl::sycl::half>::value) {
     G = 64;

sycl/test/sub_group/reduce.hpp

@@ -1,5 +1,7 @@
-// UNSUPPORTED: cuda
+// UNSUPPORTED: cuda || cpu
 // CUDA compilation and runtime do not yet support sub-groups.
+// #2252 Disable until all variants of built-ins are available in OpenCL CPU
+// runtime for every supported ISA
 //
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_DEVICE_TYPE=HOST %t.out