Merge from 'sycl' to 'sycl-web' (#2)

iclsrc · iclsrc · commit 12e4dc710aba · 2020-11-11T04:28:39.000-08:00
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
@@ -1236,6 +1236,19 @@ def SYCLIntelNumSimdWorkItems : InheritableAttr {
   let PragmaAttributeSupport = 0;
 }
 
+def SYCLIntelStallEnable : InheritableAttr {
+  let Spellings = [CXX11<"intel","stall_enable">];
+  let LangOpts = [SYCLIsHost, SYCLIsDevice];
+  let Subjects = SubjectList<[Function], ErrorDiag>;
+  let AdditionalMembers = [{
+    static const char *getName() {
+      return "stall_enable";
+    }
+  }];
+  let Documentation = [SYCLIntelStallEnableAttrDocs];
+  let PragmaAttributeSupport = 0;
+}
+
 def SYCLIntelSchedulerTargetFmaxMhz : InheritableAttr {
   let Spellings = [CXX11<"intelfpga","scheduler_target_fmax_mhz">,
                    CXX11<"intel","scheduler_target_fmax_mhz">];
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
@@ -2229,6 +2229,39 @@ device kernel, the attribute is ignored and it is not propagated to a kernel.
   }];
 }
 
+def SYCLIntelStallEnableAttrDocs : Documentation {
+  let Category = DocCatFunction;
+  let Heading = "intel::stall_enable";
+  let Content = [{
+When applied to a lambda or function call operator (of a function object)
+on device, this requests, to the extent possible, that statically-scheduled
+clusters handle stalls using a stall-enable signal to freeze computation
+within the cluster. This attribute is ignored on the host.
+
+If ``intel::stall_enable`` is applied to a function called from a device
+kernel, the attribute is ignored and it is not propagated to a kernel.
+
+The ``intel::stall_enable`` attribute takes no argument and has an effect
+when applied to a function, and no effect otherwise.
+
+.. code-block:: c++
+
+  class Functor
+  {
+      [[intel::stall_enable]] void operator()(item<1> item)
+      {
+          /* kernel code */
+      }
+  }
+
+  kernel<class kernel_name>(
+  []() [[intel::stall_enable]] {
+       /* kernel code */
+  });
+
+  }];
+}
+
 def ReqdWorkGroupSizeAttrDocs : Documentation {
   let Category = DocCatFunction;
   let Heading = "reqd_work_group_size";
diff --git a/clang/include/clang/Basic/AttributeCommonInfo.h b/clang/include/clang/Basic/AttributeCommonInfo.h
@@ -165,7 +165,8 @@ class AttributeCommonInfo {
         ParsedAttr == AT_SYCLIntelSchedulerTargetFmaxMhz ||
         ParsedAttr == AT_SYCLIntelMaxWorkGroupSize ||
         ParsedAttr == AT_SYCLIntelMaxGlobalWorkDim ||
-        ParsedAttr == AT_SYCLIntelNoGlobalWorkOffset)
+        ParsedAttr == AT_SYCLIntelNoGlobalWorkOffset ||
+        ParsedAttr == AT_SYCLIntelStallEnable)
       return true;
 
     return false;
diff --git a/clang/lib/CodeGen/CodeGenFunction.cpp b/clang/lib/CodeGen/CodeGenFunction.cpp
@@ -693,6 +693,12 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
     if (A->getEnabled())
       Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {}));
   }
+
+  if (FD->hasAttr<SYCLIntelStallEnableAttr>()) {
+    llvm::Metadata *AttrMDArgs[] = {
+        llvm::ConstantAsMetadata::get(Builder.getInt32(1))};
+    Fn->setMetadata("stall_enable", llvm::MDNode::get(Context, AttrMDArgs));
+  }
 }
 
 /// Determine whether the function F ends with a return stmt.
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -3029,6 +3029,20 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D,
                                                                      E);
 }
 
+// Handles stall_enable
+static void handleStallEnableAttr(Sema &S, Decl *D, const ParsedAttr &Attr) {
+  if (D->isInvalidDecl())
+    return;
+
+  unsigned NumArgs = Attr.getNumArgs();
+  if (NumArgs > 0) {
+    S.Diag(Attr.getLoc(), diag::warn_attribute_too_many_arguments) << Attr << 0;
+    return;
+  }
+
+  handleSimpleAttribute<SYCLIntelStallEnableAttr>(S, D, Attr);
+}
+
 // Add scheduler_target_fmax_mhz
 void Sema::addSYCLIntelSchedulerTargetFmaxMhzAttr(
     Decl *D, const AttributeCommonInfo &Attr, Expr *E) {
@@ -8401,6 +8415,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
     handleNoGlobalWorkOffsetAttr(S, D, AL);
     break;
+  case ParsedAttr::AT_SYCLIntelStallEnable:
+    handleStallEnableAttr(S, D, AL);
+    break;
   case ParsedAttr::AT_VecTypeHint:
     handleVecTypeHint(S, D, AL);
     break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
@@ -545,6 +545,20 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
 
       if (auto *A = FD->getAttr<SYCLSimdAttr>())
         Attrs.insert(A);
+
+      // Allow the kernel attribute "stall_enable" only on lambda functions
+      // and function objects that are called directly from a kernel
+      // (i.e. the one passed to the single_task or parallel_for functions).
+      // For all other cases, emit a warning and ignore.
+      if (auto *A = FD->getAttr<SYCLIntelStallEnableAttr>()) {
+        if (ParentFD == SYCLKernel) {
+          Attrs.insert(A);
+        } else {
+          SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
+          FD->dropAttr<SYCLIntelStallEnableAttr>();
+        }
+      }
+
       // Propagate the explicit SIMD attribute through call graph - it is used
       // to distinguish ESIMD code in ESIMD LLVM passes.
       if (KernelBody && KernelBody->hasAttr<SYCLSimdAttr>() &&
@@ -3222,6 +3236,7 @@ void Sema::MarkDevice(void) {
         case attr::Kind::SYCLIntelSchedulerTargetFmaxMhz:
         case attr::Kind::SYCLIntelMaxGlobalWorkDim:
         case attr::Kind::SYCLIntelNoGlobalWorkOffset:
+        case attr::Kind::SYCLIntelStallEnable:
         case attr::Kind::SYCLSimd: {
           if ((A->getKind() == attr::Kind::SYCLSimd) && KernelBody &&
               !KernelBody->getAttr<SYCLSimdAttr>()) {
diff --git a/clang/test/CodeGenSYCL/stall_enable.cpp b/clang/test/CodeGenSYCL/stall_enable.cpp
@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+class Foo {
+public:
+  [[intel::stall_enable]] void operator()() const {}
+};
+
+int main() {
+  q.submit([&](handler &h) {
+    Foo f;
+    h.single_task<class test_kernel1>(f);
+
+    h.single_task<class test_kernel2>(
+        []() [[intel::stall_enable]]{});
+  });
+  return 0;
+}
+
+// CHECK: define spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !stall_enable ![[NUM5:[0-9]+]]
+// CHECK: define spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} !stall_enable ![[NUM5]]
+// CHECK: ![[NUM5]] = !{i32 1}
diff --git a/clang/test/SemaSYCL/stall_enable.cpp b/clang/test/SemaSYCL/stall_enable.cpp
@@ -0,0 +1,38 @@
+// RUN: %clang_cc1 %s -fsyntax-only -fsycl -internal-isystem %S/Inputs -fsycl-is-device -Wno-sycl-2017-compat -DTRIGGER_ERROR -verify
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -Wno-sycl-2017-compat %s | FileCheck %s
+
+#include "sycl.hpp"
+
+using namespace cl::sycl;
+queue q;
+
+[[intel::stall_enable]] void test() {} //expected-warning{{'stall_enable' attribute ignored}}
+
+#ifdef TRIGGER_ERROR
+[[intel::stall_enable(1)]] void bar1() {} // expected-error{{'stall_enable' attribute takes no arguments}}
+[[intel::stall_enable]] int N;            // expected-error{{'stall_enable' attribute only applies to functions}}
+#endif
+
+struct FuncObj {
+  [[intel::stall_enable]] void operator()() const {}
+};
+
+int main() {
+  q.submit([&](handler &h) {
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
+    // CHECK:       SYCLIntelStallEnableAttr {{.*}}
+    h.single_task<class test_kernel1>(
+        FuncObj());
+
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
+    // CHECK:       SYCLIntelStallEnableAttr {{.*}}
+    h.single_task<class test_kernel2>(
+        []() [[intel::stall_enable]]{});
+
+    // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
+    // CHECK-NOT:   SYCLIntelStallEnableAttr {{.*}}
+    h.single_task<class test_kernel3>(
+        []() { test(); });
+  });
+  return 0;
+}
diff --git a/llvm/lib/SYCLLowerIR/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/LowerESIMD.cpp
@@ -252,6 +252,15 @@ class ESIMDIntrinDescTable {
         {"flat_write", {"svm.scatter", {ai1(3), a(2), a(0), a(1)}}},
         {"flat_write4",
          {"svm.scatter4.scaled", {ai1(2), t(2), c16(0), c64(0), a(0), a(1)}}},
+
+        // surface index-based gather/scatter:
+        // num blocks, scale, surface index, global offset, elem offsets
+        {"surf_read", {"gather.scaled2", {t(3), c16(0), aSI(1), a(2), a(3)}}},
+        // pred, num blocks, scale, surface index, global offset, elem offsets,
+        // data to write
+        {"surf_write",
+         {"scatter.scaled", {ai1(0), t(3), c16(0), aSI(2), a(3), a(4), a(5)}}},
+
         // intrinsics to query thread's coordinates:
         {"group_id_x", {"group.id.x", {}}},
         {"group_id_y", {"group.id.y", {}}},
diff --git a/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp b/sycl/include/CL/sycl/INTEL/esimd/detail/esimd_memory_intrin.hpp
@@ -81,6 +81,92 @@ SYCL_EXTERNAL void __esimd_flat_write4(
     sycl::INTEL::gpu::vector_type_t<Ty, N * NumChannels(Mask)> vals,
     sycl::INTEL::gpu::vector_type_t<uint16_t, N> pred = 1);
 
+// Low-level surface-based gather. Collects elements located at given offsets in
+// a surface and returns them as a single \ref simd object. Element can be
+// 1, 2 or 4-byte value, but is always returned as a 4-byte value within the
+// resulting simd object, with upper 2 or 3 bytes undefined.
+// Template (compile-time constant) parameters:
+// @tparam Ty - element type; can only be a 4-byte integer or \c float,
+// @tparam N  - the number of elements
+// @tparam SurfIndAliasTy - "surface index alias" type - internal type in the
+//   accessor used to denote the surface
+// @tparam TySizeLog2 - Log2 of the number of bytes read per element:
+//   0 - 1 byte, 1 - 2 bytes, 2 - 4 bytes
+// @tparam L1H - L1 cache hint
+// @tparam L3H - L3 cache hint
+//
+// Formal parameters:
+// @param scale - the scale; must be 0
+// @param surf_ind - the surface index, taken from the SYCL memory object
+// @param global_offset - offset added to each individual element's offset to
+//   compute actual memory access offset for that element
+// @param elem_offsets - per-element offsets
+//
+template <typename Ty, int N, typename SurfIndAliasTy, int TySizeLog2,
+          sycl::INTEL::gpu::CacheHint L1H = sycl::INTEL::gpu::CacheHint::None,
+          sycl::INTEL::gpu::CacheHint L3H = sycl::INTEL::gpu::CacheHint::None>
+SYCL_EXTERNAL sycl::INTEL::gpu::vector_type_t<Ty, N>
+__esimd_surf_read(int16_t scale, SurfIndAliasTy surf_ind,
+                  uint32_t global_offset,
+                  sycl::INTEL::gpu::vector_type_t<uint32_t, N> elem_offsets)
+#ifdef __SYCL_DEVICE_ONLY__
+    ;
+#else
+{
+  static_assert(N == 1 || N == 8 || N == 16);
+  static_assert(TySizeLog2 <= 2);
+  static_assert(std::is_integral<Ty>::value || TySizeLog2 == 2);
+  throw cl::sycl::feature_not_supported();
+}
+#endif // __SYCL_DEVICE_ONLY__
+
+// Low-level surface-based scatter. Writes elements of a \ref simd object into a
+// surface at given offsets. Element can be a 1, 2 or 4-byte value, but it is
+// always represented as a 4-byte value within the input simd object,
+// unused (not written) upper bytes are ignored.
+// Template (compile-time constant) parameters:
+// @tparam Ty - element type; can only be a 4-byte integer or \c float,
+// @tparam N  - the number of elements to write
+// @tparam SurfIndAliasTy - "surface index alias" type - internal type in the
+//   accessor used to denote the surface
+// @tparam TySizeLog2 - Log2 of the number of bytes written per element:
+//   0 - 1 byte, 1 - 2 bytes, 2 - 4 bytes
+// @tparam L1H - L1 cache hint
+// @tparam L3H - L3 cache hint
+//
+// Formal parameters:
+// @param pred - per-element predicates; elements with zero corresponding
+//   predicates are not written
+// @param scale - the scale; must be 0
+// @param surf_ind - the surface index, taken from the SYCL memory object
+// @param global_offset - offset added to each individual element's offset to
+//   compute actual memory access offset for that element
+// @param elem_offsets - per-element offsets
+// @param vals - values to write
+//
+template <typename Ty, int N, typename SurfIndAliasTy, int TySizeLog2,
+          sycl::INTEL::gpu::CacheHint L1H = sycl::INTEL::gpu::CacheHint::None,
+          sycl::INTEL::gpu::CacheHint L3H = sycl::INTEL::gpu::CacheHint::None>
+SYCL_EXTERNAL void
+__esimd_surf_write(sycl::INTEL::gpu::vector_type_t<uint16_t, N> pred,
+                   int16_t scale, SurfIndAliasTy surf_ind,
+                   uint32_t global_offset,
+                   sycl::INTEL::gpu::vector_type_t<uint32_t, N> elem_offsets,
+                   sycl::INTEL::gpu::vector_type_t<Ty, N> vals)
+#ifdef __SYCL_DEVICE_ONLY__
+    ;
+#else
+{
+  static_assert(N == 1 || N == 8 || N == 16);
+  static_assert(TySizeLog2 <= 2);
+  static_assert(std::is_integral<Ty>::value || TySizeLog2 == 2);
+  throw cl::sycl::feature_not_supported();
+}
+#endif // __SYCL_DEVICE_ONLY__
+
+// TODO bring the parameter order of __esimd* intrinsics in accordance with the
+// correponsing BE intrinsicics parameter order.
+
 // flat_atomic: flat-address atomic
 template <sycl::INTEL::gpu::EsimdAtomicOpType Op, typename Ty, int N,
           sycl::INTEL::gpu::CacheHint L1H = sycl::INTEL::gpu::CacheHint::None,
diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_memory.hpp
diff --git a/sycl/test/esimd/intrins_trans.cpp b/sycl/test/esimd/intrins_trans.cpp

Original file line number	Diff line number	Diff line change
`@@ -693,6 +693,12 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,`
`693`	`693`	`if (A->getEnabled())`
`694`	`694`	`Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {}));`
`695`	`695`	`}`
	`696`	`+`
	`697`	`+ if (FD->hasAttr<SYCLIntelStallEnableAttr>()) {`
	`698`	`+ llvm::Metadata *AttrMDArgs[] = {`
	`699`	`+ llvm::ConstantAsMetadata::get(Builder.getInt32(1))};`
	`700`	`+ Fn->setMetadata("stall_enable", llvm::MDNode::get(Context, AttrMDArgs));`
	`701`	`+ }`
`696`	`702`	`}`
`697`	`703`
`698`	`704`	`/// Determine whether the function F ends with a return stmt.`