Add MangledName attribute to sycl.constructor and sycl.call (#58)

We noticed that not all templated function can be register in the registry, e.g., function with template field `Type`, where `Type` can be user defined type.

Will create a separate PR for cleaning up `SYCLFuncRegistry`.

Signed-off-by: Tsang, Whitney <[email protected]>

Author: whitneywhtsang

mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td

@@ -74,6 +74,7 @@ def SYCLConstructorOp : SYCL_Op<"constructor", []> {
 
   let arguments = (ins
     FlatSymbolRefAttr:$Type,
+    FlatSymbolRefAttr:$MangledName,
     Variadic<ConstructorArgs>:$Args
   );
   let results = (outs);
@@ -122,6 +123,7 @@ def SYCLCallOp : SYCL_Op<"call", []> {
   let arguments = (ins
     OptionalAttr<FlatSymbolRefAttr>:$Type,
     FlatSymbolRefAttr:$Function,
+    FlatSymbolRefAttr:$MangledName,
     Variadic<AnyType>:$Args
   );
   let results = (outs Optional<AnyType>:$Result);
@@ -132,12 +134,14 @@ def SYCLCallOp : SYCL_Op<"call", []> {
     "::llvm::Optional<::mlir::Type>":$Result,
     "::llvm::Optional<::llvm::StringRef>":$Type,
     "::llvm::StringRef":$Function,
+    "::llvm::StringRef":$MangledName,
     "::mlir::ValueRange":$Args), [{
       odsState.addOperands(Args);
       if (Type.hasValue()) {
         odsState.addAttribute(TypeAttrName(odsState.name), ::mlir::SymbolRefAttr::get(odsBuilder.getContext(), Type.getValue()));
       }
       odsState.addAttribute(FunctionAttrName(odsState.name), ::mlir::SymbolRefAttr::get(odsBuilder.getContext(), Function));
+      odsState.addAttribute(MangledNameAttrName(odsState.name), ::mlir::SymbolRefAttr::get(odsBuilder.getContext(), MangledName));
       if (Result.hasValue()) {
         odsState.addTypes(Result.getValue());
       }

mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLToLLVM.cpp

@@ -237,17 +237,17 @@ class CallPattern final : public SYCLToLLVMConversion<sycl::SYCLCallOp> {
                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());
+    LLVMBuilder builder(rewriter, op.getLoc());
+    SmallVector<Type> operandTypes(opAdaptor.Args().getTypes());
+    FlatSymbolRefAttr funcRef = builder.getOrInsertFuncDecl(
+        opAdaptor.MangledName(), retType, operandTypes, module);
+    builder.genCall(funcRef, {}, opAdaptor.getOperands());
 
     LLVM_DEBUG({
       Operation *func = op->getParentOfType<LLVM::LLVMFuncOp>();
@@ -294,14 +294,13 @@ class ConstructorPattern final
                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.
-    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());
+
+    LLVMBuilder builder(rewriter, op.getLoc());
+    SmallVector<Type> operandTypes(opAdaptor.Args().getTypes());
+    FlatSymbolRefAttr funcRef = builder.getOrInsertFuncDecl(
+        opAdaptor.MangledName(), LLVM::LLVMVoidType::get(module.getContext()),
+        operandTypes, module);
+    builder.genCall(funcRef, {}, opAdaptor.getOperands());
 
     LLVM_DEBUG({
       Operation *func = op->getParentOfType<LLVM::LLVMFuncOp>();

mlir-sycl/test/Conversion/SYCLToLLVM/sycl-call-to-llvm.mlir

@@ -9,7 +9,7 @@
 // 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>) -> ()
+  sycl.call(%arg0, %arg1, %arg2, %arg3, %arg4) {Function = @__init, MangledName = @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEE6__initEPU3AS1iNS0_5rangeILi1EEESE_NS0_2idILi1EEE, Type = @accessor} : (memref<?x!sycl_accessor_1_i32_read_write_global_buffer>, memref<?xi32>, !sycl.range<1>, !sycl.range<1>, !sycl.id<1>) -> ()
   return
 }