Fix GEP for pointer of struct type (#48)

When getelementptr of a pointer of struct type, the first index 0 is needed to first get a pointer to the struct, then the second index is the index within the struct.

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

Author: whitneywhtsang

polygeist/lib/polygeist/Passes/ConvertPolygeistToLLVM.cpp

@@ -114,6 +114,15 @@ struct SubIndexOpLowering : public ConvertOpToLLVMPattern<SubIndexOp> {
 
     if (auto ST = sourceMemRefType.getElementType()
                       .dyn_cast<mlir::LLVM::LLVMStructType>()) {
+      assert(sourceMemRefType.getShape().size() ==
+                 viewMemRefType.getShape().size() &&
+             "Expecting the input and output MemRef size to be the same");
+      // The first index (zero) takes a pointer to the structure.
+      Value zero = rewriter.create<LLVM::ConstantOp>(
+          loc, idxs[0].getType(),
+          rewriter.getIntegerAttr(idxs[0].getType(), 0));
+      Value idxs[] = {zero, transformed.index()};
+
       // 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

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

@@ -1,6 +1,7 @@
 // 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: [[ZERO:%.*]] = llvm.mlir.constant(0 : i64) : i64
+// CHECK: [[GEP:%.*]] = llvm.getelementptr {{.*}}[[[ZERO]], 0] : (!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]], {{.*}}