Add support of memref of struct type (#44)

`SubIndexOpLowering` was unable to correctly handle memref of struct type.
memref of struct type is only generated for struct that has at least one entry of SYCL type, otherwise a llvm pointer type is generated instead of a memref.  (ref: https://github.com/InteonCo/Polygeist/blob/main/tools/cgeist/Lib/clang-mlir.cc#L5421)
When a memref element type is a struct type, the return type of a `polygeist.subindex` should be a memref of the element type of the struct. (ref: https://github.com/InteonCo/Polygeist/blob/main/tools/cgeist/Lib/clang-mlir.cc#L703)

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

Author: whitneywhtsang

polygeist/lib/polygeist/Passes/ConvertPolygeistToLLVM.cpp

@@ -29,6 +29,7 @@
 #include "mlir/Dialect/LLVMIR/LLVMTypes.h"
 #include "mlir/Dialect/OpenMP/OpenMPDialect.h"
 #include "mlir/Dialect/SCF/IR/SCF.h"
+#include "mlir/Dialect/SYCL/IR/SYCLOpsTypes.h"
 #include "mlir/IR/BlockAndValueMapping.h"
 #include "mlir/IR/ImplicitLocOpBuilder.h"
 #include "mlir/Transforms/RegionUtils.h"
@@ -111,6 +112,38 @@ struct SubIndexOpLowering : public ConvertOpToLLVMPattern<SubIndexOp> {
       targetMemRef.setOffset(rewriter, loc, baseOffset);
     }
 
+    if (auto ST = sourceMemRefType.getElementType()
+                      .dyn_cast<mlir::LLVM::LLVMStructType>()) {
+      // According to MLIRASTConsumer::getMLIRType() in clang-mlir.cc, memref of
+      // struct type is only generated for struct that has at least one entry of
+      // SYCL type, otherwise a llvm pointer type is generated instead of a
+      // memref.
+      assert(any_of(ST.getBody(),
+                    [](Type Element) {
+                      return Element
+                          .isa<mlir::sycl::IDType, mlir::sycl::AccessorType,
+                               mlir::sycl::RangeType,
+                               mlir::sycl::AccessorImplDeviceType,
+                               mlir::sycl::ArrayType, mlir::sycl::ItemType,
+                               mlir::sycl::ItemBaseType, mlir::sycl::NdItemType,
+                               mlir::sycl::GroupType>();
+                    }) &&
+             "Expecting at least one element type of the struct to be a SYCL "
+             "type");
+      // According to MLIRScanner::InitializeValueByInitListExpr() in
+      // clang-mlir.cc, when a memref element type is a struct type, the return
+      // type of a polygeist.subindex should be a memref of the element type of
+      // the struct.
+      auto elemPtrTy = LLVM::LLVMPointerType::get(
+          getTypeConverter()->convertType(viewMemRefType.getElementType()));
+      auto gep = rewriter.create<LLVM::GEPOp>(loc, elemPtrTy, prev, idxs);
+      MemRefDescriptor nexRef = createMemRefDescriptor(
+          loc, viewMemRefType, gep, gep, sizes, strides, rewriter);
+
+      rewriter.replaceOp(subViewOp, {nexRef});
+      return success();
+    }
+
     assert(getTypeConverter()->convertType(viewMemRefType.getElementType()) ==
                prev.getType().cast<LLVM::LLVMPointerType>().getElementType() &&
            "Expecting the element types to match");

polygeist/test/polygeist-opt/sycl/subindex.mlir

@@ -0,0 +1,13 @@
+// RUN: polygeist-opt --convert-polygeist-to-llvm %s | FileCheck %s
+
+// CHECK: [[GEP:%.*]] = llvm.getelementptr {{.*}} : (!llvm.ptr<struct<([[SYCLIDSTRUCT:struct<"class.cl::sycl::id.1"]], {{.*}} -> !llvm.ptr<[[SYCLIDSTRUCT]], {{.*}}
+// CHECK: [[MEMREF:%.*]] = llvm.mlir.undef : !llvm.struct<(ptr<[[SYCLIDSTRUCT]], {{.*}}
+// CHECK: {{.*}} = llvm.insertvalue [[GEP]], [[MEMREF]][0] : !llvm.struct<(ptr<[[SYCLIDSTRUCT]], {{.*}}
+
+module {
+  func.func @test(%arg0: memref<?x!llvm.struct<(!sycl.id<1>)>>) -> memref<?x!sycl.id<1>> {
+    %c0 = arith.constant 0 : index
+    %0 = "polygeist.subindex"(%arg0, %c0) : (memref<?x!llvm.struct<(!sycl.id<1>)>>, index) -> memref<?x!sycl.id<1>>
+    return %0 : memref<?x!sycl.id<1>>
+  }
+}