Conversion of 'sycl.constructor(%0, %1) {type = @Accessor}' (DefaultCtor) (#57)

Add conversion of `sycl::accessor<int, 1, read_write, global_buffer, (placeholder)0>:: accessor()`.

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

Author: whitneywhtsang

mlir-sycl/include/mlir/Conversion/SYCLToLLVM/SYCLFuncRegistry.h

@@ -41,8 +41,10 @@ class SYCLFuncDescriptor {
     Unknown, 
 
     // Member functions for the sycl:accessor class.
-    AccessorInt1ReadWriteGlobalBufferFalseInit,  // sycl::accessor<int, 1, read_write, global_buffer, (placeholder)0>::
-                                                 //   __init(int AS1*, sycl::range<1>, sycl::range<1>, sycl::id<1>)
+    AccessorInt1ReadWriteGlobalBufferFalseCtorDefault, // sycl::accessor<int, 1, read_write, global_buffer, (placeholder)0>::
+                                                       //   accessor()
+    AccessorInt1ReadWriteGlobalBufferFalseInit,        // sycl::accessor<int, 1, read_write, global_buffer, (placeholder)0>::
+                                                       //   __init(int AS1*, sycl::range<1>, sycl::range<1>, sycl::id<1>)
 
     // Member functions for the sycl:id<n> class.
     Id1CtorDefault, // sycl::id<1>::id()

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

@@ -80,6 +80,7 @@ void SYCLFuncDescriptor::declareFunction(ModuleOp &module, OpBuilder &b) {
 
 bool SYCLAccessorFuncDescriptor::isValid(FuncId funcId) const {
   switch (funcId) {
+  case FuncId::AccessorInt1ReadWriteGlobalBufferFalseCtorDefault:
   case FuncId::AccessorInt1ReadWriteGlobalBufferFalseInit:
     return true;
   default:
@@ -180,11 +181,12 @@ SYCLFuncRegistry::getFuncId(SYCLFuncDescriptor::FuncKind funcKind, Type retType,
       continue;
     }
     if (desc.argTys.size() != argTypes.size()) {
-      LLVM_DEBUG(llvm::dbgs() << "\tskip, number of arguments does not match\n");
+      LLVM_DEBUG(llvm::dbgs()
+                 << "\tskip, number of arguments does not match\n");
       continue;
     }
     if (!std::equal(argTypes.begin(), argTypes.end(), desc.argTys.begin())) {
-      LLVM_DEBUG(llvm::dbgs() << "\tskip, arguments types do not match\n");      
+      LLVM_DEBUG(llvm::dbgs() << "\tskip, arguments types do not match\n");
       continue;
     }
 
@@ -212,19 +214,35 @@ void SYCLFuncRegistry::declareAccessorFuncDescriptors(
   MLIRContext *context = module.getContext();
   auto voidTy = LLVM::LLVMVoidType::get(context);
   auto i32Ty = IntegerType::get(context, 32);
-  auto i32PtrTy = LLVM::LLVMPointerType::get(i32Ty);
-  Type accessorInt1ReadWriteGlobalBufferPtrTy =
-      converter.convertType(MemRefType::get(
-          -1, AccessorType::get(context, i32Ty, 1, MemoryAccessMode::ReadWrite,
-                                MemoryTargetMode::GlobalBuffer, {})));
-  Type id1Ty = converter.convertType(IDType::get(context, 1));
-  Type range1Ty = converter.convertType(RangeType::get(context, 1));
+  auto i32PtrTy = converter.convertType(MemRefType::get(-1, i32Ty));
+
+  constexpr unsigned int dim = 1;
+  Type id1SYCLTy = IDType::get(context, dim);
+  Type range1SYCLTy = RangeType::get(context, dim);
+  Type accessorImplDeviceSYCLTy = AccessorImplDeviceType::get(
+      context, dim, {id1SYCLTy, range1SYCLTy, range1SYCLTy});
+  Type accessorInt1ReadWriteGlobalBufferSYCLTy = AccessorType::get(
+      context, i32Ty, dim, MemoryAccessMode::ReadWrite,
+      MemoryTargetMode::GlobalBuffer, {accessorImplDeviceSYCLTy});
+  Type accessorInt1ReadWriteGlobalBufferPtrTy = converter.convertType(
+      MemRefType::get(-1, accessorInt1ReadWriteGlobalBufferSYCLTy));
+
+  Type id1Ty = converter.convertType(id1SYCLTy);
+  Type range1Ty = converter.convertType(range1SYCLTy);
 
   // Construct the SYCL functions descriptors for the sycl::accessor<n> type.
   // Descriptor format: (enum, function name, signature).
   // clang-format off
   std::vector<SYCLFuncDescriptor> descriptors = {
       // sycl::accessor<int, 1, read_write, global_buffer, (placeholder)0>::
+      //   accessor()
+      SYCLAccessorFuncDescriptor(
+          FuncId::AccessorInt1ReadWriteGlobalBufferFalseCtorDefault,
+          "_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_"
+          "6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_"
+          "property_listIJEEEEC2Ev",
+          voidTy, {accessorInt1ReadWriteGlobalBufferPtrTy}),
+      // sycl::accessor<int, 1, read_write, global_buffer, (placeholder)0>::
       //   __init(int AS1*, sycl::range<1>, sycl::range<1>, sycl::id<1>)
       SYCLAccessorFuncDescriptor(
           FuncId::AccessorInt1ReadWriteGlobalBufferFalseInit,

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

@@ -97,7 +97,7 @@ static Optional<Type> convertBodyType(StringRef name,
 static Optional<Type>
 convertAccessorImplDeviceType(sycl::AccessorImplDeviceType type,
                               LLVMTypeConverter &converter) {
-  return convertBodyType("class.cl::sycl::detail::AccessorImplDevice" +
+  return convertBodyType("class.cl::sycl::detail::AccessorImplDevice." +
                              std::to_string(type.getDimension()),
                          type.getBody(), converter);
 }
@@ -107,7 +107,7 @@ static Optional<Type> convertAccessorType(sycl::AccessorType type,
                                           LLVMTypeConverter &converter) {
   auto convertedTy = LLVM::LLVMStructType::getIdentified(
       &converter.getContext(),
-      "class.cl::sycl::accessor" + std::to_string(type.getDimension()));
+      "class.cl::sycl::accessor." + std::to_string(type.getDimension()));
   if (!convertedTy.isInitialized()) {
     SmallVector<Type> convertedElemTypes;
     convertedElemTypes.reserve(type.getBody().size());

mlir-sycl/test/Conversion/SYCLToLLVM/func-ops-to-llvm.mlir

@@ -1,5 +1,19 @@
 // RUN: sycl-mlir-opt -split-input-file -convert-sycl-to-llvm -verify-diagnostics %s | FileCheck %s
 
+//===-------------------------------------------------------------------------------------------------===//
+// Constructors for sycl::accessor<t, d, m, t, p>::accessor()
+//===-------------------------------------------------------------------------------------------------===//
+
+// CHECK: llvm.func @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev([[THIS_PTR_TYPE:!llvm.struct<\(ptr<struct<"class.cl::sycl::accessor.1",.*]])
+!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>)>)>
+func.func @accessorInt1ReadWriteGlobalBufferFalseCtor(%arg0: memref<?x!sycl_accessor_1_i32_read_write_global_buffer>) {
+  // CHECK: llvm.call @_ZN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev({{.*}}) : ([[THIS_PTR_TYPE]]) -> ()
+  sycl.constructor(%arg0) {Type = @accessor} : (memref<?x!sycl_accessor_1_i32_read_write_global_buffer>) -> () 
+  return
+}
+
+// -----
+
 //===-------------------------------------------------------------------------------------------------===//
 // Constructors for sycl::id<n>::id()
 //===-------------------------------------------------------------------------------------------------===//