Conversion for sycl.call operator (#56)

etiotto · etiotto · commit 09df5981e365 · 2022-09-06T13:08:17.000-04:00
This PR adds a conversion pattern named CallPattern to the SYCLToLLVM conversion pass. This new conversion pattern lowers a sycl.call operator to an LLVM call to the SYCL member function named by the operation.

Signed-off-by: Tiotto, Ettore &lt;ettore.tiotto@intel.com&gt;
diff --git a/mlir-sycl/include/mlir/Conversion/SYCLToLLVM/SYCLFuncRegistry.h b/mlir-sycl/include/mlir/Conversion/SYCLToLLVM/SYCLFuncRegistry.h
@@ -21,6 +21,7 @@
 
 namespace mlir {
 class LLVMTypeConverter;
+
 namespace sycl {
 class SYCLFuncRegistry;
 
@@ -79,39 +80,37 @@ class SYCLFuncDescriptor {
     Range1CopyCtor,    // sycl::range<1>::range(sycl::range<1> const&)
     Range2CopyCtor,    // sycl::range<2>::range(sycl::range<2> const&)
     Range3CopyCtor,    // sycl::range<3>::range(sycl::range<3> const&)
-
-    Arr1CtorSizeT,  // sycl::detail::array<1>::array<1>(std::enable_if<(1)==(1), unsigned long>::type)
   };
   // clang-format on
 
-  /// Enumerates the kind of FuncId.
-  enum class FuncKind {
+  /// Enumerates the descriptor kind.
+  enum class Kind {
     Unknown,
     Accessor, // sycl::accessor class
     Id,       // sycl::id<n> class
     Range,    // sycl::range<n> class
   };
 
-  /// Each descriptor is uniquely identified by the pair {FuncId, FuncKind}.
+  /// Each descriptor is uniquely identified by the pair {FuncId, Kind}.
   class Id {
   public:
     friend class SYCLFuncRegistry;
     friend llvm::raw_ostream &operator<<(llvm::raw_ostream &, const Id &);
 
-    Id(FuncId id, FuncKind kind) : funcId(id), funcKind(kind) {
+    Id(FuncId id, Kind kind) : funcId(id), kind(kind) {
       assert(funcId != FuncId::Unknown && "Illegal function id");
-      assert(funcKind != FuncKind::Unknown && "Illegal function id kind");
+      assert(kind != Kind::Unknown && "Illegal descriptor kind");
     }
 
-    /// Maps a FuncKind to a descriptive name.
-    static std::map<SYCLFuncDescriptor::FuncKind, std::string> funcKindToName;
+    /// Maps a Kind to a descriptive name.
+    static std::map<SYCLFuncDescriptor::Kind, std::string> kindToName;
 
-    /// Maps a descriptive name to a FuncKind.
-    static std::map<std::string, SYCLFuncDescriptor::FuncKind> nameToFuncKind;
+    /// Maps a descriptive name to a Kind.
+    static std::map<std::string, SYCLFuncDescriptor::Kind> nameToKind;
 
   private:
     FuncId funcId = FuncId::Unknown;
-    FuncKind funcKind = FuncKind::Unknown;
+    Kind kind = Kind::Unknown;
   };
 
   /// Returns true if the given \p funcId is valid.
@@ -123,7 +122,7 @@ class SYCLFuncDescriptor {
                     Location loc);
 
 protected:
-  SYCLFuncDescriptor(FuncId funcId, FuncKind kind, StringRef name,
+  SYCLFuncDescriptor(FuncId funcId, Kind kind, StringRef name,
                      Type outputTy, ArrayRef<Type> argTys)
       : descId(funcId, kind), name(name), outputTy(outputTy),
         argTys(argTys.begin(), argTys.end()) {}
@@ -142,7 +141,7 @@ class SYCLFuncDescriptor {
 inline llvm::raw_ostream &operator<<(llvm::raw_ostream &os,
                                      const SYCLFuncDescriptor::Id &id) {
   os << "funcId=" << (int)id.funcId
-     << ", funcKind=" << SYCLFuncDescriptor::Id::funcKindToName[id.funcKind];
+     << ", kind=" << SYCLFuncDescriptor::Id::kindToName[id.kind];
   return os;
 }
 
@@ -161,7 +160,7 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &os,
   public:                                                                      \
     friend class SYCLFuncRegistry;                                             \
     using FuncId = SYCLFuncDescriptor::FuncId;                                 \
-    using FuncKind = SYCLFuncDescriptor::FuncKind;                             \
+    using Kind = SYCLFuncDescriptor::Kind;                                     \
                                                                                \
   private:                                                                     \
     ClassName(FuncId funcId, StringRef name, Type outputTy,                    \
@@ -171,17 +170,17 @@ inline llvm::raw_ostream &operator<<(llvm::raw_ostream &os,
     }                                                                          \
     bool isValid(FuncId) const override;                                       \
   };
-DEFINE_CLASS(SYCLAccessorFuncDescriptor, FuncKind::Accessor)
-DEFINE_CLASS(SYCLIdFuncDescriptor, FuncKind::Id)
-DEFINE_CLASS(SYCLRangeFuncDescriptor, FuncKind::Range)
+DEFINE_CLASS(SYCLAccessorFuncDescriptor, Kind::Accessor)
+DEFINE_CLASS(SYCLIdFuncDescriptor, Kind::Id)
+DEFINE_CLASS(SYCLRangeFuncDescriptor, Kind::Range)
 #undef DEFINE_CLASS
 
 /// \class SYCLFuncRegistry
 /// Singleton class representing the set of SYCL functions callable from the
 /// compiler.
 class SYCLFuncRegistry {
   using FuncId = SYCLFuncDescriptor::FuncId;
-  using FuncKind = SYCLFuncDescriptor::FuncKind;
+  using Kind = SYCLFuncDescriptor::Kind;
   using Registry = std::map<FuncId, SYCLFuncDescriptor>;
 
 public:
@@ -199,8 +198,8 @@ class SYCLFuncRegistry {
   }
 
   /// Returns the SYCLFuncDescriptor::Id::FuncId corresponding to the function
-  /// descriptor that matches the given \p funcKind and signature.
-  FuncId getFuncId(FuncKind funcKind, Type retType, TypeRange argTypes) const;
+  /// descriptor that matches the given \p kind and signature.
+  FuncId getFuncId(Kind kind, Type retType, TypeRange argTypes) const;
 
 private:
   SYCLFuncRegistry(ModuleOp &module, OpBuilder &builder);
diff --git a/mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLFuncRegistry.cpp b/mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLFuncRegistry.cpp
@@ -28,20 +28,20 @@ using namespace mlir::sycl;
 // SYCLFuncDescriptor::Id
 //===----------------------------------------------------------------------===//
 
-std::map<SYCLFuncDescriptor::FuncKind, std::string>
-    SYCLFuncDescriptor::Id::funcKindToName = {
-        {FuncKind::Accessor, "accessor"},
-        {FuncKind::Id, "id"},
-        {FuncKind::Range, "range"},
-        {FuncKind::Unknown, "unknown"},
+std::map<SYCLFuncDescriptor::Kind, std::string>
+    SYCLFuncDescriptor::Id::kindToName = {
+        {Kind::Accessor, "accessor"},
+        {Kind::Id, "id"},
+        {Kind::Range, "range"},
+        {Kind::Unknown, "unknown"},
 };
 
-std::map<std::string, SYCLFuncDescriptor::FuncKind>
-    SYCLFuncDescriptor::Id::nameToFuncKind = {
-        {"accessor", FuncKind::Accessor},
-        {"id", FuncKind::Id},
-        {"range", FuncKind::Range},
-        {"unknown", FuncKind::Unknown},
+std::map<std::string, SYCLFuncDescriptor::Kind>
+    SYCLFuncDescriptor::Id::nameToKind = {
+        {"accessor", Kind::Accessor},
+        {"id", Kind::Id},
+        {"range", Kind::Range},
+        {"unknown", Kind::Unknown},
 };
 
 //===----------------------------------------------------------------------===//
@@ -157,20 +157,19 @@ const SYCLFuncRegistry SYCLFuncRegistry::create(ModuleOp &module,
 }
 
 SYCLFuncDescriptor::FuncId
-SYCLFuncRegistry::getFuncId(SYCLFuncDescriptor::FuncKind funcKind, Type retType,
+SYCLFuncRegistry::getFuncId(SYCLFuncDescriptor::Kind kind, Type retType,
                             TypeRange argTypes) const {
-  assert(funcKind != SYCLFuncDescriptor::FuncKind::Unknown &&
-         "Invalid funcKind");
+  assert(kind != Kind::Unknown && "Invalid descriptor kind");
   LLVM_DEBUG(llvm::dbgs() << "Looking up function of kind: "
-                          << SYCLFuncDescriptor::Id::funcKindToName[funcKind]
+                          << SYCLFuncDescriptor::Id::kindToName.at(kind)
                           << "\n";);
 
   for (const auto &entry : registry) {
     const SYCLFuncDescriptor &desc = entry.second;
     LLVM_DEBUG(llvm::dbgs() << desc << "\n");
 
-    // Skip through entries that do not match the requested funcIdKind.
-    if (desc.descId.funcKind != funcKind) {
+    // Skip through entries that do not match the requested kind.
+    if (desc.descId.kind != kind) {
       LLVM_DEBUG(llvm::dbgs() << "\tskip, kind does not match\n");
       continue;
     }
@@ -186,7 +185,12 @@ SYCLFuncRegistry::getFuncId(SYCLFuncDescriptor::FuncKind funcKind, Type retType,
       continue;
     }
     if (!std::equal(argTypes.begin(), argTypes.end(), desc.argTys.begin())) {
-      LLVM_DEBUG(llvm::dbgs() << "\tskip, arguments types do not match\n");
+      LLVM_DEBUG({
+        auto pair = std::mismatch(argTypes.begin(), argTypes.end(),
+                                  desc.argTys.begin());
+        llvm::dbgs() << "\tskip, arguments " << *pair.first << " and "
+                     << *pair.second << " do not match\n";
+      });
       continue;
     }
 
@@ -268,14 +272,8 @@ void SYCLFuncRegistry::declareIdFuncDescriptors(LLVMTypeConverter &converter,
       converter.convertType(MemRefType::get(-1, IDType::get(context, 2)));
   Type id3PtrTy =
       converter.convertType(MemRefType::get(-1, IDType::get(context, 3)));
-
   auto voidTy = LLVM::LLVMVoidType::get(context);
   auto i64Ty = IntegerType::get(context, 64);
-  auto indexTy = IndexType::get(context);
-
-  auto arrayMemref = mlir::MemRefType::get(1, indexTy);
-  Type arr1PtrTy =
-      converter.convertType(mlir::MemRefType::get(-1, arrayMemref));
 
   // Construct the SYCL functions descriptors for the sycl::id<n> type.
   // Descriptor format: (enum, function name, signature).
@@ -329,6 +327,7 @@ void SYCLFuncRegistry::declareIdFuncDescriptors(LLVMTypeConverter &converter,
       SYCLIdFuncDescriptor(FuncId::Id3Ctor3SizeT,
           "_ZN2cl4sycl2idILi3EEC2ILi3EEENSt9enable_ifIXeqT_Li3EEmE4typeEmm",
           voidTy, {id3PtrTy, i64Ty, i64Ty, i64Ty}),
+      
       // sycl::id<1>::id(sycl::id<1> const&)
       SYCLIdFuncDescriptor(FuncId::Id1CopyCtor,
           "_ZN2cl4sycl2idILi1EEC1ERKS2_", voidTy, {id1PtrTy, id1PtrTy}),
diff --git a/mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLToLLVM.cpp b/mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLToLLVM.cpp
@@ -206,6 +206,64 @@ static Optional<Type> convertRangeType(sycl::RangeType type,
                             converter);
 }
 
+//===----------------------------------------------------------------------===//
+// CallPattern - Converts `sycl.call` to LLVM.
+//===----------------------------------------------------------------------===//
+
+class CallPattern final : public SYCLToLLVMConversion<sycl::SYCLCallOp> {
+public:
+  using SYCLToLLVMConversion<sycl::SYCLCallOp>::SYCLToLLVMConversion;
+  using FuncId = SYCLFuncDescriptor::FuncId;
+  using Kind = SYCLFuncDescriptor::Kind;
+
+  LogicalResult
+  matchAndRewrite(sycl::SYCLCallOp op, OpAdaptor opAdaptor,
+                  ConversionPatternRewriter &rewriter) const override {
+    assert(op.Type().has_value() &&
+           "Expecting op.Type() to have a valid value");
+    StringRef typeName = op.Type().value();
+    Kind kind = SYCLFuncDescriptor::Id::nameToKind.at(typeName.str());
+    assert((kind != Kind::Unknown) && "unknown descriptor kind");
+    return rewriteCall(kind, op, opAdaptor, rewriter);
+  }
+
+private:
+  /// Rewrite sycl.call() {Function = *, Type = *} to a LLVM call to the
+  /// appropriate member function.
+  LogicalResult rewriteCall(Kind kind, SYCLCallOp op, OpAdaptor opAdaptor,
+                            ConversionPatternRewriter &rewriter) const {
+    assert((kind != Kind::Unknown) && "Unexpected descriptor kind");
+    LLVM_DEBUG(llvm::dbgs() << "CallPattern: Rewriting op: "; op.dump();
+               llvm::dbgs() << "\n");
+
+    ModuleOp module = op.getOperation()->getParentOfType<ModuleOp>();
+    const auto &registry = SYCLFuncRegistry::create(module, rewriter);
+
+    /// Lookup the FuncId corresponding to the member function to use.
+    Type retType = op.getODSResults(0).empty()
+                       ? LLVM::LLVMVoidType::get(module.getContext())
+                       : op.Result().getType();
+
+    FuncId funcId =
+        registry.getFuncId(kind, retType, opAdaptor.Args().getTypes());
+    SYCLFuncDescriptor::call(funcId, opAdaptor.getOperands(), registry,
+                             rewriter, op.getLoc());
+
+    LLVM_DEBUG({
+      Operation *func = op->getParentOfType<LLVM::LLVMFuncOp>();
+      if (!func)
+        func = op->getParentOfType<func::FuncOp>();
+
+      assert(func && "Could not find parent function");
+      llvm::dbgs() << "ConstructorPattern: Function after rewrite:\n"
+                   << *func << "\n";
+    });
+
+    rewriter.eraseOp(op);
+    return success();
+  }
+};
+
 //===----------------------------------------------------------------------===//
 // ConstructorPattern - Converts `sycl.constructor` to LLVM.
 //===----------------------------------------------------------------------===//
@@ -214,36 +272,34 @@ class ConstructorPattern final
     : public SYCLToLLVMConversion<sycl::SYCLConstructorOp> {
 public:
   using SYCLToLLVMConversion<sycl::SYCLConstructorOp>::SYCLToLLVMConversion;
+  using FuncId = SYCLFuncDescriptor::FuncId;
+  using Kind = SYCLFuncDescriptor::Kind;
 
   LogicalResult
-  matchAndRewrite(sycl::SYCLConstructorOp op, OpAdaptor opAdaptor,
+  matchAndRewrite(SYCLConstructorOp op, OpAdaptor opAdaptor,
                   ConversionPatternRewriter &rewriter) const override {
     return rewriteConstructor(
-        SYCLFuncDescriptor::Id::nameToFuncKind.at(op.Type().str()), op,
-        opAdaptor, rewriter);
+        SYCLFuncDescriptor::Id::nameToKind.at(op.Type().str()), op, opAdaptor,
+        rewriter);
   }
 
 private:
   /// Rewrite sycl.constructor() { type = * } to a LLVM call to the appropriate
   /// constructor function.
-  LogicalResult rewriteConstructor(SYCLFuncDescriptor::FuncKind ctorKind,
-                                   SYCLConstructorOp op, OpAdaptor opAdaptor,
+  LogicalResult rewriteConstructor(Kind kind, SYCLConstructorOp op,
+                                   OpAdaptor opAdaptor,
                                    ConversionPatternRewriter &rewriter) const {
-    assert((ctorKind != SYCLFuncDescriptor::FuncKind::Unknown) &&
-           "Unexpected ctorKind");
+    assert((kind != Kind::Unknown) && "Unexpected descriptor kind");
     LLVM_DEBUG(llvm::dbgs() << "ConstructorPattern: Rewriting op: "; op.dump();
                llvm::dbgs() << "\n");
 
     ModuleOp module = op.getOperation()->getParentOfType<ModuleOp>();
     const auto &registry = SYCLFuncRegistry::create(module, rewriter);
 
-    /// Lookup the FuncId corresponding to the ctor function to use, which is
-    /// determined based on 'ctorKind) the kind of constructor to search for, and
-    /// the LLVM types of the sycl.constructor arguments.
-    SYCLFuncDescriptor::FuncId funcId = registry.getFuncId(
-        ctorKind, LLVM::LLVMVoidType::get(module.getContext()),
-        opAdaptor.Args().getTypes());
-
+    /// Lookup the FuncId corresponding to the ctor function to use.
+    auto retType = LLVM::LLVMVoidType::get(module.getContext());
+    FuncId funcId =
+        registry.getFuncId(kind, retType, opAdaptor.Args().getTypes());
     SYCLFuncDescriptor::call(funcId, opAdaptor.getOperands(), registry,
                              rewriter, op.getLoc());
 
@@ -300,5 +356,6 @@ void mlir::sycl::populateSYCLToLLVMConversionPatterns(
     LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
   populateSYCLToLLVMTypeConversion(typeConverter);
 
+  patterns.add<CallPattern>(patterns.getContext(), typeConverter);
   patterns.add<ConstructorPattern>(patterns.getContext(), typeConverter);
 }
diff --git a/mlir-sycl/test/Conversion/SYCLToLLVM/sycl-call-to-llvm.mlir b/mlir-sycl/test/Conversion/SYCLToLLVM/sycl-call-to-llvm.mlir
@@ -0,0 +1,16 @@
+// RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm -verify-diagnostics %s | FileCheck %s
+
+//===-------------------------------------------------------------------------------------------------===//
+// Member functions for sycl::accessor
+//===-------------------------------------------------------------------------------------------------===//
+
+!sycl_accessor_1_i32_read_write_global_buffer = !sycl.accessor<[1, i32, read_write, global_buffer], (!sycl.accessor_impl_device<[1], (!sycl.id<1>, !sycl.range<1>, !sycl.range<1>)>, !llvm.struct<(ptr<i32, 1>)>)>
+
+// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE([[ARG_TYPES:!llvm.struct<\(ptr<struct<"class.cl::sycl::accessor.1",.*]])
+func.func @accessorInit1(%arg0: memref<?x!sycl_accessor_1_i32_read_write_global_buffer>, %arg1: memref<?xi32>, %arg2: !sycl.range<1>, %arg3: !sycl.range<1>, %arg4: !sycl.id<1>) {
+  // CHECK: llvm.call @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE({{.*}}) : ([[ARG_TYPES]]) -> ()
+  sycl.call(%arg0, %arg1, %arg2, %arg3, %arg4) {Function = @__init, Type = @accessor} : (memref<?x!sycl_accessor_1_i32_read_write_global_buffer>, memref<?xi32>, !sycl.range<1>, !sycl.range<1>, !sycl.id<1>) -> ()
+  return
+}
+
+// -----
diff --git a/mlir-sycl/test/Conversion/SYCLToLLVM/sycl-constructor-to-llvm.mlir b/mlir-sycl/test/Conversion/SYCLToLLVM/sycl-constructor-to-llvm.mlir
