Merge from 'sycl' to 'sycl-web'

Author: iclsrc

llvm/lib/SYCLLowerIR/LowerESIMD.cpp

@@ -789,7 +789,8 @@ static void translateUnPackMask(CallInst &CI) {
   if (Width > N) {
     llvm::Type *Ty = llvm::IntegerType::get(Context, N);
     Arg0 = Builder.CreateTrunc(Arg0, Ty);
-    cast<llvm::Instruction>(Arg0)->setDebugLoc(CI.getDebugLoc());
+    if (auto *Trunc = dyn_cast<llvm::Instruction>(Arg0))
+      Trunc->setDebugLoc(CI.getDebugLoc());
   }
   assert(Arg0->getType()->getPrimitiveSizeInBits() == N);
   Arg0 = Builder.CreateBitCast(

llvm/test/SYCLLowerIR/esimd_lower_crash_zext.ll

@@ -13,3 +13,13 @@ define spir_func <32 x i16> @_Z3foov() {
 }
 
 declare dso_local spir_func <32 x i16> @_Z19__esimd_unpack_maskILi32EEN2cl4sycl5INTEL3gpu11vector_typeItXT_EE4typeEj(i32)
+
+define spir_func <16 x i16> @_Z3barv() {
+; CHECK-LABEL: @_Z3barv(
+; CHECK:    ret <16 x i16> zext (<16 x i1> bitcast (<1 x i16> <i16 15> to <16 x i1>) to <16 x i16>)
+;
+  %call.i.i = call spir_func <16 x i16> @_Z19__esimd_unpack_maskILi16EEN2cl4sycl3ext5intel12experimental5esimd6detail11vector_typeItXT_EE4typeEj(i32 15)
+  ret <16 x i16> %call.i.i
+}
+
+declare dso_local spir_func <16 x i16> @_Z19__esimd_unpack_maskILi16EEN2cl4sycl3ext5intel12experimental5esimd6detail11vector_typeItXT_EE4typeEj(i32)

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -319,7 +319,7 @@ static void collectKernelModuleMap(
     }
     case Scope_Global:
       // the map key is not significant here
-      ResKernelModuleMap["<GLOBAL>"].push_back(&F);
+      ResKernelModuleMap[GLOBAL_SCOPE_NAME].push_back(&F);
       break;
     }
   }

sycl/doc/extensions/Uniform/Uniform.asciidoc

@@ -110,8 +110,28 @@ namespace experimental {
 
 template <class T>
 class uniform {
+public:
   explicit uniform(T x) noexcept;
-  operator T() const;
+  operator const T() const;
+
+  uniform& operator=(const uniform&) = delete;
+
+  /* Other explicitly deleted operators improve error messages
+     if a user incorrectly attempts to modify a uniform */
+  uniform& operator+=(const T&) = delete;
+  uniform& operator-=(const T&) = delete;
+  uniform& operator*=(const T&) = delete;
+  uniform& operator/=(const T&) = delete;
+  uniform& operator%=(const T&) = delete;
+  uniform& operator&=(const T&) = delete;
+  uniform& operator|=(const T&) = delete;
+  uniform& operator^=(const T&) = delete;
+  uniform& operator<<=(const T&) = delete;
+  uniform& operator>>=(const T&) = delete;
+  uniform& operator++() = delete;
+  uniform& operator++(int) = delete;
+  uniform& operator--() = delete;
+  uniform& operator--(int) = delete;
 };
 
 } // namespace experimental
@@ -181,6 +201,7 @@ float x = array[sycl::ext::oneapi::experimental::uniform(index)];
 |========================================
 |Rev|Date|Author|Changes
 |1|2021-04-23|John Pennycook|*Initial public working draft*
+|2|2021-08-03|John Pennycook|*Add const and deleted operators*
 |========================================
 
 //************************************************************************

sycl/include/CL/sycl/queue.hpp

@@ -22,7 +22,11 @@
 #include <CL/sycl/property_list.hpp>
 #include <CL/sycl/stl.hpp>
 
-#include <inttypes.h>
+// Explicitly request format macros
+#ifndef __STDC_FORMAT_MACROS
+#define __STDC_FORMAT_MACROS 1
+#endif
+#include <cinttypes>
 #include <utility>
 
 // having _TWO_ mid-param #ifdefs makes the functions very difficult to read.