Merge from 'sycl' to 'sycl-web'

Author: iclsrc

buildbot/dependency.conf

@@ -4,8 +4,8 @@ ocl_cpu_rt_ver=2021.11.3.0.09
 # https://github.com/intel/llvm/releases/download/2021-WW11/win-oclcpuexp-2021.11.3.0.09_rel.zip
 ocl_cpu_rt_ver_win=2021.11.3.0.09
 # Same GPU driver supports Level Zero and OpenCL
-# https://github.com/intel/compute-runtime/releases/tag/21.12.19358
-ocl_gpu_rt_ver=21.12.19358
+# https://github.com/intel/compute-runtime/releases/tag/21.13.19438
+ocl_gpu_rt_ver=21.13.19438
 # Same GPU driver supports Level Zero and OpenCL
 # https://downloadmirror.intel.com/30266/a08/igfx_win10_100.9316.zip
 ocl_gpu_rt_ver_win=27.20.100.9316
@@ -30,7 +30,7 @@ ocloc_ver_win=27.20.100.9168
 [DRIVER VERSIONS]
 cpu_driver_lin=2021.11.3.0.09
 cpu_driver_win=2021.11.3.0.09
-gpu_driver_lin=21.12.19358
+gpu_driver_lin=21.13.19438
 gpu_driver_win=27.20.100.9316
 fpga_driver_lin=2021.11.3.0.09
 fpga_driver_win=2021.11.3.0.09

llvm/lib/SYCLLowerIR/LowerESIMD.cpp

@@ -730,8 +730,7 @@ static bool translateVLoad(CallInst &CI, SmallPtrSet<Type *, 4> &GVTS) {
   if (GVTS.find(CI.getType()) != GVTS.end())
     return false;
   IRBuilder<> Builder(&CI);
-  auto *ElemT = CI.getArgOperand(0)->getType()->getPointerElementType();
-  auto LI = Builder.CreateLoad(ElemT, CI.getArgOperand(0), CI.getName());
+  auto LI = Builder.CreateLoad(CI.getType(), CI.getArgOperand(0), CI.getName());
   LI->setDebugLoc(CI.getDebugLoc());
   CI.replaceAllUsesWith(LI);
   return true;

llvm/lib/SYCLLowerIR/LowerESIMDVLoadVStore.cpp

@@ -89,8 +89,8 @@ PreservedAnalyses ESIMDLowerLoadStorePass::run(Function &F,
       if (GenXIntrinsic::isVStore(&Inst))
         Builder.CreateStore(Inst.getOperand(0), Inst.getOperand(1));
       else {
-        auto *ElemT = Inst.getOperand(0)->getType()->getPointerElementType();
-        auto LI = Builder.CreateLoad(ElemT, Inst.getOperand(0), Inst.getName());
+        auto LI = Builder.CreateLoad(Inst.getType(), Inst.getOperand(0),
+                                     Inst.getName());
         LI->setDebugLoc(Inst.getDebugLoc());
         Inst.replaceAllUsesWith(LI);
       }

llvm/lib/SYCLLowerIR/LowerWGScope.cpp

@@ -303,8 +303,7 @@ shareOutputViaLocalMem(Instruction &I, BasicBlock &BBa, BasicBlock &BBb,
   Bld.CreateStore(&I, WGLocal);
   // 3) Generate a load in the "worker" BB of the value stored by the leader
   Bld.SetInsertPoint(&BBb.front());
-  auto *WGValT = WGLocal->getType()->getPointerElementType();
-  auto *WGVal = Bld.CreateLoad(WGValT, WGLocal, "wg_val_" + Twine(I.getName()));
+  auto *WGVal = Bld.CreateLoad(T, WGLocal, "wg_val_" + Twine(I.getName()));
   // 4) Finally, replace usages of I outside the scope
   for (auto *U : Users)
     U->replaceUsesOfWith(&I, WGVal);
@@ -417,8 +416,7 @@ static void copyBetweenPrivateAndShadow(Value *L, GlobalVariable *Shadow,
 
     if (!Loc2Shadow)
       std::swap(Src, Dst);
-    auto *SrcT = Src->getType()->getPointerElementType();
-    Value *LocalVal = Builder.CreateLoad(SrcT, Src, "mat_ld");
+    Value *LocalVal = Builder.CreateLoad(T, Src, "mat_ld");
     Builder.CreateStore(LocalVal, Dst);
   }
 }

sycl/include/CL/sycl/INTEL/esimd/esimd.hpp

@@ -29,6 +29,8 @@ using namespace sycl::INTEL::gpu::detail;
 ///
 /// \ingroup sycl_esimd
 template <typename Ty, int N> class simd {
+  template <typename, typename> friend class simd_view;
+
 public:
   /// The underlying builtin data type.
   using vector_type = detail::vector_type_t<Ty, N>;

sycl/include/CL/sycl/INTEL/esimd/esimd_view.hpp

@@ -24,6 +24,9 @@ namespace gpu {
 ///
 /// \ingroup sycl_esimd
 template <typename BaseTy, typename RegionTy> class simd_view {
+  template <typename, int> friend class simd;
+  template <typename, typename> friend class simd_view;
+
 public:
   static_assert(!is_simd_view_v<BaseTy>::value);
   // Deduce the corresponding value type from its region type.
@@ -44,31 +47,17 @@ template <typename BaseTy, typename RegionTy> class simd_view {
   /// type of the base object type.
   using element_type = typename ShapeTy::element_type;
 
-  // TODO @rolandschulz
-  // {quote}
-  // Why is this and the next constructor public ? Those should only be called
-  // internally by e.g.select, correct ?
-  // {/quote}
-  //
   /// @{
   /// Constructors.
+
+private:
   simd_view(BaseTy &Base, RegionTy Region) : M_base(Base), M_region(Region) {}
   simd_view(BaseTy &&Base, RegionTy Region) : M_base(Base), M_region(Region) {}
 
-  // TODO @rolandschulz
-  // {quote}
-  // Is this intentional not a correct copy constructor (would need to be const
-  // for that)? I believe we agreed that simd_view would have a deleted copy and
-  // move constructor.Why are they suddenly back ?
-  // {/quote}
-  // TODO @kbobrovs
-  // copy constructor is still incorrect (no 'const'), move constructor is still
-  // present.
-  //
-  // Disallow copy constructor for simd_view.
-  simd_view(simd_view &Other) = delete;
-  simd_view(simd_view &&Other)
-      : M_base(Other.M_base), M_region(Other.M_region) {}
+public:
+  // Disallow copy and move constructors for simd_view.
+  simd_view(const simd_view &Other) = delete;
+  simd_view(simd_view &&Other) = delete;
   /// @}
 
   /// Conversion to simd type.

sycl/include/CL/sycl/ONEAPI/intel_matrix/matrix.hpp

@@ -14,6 +14,10 @@
 
 #pragma once
 
+#include <CL/sycl/feature_test.hpp>
+
+#if (SYCL_EXT_ONEAPI_MATRIX == 1)
 #if defined(__AMXTILE__) && defined(__AMXINT8__) && defined(__AMXBF16__)
 #include <CL/sycl/ONEAPI/intel_matrix/matrix-amx.hpp>
 #endif
+#endif

sycl/include/CL/sycl/feature_test.hpp

@@ -12,7 +12,9 @@ namespace sycl {
 
 // Feature test macro definitions
 
+// TODO: Move these feature-test macros to compiler driver.
 #define SYCL_EXT_INTEL_DEVICE_INFO 1
+#define SYCL_EXT_ONEAPI_MATRIX 1
 
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/plugins/cuda/pi_cuda.cpp

@@ -3850,9 +3850,10 @@ static size_t imageElementByteSize(CUDA_ARRAY_DESCRIPTOR array_desc) {
   case CU_AD_FORMAT_SIGNED_INT32:
   case CU_AD_FORMAT_FLOAT:
     return 4;
+  default:
+    cl::sycl::detail::pi::die("Invalid image format.");
+    return 0;
   }
-  cl::sycl::detail::pi::die("Invalid iamge format.");
-  return 0;
 }
 
 /// General ND memory copy operation for images (where N > 1).

sycl/test/esimd/simd_view.cpp

@@ -7,18 +7,6 @@
 
 using namespace sycl::INTEL::gpu;
 
-bool test_simd_view_ctors() __attribute__((sycl_device)) {
-  simd<int, 16> v0(0, 1);
-
-  region1d_t<int, 4, 1> r0(4);
-  simd_view<simd<int, 16>, region1d_t<int, 4, 1>> ref0(v0, r0);
-  simd_view<simd<int, 16>, region1d_t<int, 4, 1>> ref1(std::move(ref0));
-  simd_view<simd<int, 16>, region1d_t<int, 4, 1>> ref2(
-      std::move(v0.select<4, 1>(8)));
-
-  return (ref0[0] == 4) && (ref1[1] == 5) && (ref2[0] == 8);
-}
-
 bool test_simd_view_bin_ops() __attribute__((sycl_device)) {
   simd<int, 16> v0 = 1;
   simd<int, 16> v1 = 2;