From 3bfbd54ee59234d67eff1bc85308f5722ead065b Mon Sep 17 00:00:00 2001
From: "William S. Moses" <gh@wsmoses.com>
Date: Wed, 29 Dec 2021 18:54:19 -0500
Subject: [PATCH] Fix cast error

---
 include/polygeist/PolygeistOps.td             |   3 +
 lib/polygeist/Ops.cpp                         |  98 +++++++-----
 lib/polygeist/Passes/ParallelLower.cpp        |  33 ++--
 llvm-project                                  |   2 +-
 test/polygeist-opt/cudalower.mlir             |  24 +++
 tools/mlir-clang/Lib/clang-mlir.cc            | 145 ++++++++++--------
 tools/mlir-clang/Test/Verification/dynalloc.c |   9 +-
 .../Test/Verification/gettimeofday.c          |  23 ++-
 .../Test/Verification/memcpystruct.c          |  19 +++
 tools/mlir-clang/Test/Verification/redstore.c |  21 ++-
 .../mlir-clang/Test/Verification/redstore2.c  |  17 +-
 tools/mlir-clang/Test/Verification/virt.cpp   |  11 +-
 12 files changed, 249 insertions(+), 156 deletions(-)
 create mode 100644 test/polygeist-opt/cudalower.mlir
 create mode 100644 tools/mlir-clang/Test/Verification/memcpystruct.c

diff --git a/include/polygeist/PolygeistOps.td b/include/polygeist/PolygeistOps.td
index 8bebc99..33c5716 100644
--- a/include/polygeist/PolygeistOps.td
+++ b/include/polygeist/PolygeistOps.td
@@ -36,6 +36,7 @@ def SubIndexOp : Polygeist_Op<"subindex", [
   let arguments = (ins AnyMemRef : $source, Index : $index);
   let results = (outs AnyMemRef : $result);
 
+  let hasFolder = 1;
   let hasCanonicalizer = 1;
 
   let extraClassDeclaration = [{
@@ -55,6 +56,7 @@ def Memref2PointerOp : Polygeist_Op<"memref2pointer", [
   let arguments = (ins AnyMemRef : $source);
   let results = (outs LLVM_AnyPointer : $result);
 
+  let hasFolder = 1;
   let hasCanonicalizer = 1;
   
   let extraClassDeclaration = [{
@@ -70,6 +72,7 @@ def Pointer2MemrefOp : Polygeist_Op<"pointer2memref", [
   let arguments = (ins LLVM_AnyPointer : $source);
   let results = (outs AnyMemRef : $result);
 
+  let hasFolder = 1;
   let hasCanonicalizer = 1;
   
   let extraClassDeclaration = [{
diff --git a/lib/polygeist/Ops.cpp b/lib/polygeist/Ops.cpp
index 5bbccdd..6eda43d 100644
--- a/lib/polygeist/Ops.cpp
+++ b/lib/polygeist/Ops.cpp
@@ -103,6 +103,8 @@ public:
 
     if (!memref::CastOp::canFoldIntoConsumerOp(castOp))
       return failure();
+    if (subViewOp.getType().getElementType() != subViewOp.result().getType().cast<MemRefType>().getElementType())
+      return failure();
 
     rewriter.replaceOpWithNewOp<SubIndexOp>(
         subViewOp, subViewOp.result().getType().cast<MemRefType>(),
@@ -126,6 +128,8 @@ public:
     if (castOp.getType().cast<MemRefType>().getShape().size() !=
         subindexOp.getType().cast<MemRefType>().getShape().size())
       return failure();
+    if (castOp.getType().cast<MemRefType>().getElementType() != subindexOp.result().getType().cast<MemRefType>().getElementType())
+      return failure();
 
     rewriter.replaceOpWithNewOp<SubIndexOp>(
         castOp, castOp.getType(), subindexOp.source(), subindexOp.index());
@@ -182,7 +186,7 @@ public:
       if (!cidx)
         return failure();
 
-      if (cidx.value() != 0 && cidx.value() != -1)
+      if (cidx.value() != 0)
         return failure();
 
       rewriter.replaceOpWithNewOp<memref::CastOp>(subViewOp, subViewOp.source(),
@@ -686,23 +690,6 @@ void SubIndexOp::getCanonicalizationPatterns(OwningRewritePatternList &results,
   // Disabled: SubToSubView
 }
 
-/// Simplify memref2pointer(cast(x)) to memref2pointer(x)
-class Memref2PointerCast final : public OpRewritePattern<Memref2PointerOp> {
-public:
-  using OpRewritePattern<Memref2PointerOp>::OpRewritePattern;
-
-  LogicalResult matchAndRewrite(Memref2PointerOp op,
-                                PatternRewriter &rewriter) const override {
-    auto src = op.source().getDefiningOp<memref::CastOp>();
-    if (!src)
-      return failure();
-
-    rewriter.replaceOpWithNewOp<polygeist::Memref2PointerOp>(op, op.getType(),
-                                                             src.source());
-    return success();
-  }
-};
-
 /// Simplify pointer2memref(memref2pointer(x)) to cast(x)
 class Memref2Pointer2MemrefCast final
     : public OpRewritePattern<Pointer2MemrefOp> {
@@ -719,9 +706,46 @@ public:
     return success();
   }
 };
+/// Simplify pointer2memref(memref2pointer(x)) to cast(x)
+class Memref2PointerIndex final
+    : public OpRewritePattern<Memref2PointerOp> {
+public:
+  using OpRewritePattern<Memref2PointerOp>::OpRewritePattern;
+
+  LogicalResult matchAndRewrite(Memref2PointerOp op,
+                                PatternRewriter &rewriter) const override {
+    auto src = op.source().getDefiningOp<SubIndexOp>();
+    if (!src)
+      return failure();
+
+    if (src.source().getType().cast<MemRefType>().getShape().size() != 1) return failure();
+
+    rewriter.replaceOpWithNewOp<LLVM::GEPOp>(op, op.getType(), rewriter.create<Memref2PointerOp>(op.getLoc(), op.getType(), src.source()),
+            std::vector<Value>({rewriter.create<arith::IndexCastOp>(op.getLoc(), rewriter.getI64Type(), src.index())}));
+    return success();
+  }
+};
+
+OpFoldResult Memref2PointerOp::fold(ArrayRef<Attribute> operands) {
+    if (auto subindex = source().getDefiningOp<SubIndexOp>()) {
+        if (auto cop = subindex.index().getDefiningOp<ConstantIntOp>()) {
+            if (cop.getValue() == 0) {
+                sourceMutable().assign(subindex.source());
+                return result();
+            }
+        }
+    }
+    /// Simplify memref2pointer(cast(x)) to memref2pointer(x)
+    if (auto mc = source().getDefiningOp<memref::CastOp>()) {
+        sourceMutable().assign(mc.source());
+        return result();
+    }
+    return nullptr;
+}
+
 void Memref2PointerOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
-  results.insert<Memref2PointerCast, Memref2Pointer2MemrefCast>(context);
+  results.insert<Memref2Pointer2MemrefCast, Memref2PointerIndex>(context);
 }
 
 /// Simplify cast(pointer2memref(x)) to pointer2memref(x)
@@ -835,26 +859,26 @@ public:
   }
 };
 
-/// Simplify pointer2memref(cast(x)) to pointer2memref(x)
-class BCPointer2Memref final : public OpRewritePattern<Pointer2MemrefOp> {
-public:
-  using OpRewritePattern<Pointer2MemrefOp>::OpRewritePattern;
-
-  LogicalResult matchAndRewrite(Pointer2MemrefOp op,
-                                PatternRewriter &rewriter) const override {
-    auto src = op.source().getDefiningOp<LLVM::BitcastOp>();
-    if (!src)
-      return failure();
-
-    rewriter.replaceOpWithNewOp<Pointer2MemrefOp>(op, op.getType(),
-                                                  src.getArg());
-    return success();
-  }
-};
-
 void Pointer2MemrefOp::getCanonicalizationPatterns(
     OwningRewritePatternList &results, MLIRContext *context) {
   results.insert<Pointer2MemrefCast, Pointer2Memref2PointerCast,
-                 Pointer2MemrefLoad, Pointer2MemrefStore, BCPointer2Memref>(
+                 Pointer2MemrefLoad, Pointer2MemrefStore>(
       context);
 }
+
+OpFoldResult Pointer2MemrefOp::fold(ArrayRef<Attribute> operands) {
+    /// Simplify pointer2memref(cast(x)) to pointer2memref(x)
+    if (auto mc = source().getDefiningOp<LLVM::BitcastOp>()) {
+        sourceMutable().assign(mc.getArg());
+        return result();
+    }
+    return nullptr;
+}
+
+OpFoldResult SubIndexOp::fold(ArrayRef<Attribute> operands) {
+    if (result().getType() == source().getType()) {
+        if (matchPattern(index(), m_Zero()))
+            return source();
+    }
+    return nullptr;
+}
diff --git a/lib/polygeist/Passes/ParallelLower.cpp b/lib/polygeist/Passes/ParallelLower.cpp
index 898991b..d2f3cee 100644
--- a/lib/polygeist/Passes/ParallelLower.cpp
+++ b/lib/polygeist/Passes/ParallelLower.cpp
@@ -392,9 +392,11 @@ void ParallelLower::runOnOperation() {
           storeOp.getLoc(), storeOp.memref(), indices));
       storeOp.erase();
     });
+    launchOp.erase();
+  });
 
-    container.walk([&](LLVM::CallOp call) {
-      if (call.getCallee().getValue() == "cudaMemcpy") {
+    getOperation().walk([&](LLVM::CallOp call) {
+      if (call.getCallee().getValue() == "cudaMemcpy" || call.getCallee().getValue() == "cudaMemcpyAsync") {
         OpBuilder bz(call);
         auto falsev = bz.create<ConstantIntOp>(call.getLoc(), false, 1);
         bz.create<LLVM::MemcpyOp>(call.getLoc(), call.getOperand(0),
@@ -403,8 +405,17 @@ void ParallelLower::runOnOperation() {
         call.replaceAllUsesWith(
             bz.create<ConstantIntOp>(call.getLoc(), 0, call.getType(0)));
         call.erase();
-      }
-      if (call.getCallee().getValue() == "cudaMemset") {
+      } else if (call.getCallee().getValue() == "cudaMemcpyToSymbol") {
+        OpBuilder bz(call);
+        auto falsev = bz.create<ConstantIntOp>(call.getLoc(), false, 1);
+        bz.create<LLVM::MemcpyOp>(call.getLoc(), 
+                                bz.create<LLVM::GEPOp>(call.getLoc(), call.getOperand(0).getType(), call.getOperand(0), std::vector<Value>({call.getOperand(3)})),
+                                  call.getOperand(1), call.getOperand(2),
+                                  /*isVolatile*/ falsev);
+        call.replaceAllUsesWith(
+            bz.create<ConstantIntOp>(call.getLoc(), 0, call.getType(0)));
+        call.erase();
+      } else if (call.getCallee().getValue() == "cudaMemset") {
         OpBuilder bz(call);
         auto falsev = bz.create<ConstantIntOp>(call.getLoc(), false, 1);
         bz.create<LLVM::MemsetOp>(call.getLoc(), call.getOperand(0),
@@ -413,17 +424,21 @@ void ParallelLower::runOnOperation() {
         Value vals[] = {call.getOperand(2)};
         call.replaceAllUsesWith(ArrayRef<Value>(vals));
         call.erase();
-      }
-      if (call.getCallee().getValue() == "cudaMalloc") {
-
+        /*
+      } else if (call.getCallee().getValue() == "cudaMalloc") {
         Value vals[] = {call.getOperand(0)};
         call.replaceAllUsesWith(ArrayRef<Value>(vals));
         call.erase();
+        */
+      } else if (call.getCallee().getValue() == "cudaDeviceSynchronize") {
+        OpBuilder bz(call);
+        auto retv = bz.create<ConstantIntOp>(call.getLoc(), 0, call.getResult(0).getType().cast<IntegerType>().getWidth());
+        Value vals[] = {retv};
+        call.replaceAllUsesWith(ArrayRef<Value>(vals));
+        call.erase();
       }
     });
 
-    launchOp.erase();
-  });
 
   for (auto f : toErase)
     if (f->use_empty())
diff --git a/llvm-project b/llvm-project
index 423ba12..4ce3bb1 160000
--- a/llvm-project
+++ b/llvm-project
@@ -1 +1 @@
-Subproject commit 423ba12971bac8397c87fcf975ba6a4b7530ed28
+Subproject commit 4ce3bb1143f2b6dd72bee09f045ca1bce509adb9
diff --git a/test/polygeist-opt/cudalower.mlir b/test/polygeist-opt/cudalower.mlir
new file mode 100644
index 0000000..5eb13d5
--- /dev/null
+++ b/test/polygeist-opt/cudalower.mlir
@@ -0,0 +1,24 @@
+// RUN: polygeist-opt --parallel-lower --split-input-file %s | FileCheck %s
+
+module attributes {llvm.data_layout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64", llvm.target_triple = "nvptx64-nvidia-cuda"}  {
+  llvm.func @cudaMemcpy(!llvm.ptr<i8>, !llvm.ptr<i8>, i64, i32) -> i32
+  func @_Z1aPiS_(%arg0: memref<?xi32>, %arg1: memref<?xi32>) -> i32 attributes {llvm.linkage = #llvm.linkage<external>} {
+    %c1_i32 = arith.constant 1 : i32
+    %c64_i64 = arith.constant 64 : i64
+    %0 = "polygeist.memref2pointer"(%arg0) : (memref<?xi32>) -> !llvm.ptr<i8>
+    %1 = "polygeist.memref2pointer"(%arg1) : (memref<?xi32>) -> !llvm.ptr<i8>
+    %2 = llvm.call @cudaMemcpy(%0, %1, %c64_i64, %c1_i32) : (!llvm.ptr<i8>, !llvm.ptr<i8>, i64, i32) -> i32
+    return %2 : i32
+  }
+}
+
+// CHECK:   func @_Z1aPiS_(%arg0: memref<?xi32>, %arg1: memref<?xi32>) -> i32 attributes {llvm.linkage = #llvm.linkage<external>} {
+// CHECK-NEXT:     %c64_i64 = arith.constant 64 : i64
+// CHECK-NEXT:     %false = arith.constant false
+// CHECK-NEXT:     %c0_i32 = arith.constant 0 : i32
+// CHECK-NEXT:     %0 = "polygeist.memref2pointer"(%arg0) : (memref<?xi32>) -> !llvm.ptr<i8>
+// CHECK-NEXT:     %1 = "polygeist.memref2pointer"(%arg1) : (memref<?xi32>) -> !llvm.ptr<i8>
+// CHECK-NEXT:     "llvm.intr.memcpy"(%0, %1, %c64_i64, %false) : (!llvm.ptr<i8>, !llvm.ptr<i8>, i64, i1) -> ()
+// CHECK-NEXT:     return %c0_i32 : i32
+// CHECK-NEXT:   }
+
diff --git a/tools/mlir-clang/Lib/clang-mlir.cc b/tools/mlir-clang/Lib/clang-mlir.cc
index 89da78c..c11563b 100644
--- a/tools/mlir-clang/Lib/clang-mlir.cc
+++ b/tools/mlir-clang/Lib/clang-mlir.cc
@@ -16,6 +16,7 @@
 #include <clang/Basic/LangStandard.h>
 #include <clang/Basic/TargetInfo.h>
 #include <clang/Basic/TargetOptions.h>
+#include <clang/Basic/OperatorKinds.h>
 #include <clang/Basic/Version.h>
 #include <clang/Driver/Compilation.h>
 #include <clang/Driver/Driver.h>
@@ -1765,6 +1766,21 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) {
   if (valEmitted.second)
     return valEmitted.first;
 
+  if (auto oc = dyn_cast<CXXOperatorCallExpr>(expr)) {
+    if (oc->getOperator() == clang::OO_EqualEqual) {
+	if (auto lhs = dyn_cast<CXXTypeidExpr>(expr->getArg(0))) {
+	  if (auto rhs = dyn_cast<CXXTypeidExpr>(expr->getArg(1))) {
+		QualType LT = lhs->isTypeOperand() ? lhs->getTypeOperand(Glob.CGM.getContext()) : lhs->getExprOperand()->getType();
+		QualType RT = rhs->isTypeOperand() ? rhs->getTypeOperand(Glob.CGM.getContext()) : rhs->getExprOperand()->getType();
+		llvm::Constant* LC = Glob.CGM.GetAddrOfRTTIDescriptor(LT);
+		llvm::Constant* RC = Glob.CGM.GetAddrOfRTTIDescriptor(RT);
+        auto postTy = getMLIRType(expr->getType()).cast<mlir::IntegerType>();
+		return ValueCategory(builder.create<arith::ConstantIntOp>(loc, LC == RC, postTy.getWidth()), false);
+	  }
+	}
+    }
+  }
+
   if (auto ic = dyn_cast<ImplicitCastExpr>(expr->getCallee()))
     if (auto sr = dyn_cast<DeclRefExpr>(ic->getSubExpr())) {
       if (sr->getDecl()->getIdentifier() &&
@@ -1877,30 +1893,40 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) {
     return val;
   };
 
-#if 0
   if (auto ic = dyn_cast<ImplicitCastExpr>(expr->getCallee()))
     if (auto sr = dyn_cast<DeclRefExpr>(ic->getSubExpr())) {
       // TODO add pow to standard dialect
       if (sr->getDecl()->getIdentifier() &&
           (sr->getDecl()->getName() == "__powf" ||
            sr->getDecl()->getName() == "pow" ||
+           sr->getDecl()->getName() == "__nv_pow" ||
+           sr->getDecl()->getName() == "__nv_powf" ||
+           sr->getDecl()->getName() == "__powi" ||
+           sr->getDecl()->getName() == "powi" ||
+           sr->getDecl()->getName() == "__nv_powi" ||
+           sr->getDecl()->getName() == "__nv_powi" ||
            sr->getDecl()->getName() == "powf")) {
         auto mlirType = getMLIRType(expr->getType());
         std::vector<mlir::Value> args;
         for (auto a : expr->arguments()) {
           args.push_back(Visit(a).getValue(builder));
         }
-        return ValueCategory(
-            builder.create<mlir::math::PowFOp>(loc, mlirType, args[0], args[1]),
+        if (args[1].getType().isa<mlir::IntegerType>())
+          return ValueCategory(
+            builder.create<LLVM::PowIOp>(loc, mlirType, args[0], args[1]),
+            /*isReference*/ false);
+        else
+          return ValueCategory(
+            builder.create<LLVM::PowOp>(loc, mlirType, args[0], args[1]),
             /*isReference*/ false);
       }
     }
-#endif
   if (auto ic = dyn_cast<ImplicitCastExpr>(expr->getCallee()))
     if (auto sr = dyn_cast<DeclRefExpr>(ic->getSubExpr())) {
       if (sr->getDecl()->getIdentifier() &&
           (sr->getDecl()->getName() == "__nv_fabsf" ||
            sr->getDecl()->getName() == "__nv_fabs" ||
+           sr->getDecl()->getName() == "__nv_abs" ||
            sr->getDecl()->getName() == "fabs" ||
            sr->getDecl()->getName() == "fabsf" ||
            sr->getDecl()->getName() == "__builtin_fabs" ||
@@ -2244,7 +2270,18 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) {
         auto srcSub = expr->getArg(1);
         while (auto BC = dyn_cast<clang::CastExpr>(srcSub))
           srcSub = BC->getSubExpr();
-
+        
+        if (sr->getDecl()->getName() == "memcpy" ||
+              sr->getDecl()->getName() == "__builtin_memcpy") {
+            std::vector<mlir::Value> args = {
+                getLLVM(expr->getArg(0)), getLLVM(expr->getArg(1)),
+                getLLVM(expr->getArg(2)), /*isVolatile*/
+                builder.create<ConstantIntOp>(loc, false, 1)};
+            builder.create<LLVM::MemcpyOp>(loc, args[0], args[1], args[2],
+                                           args[3]);
+            return ValueCategory(args[0], /*isReference*/ false);
+          }
+#if 0
         auto dstst = dstSub->getType()->getUnqualifiedDesugaredType();
         if (isa<clang::PointerType>(dstst) || isa<clang::ArrayType>(dstst)) {
 
@@ -2450,16 +2487,6 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) {
               }
             }
           }
-          if (sr->getDecl()->getName() == "memcpy" ||
-              sr->getDecl()->getName() == "__builtin_memcpy") {
-            std::vector<mlir::Value> args = {
-                getLLVM(expr->getArg(0)), getLLVM(expr->getArg(1)),
-                getLLVM(expr->getArg(2)), /*isVolatile*/
-                builder.create<ConstantIntOp>(loc, false, 1)};
-            builder.create<LLVM::MemcpyOp>(loc, args[0], args[1], args[2],
-                                           args[3]);
-            return ValueCategory(args[0], /*isReference*/ false);
-          }
           /*
           function.dump();
           expr->dump();
@@ -2470,6 +2497,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) {
           assert(0 && "unhandled cudaMemcpy");
           */
         }
+#endif
       }
     }
 
@@ -2572,6 +2600,7 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) {
       "fwrite",
       "__builtin_memcpy",
       "cudaMemcpy",
+      "cudaMemcpyAsync",
       "cudaMalloc",
       "open",
       "gettimeofday",
@@ -3196,8 +3225,6 @@ ValueCategory MLIRScanner::VisitUnaryOperator(clang::UnaryOperator *U) {
     } else if (auto mt = ty.dyn_cast<MemRefType>()) {
       auto shape = std::vector<int64_t>(mt.getShape());
       shape[0] = -1;
-      // Technically not legal per the -1
-      llvm::errs() << "postdec of memref subindex doing illegal op\n";
       auto mt0 = mlir::MemRefType::get(shape, mt.getElementType(),
                                        MemRefLayoutAttrInterface(),
                                        mt.getMemorySpace());
@@ -3367,12 +3394,7 @@ bool hasAffineArith(Operation *op, AffineExpr &expr,
 
 ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) {
   auto loc = getMLIRLocation(BO->getExprLoc());
-  auto lhs = Visit(BO->getLHS());
-  if (!lhs.val && BO->getOpcode() != clang::BinaryOperator::Opcode::BO_Comma) {
-    BO->dump();
-    BO->getLHS()->dump();
-    assert(lhs.val);
-  }
+  
   auto fixInteger = [&](mlir::Value res) {
     auto prevTy = res.getType().cast<mlir::IntegerType>();
     auto postTy = getMLIRType(BO->getType()).cast<mlir::IntegerType>();
@@ -3393,6 +3415,13 @@ ValueCategory MLIRScanner::VisitBinaryOperator(clang::BinaryOperator *BO) {
     return ValueCategory(res, /*isReference*/ false);
   };
 
+  auto lhs = Visit(BO->getLHS());
+  if (!lhs.val && BO->getOpcode() != clang::BinaryOperator::Opcode::BO_Comma) {
+    BO->dump();
+    BO->getLHS()->dump();
+    assert(lhs.val);
+  }
+
   switch (BO->getOpcode()) {
   case clang::BinaryOperator::Opcode::BO_LAnd: {
     mlir::Type types[] = {builder.getIntegerType(1)};
@@ -4115,35 +4144,17 @@ ValueCategory MLIRScanner::VisitOpaqueValueExpr(OpaqueValueExpr *E) {
   return res;
 }
 
-ValueCategory MLIRScanner::VisitCXXTypeidExpr(clang::CXXTypeidExpr *expr) {
-  assert(0 && "typeid expr unhandled");
-  llvm_unreachable("typeid expr unhandled");
-  /*
-  llvm::Type *StdTypeInfoPtrTy =
-    ConvertType(E->getType())->getPointerTo();
-
-  if (E->isTypeOperand()) {
-    llvm::Constant *TypeInfo =
-        CGM.GetAddrOfRTTIDescriptor(E->getTypeOperand(getContext()));
-    return Builder.CreateBitCast(TypeInfo, StdTypeInfoPtrTy);
-  }
-
-  // C++ [expr.typeid]p2:
-  //   When typeid is applied to a glvalue expression whose type is a
-  //   polymorphic class type, the result refers to a std::type_info object
-  //   representing the type of the most derived object (that is, the dynamic
-  //   type) to which the glvalue refers.
-  // If the operand is already most derived object, no need to look up vtable.
-  if (E->isPotentiallyEvaluated() && !E->isMostDerived(getContext())) {
-      assert(0 && "unhandled");
-    //return EmitTypeidFromVTable(*this, E->getExprOperand(),
-    //                            StdTypeInfoPtrTy);
-  }
-
-  QualType OperandTy = E->getExprOperand()->getType();
-  return Builder.CreateBitCast(CGM.GetAddrOfRTTIDescriptor(OperandTy),
-                               StdTypeInfoPtrTy);
-                               */
+ValueCategory MLIRScanner::VisitCXXTypeidExpr(clang::CXXTypeidExpr *E) {
+  QualType T;
+  if (E->isTypeOperand())
+    T = E->getTypeOperand(Glob.CGM.getContext());
+  else
+    T = E->getExprOperand()->getType();
+  llvm::Constant* C = Glob.CGM.GetAddrOfRTTIDescriptor(T);
+  C->dump();
+  auto ty = getMLIRType(E->getType());
+  llvm::errs() << ty << "\n";
+  assert(0 && "unhandled typeid");
 }
 
 ValueCategory
@@ -4372,23 +4383,27 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) {
     auto mlirty = getMLIRType(E->getType());
 
     if (auto PT = mlirty.dyn_cast<mlir::LLVM::LLVMPointerType>()) {
-      mlirty = mlir::MemRefType::get(-1, PT.getElementType());
-    } else if (!mlirty.isa<mlir::MemRefType>()) {
+      return ValueCategory(builder.create<mlir::polygeist::Memref2PointerOp>(
+                                   loc, PT, scalar),
+                               /*isReference*/ false);
+    } else if (auto mt = mlirty.dyn_cast<mlir::MemRefType>()) {
+        auto ty =
+            mlir::MemRefType::get(mt.getShape(), mt.getElementType(),
+                                  MemRefLayoutAttrInterface(), ut.getMemorySpace());
+        if (ut.getShape().size() == mt.getShape().size() + 1) {
+          return ValueCategory(builder.create<mlir::polygeist::SubIndexOp>(
+                                   loc, ty, scalar, getConstantIndex(0)),
+                               /*isReference*/ false);
+        }
+        return ValueCategory(builder.create<mlir::memref::CastOp>(loc, scalar, ty),
+                             /*isReference*/ false);
+    } else {
       E->dump();
       E->getType()->dump();
       llvm::errs() << " scalar: " << scalar << " mlirty: " << mlirty << "\n";
+      assert(0 && "illegal type for cast");
+      llvm_unreachable("illegal type for cast");
     }
-    auto mt = mlirty.cast<mlir::MemRefType>();
-    auto ty =
-        mlir::MemRefType::get(mt.getShape(), mt.getElementType(),
-                              MemRefLayoutAttrInterface(), ut.getMemorySpace());
-    if (ut.getShape().size() == mt.getShape().size() + 1) {
-      return ValueCategory(builder.create<mlir::polygeist::SubIndexOp>(
-                               loc, ty, scalar, getConstantIndex(0)),
-                           /*isReference*/ false);
-    }
-    return ValueCategory(builder.create<mlir::memref::CastOp>(loc, scalar, ty),
-                         /*isReference*/ false);
   }
   case clang::CastKind::CK_LValueToRValue: {
     if (auto dr = dyn_cast<DeclRefExpr>(E->getSubExpr())) {
diff --git a/tools/mlir-clang/Test/Verification/dynalloc.c b/tools/mlir-clang/Test/Verification/dynalloc.c
index 2db12e6..e1e82cd 100644
--- a/tools/mlir-clang/Test/Verification/dynalloc.c
+++ b/tools/mlir-clang/Test/Verification/dynalloc.c
@@ -9,14 +9,13 @@ void create_matrix(float *m, int size) {
 // CHECK:   func @create_matrix(%arg0: memref<?xf32>, %arg1: i32)
 // CHECK-DAG:     %c2 = arith.constant 2 : index
 // CHECK-DAG:     %c1 = arith.constant 1 : index
-// CHECK-NEXT:     %c0_i32 = arith.constant 0 : i32
+// CHECK-DAG:     %cst = arith.constant 0.000000e+00 : f32
 // CHECK-NEXT:     %0 = arith.index_cast %arg1 : i32 to index
 // CHECK-NEXT:     %1 = arith.muli %0, %c2 : index
 // CHECK-NEXT:     %2 = arith.subi %1, %c1 : index
 // CHECK-NEXT:     %3 = memref.alloca(%2) : memref<?xf32>
-// CHECK-NEXT:     %4 = arith.sitofp %c0_i32 : i32 to f32
-// CHECK-NEXT:     affine.store %4, %3[symbol(%0)] : memref<?xf32>
-// CHECK-NEXT:     %5 = affine.load %3[symbol(%0)] : memref<?xf32>
-// CHECK-NEXT:     affine.store %5, %arg0[symbol(%0)] : memref<?xf32>
+// CHECK-NEXT:     affine.store %cst, %3[symbol(%0)] : memref<?xf32>
+// CHECK-NEXT:     %[[i5:.+]] = affine.load %3[symbol(%0)] : memref<?xf32>
+// CHECK-NEXT:     affine.store %[[i5]], %arg0[symbol(%0)] : memref<?xf32>
 // CHECK-NEXT:     return
 // CHECK-NEXT:   }
diff --git a/tools/mlir-clang/Test/Verification/gettimeofday.c b/tools/mlir-clang/Test/Verification/gettimeofday.c
index 4e8ce6e..f1762c7 100644
--- a/tools/mlir-clang/Test/Verification/gettimeofday.c
+++ b/tools/mlir-clang/Test/Verification/gettimeofday.c
@@ -10,19 +10,16 @@ double alloc() {
 
 // CHECK:   func @alloc() -> f64
 // CHECK-NEXT:     %cst = arith.constant 9.9999999999999995E-7 : f64
-// CHECK-NEXT:     %c1_i32 = arith.constant 1 : i32
-// CHECK-NEXT:     %c0_i32 = arith.constant 0 : i32
-// CHECK-NEXT:     %c1_i64 = arith.constant 1 : i64
-// CHECK-NEXT:     %0 = llvm.alloca %c1_i64 x !llvm.struct<(i64, i64)> : (i64) -> !llvm.ptr<struct<(i64, i64)>>
-// CHECK-NEXT:     %1 = llvm.mlir.null : !llvm.ptr<struct<(i32, i32)>>
-// CHECK-NEXT:     %2 = llvm.call @gettimeofday(%0, %1) : (!llvm.ptr<struct<(i64, i64)>>, !llvm.ptr<struct<(i32, i32)>>) -> i32
-// CHECK-NEXT:     %3 = llvm.getelementptr %0[%c0_i32, %c0_i32] : (!llvm.ptr<struct<(i64, i64)>>, i32, i32) -> !llvm.ptr<i64>
-// CHECK-NEXT:     %4 = llvm.load %3 : !llvm.ptr<i64>
-// CHECK-NEXT:     %5 = llvm.getelementptr %0[%c0_i32, %c1_i32] : (!llvm.ptr<struct<(i64, i64)>>, i32, i32) -> !llvm.ptr<i64>
-// CHECK-NEXT:     %6 = llvm.load %5 : !llvm.ptr<i64>
-// CHECK-NEXT:     %7 = arith.sitofp %4 : i64 to f64
-// CHECK-NEXT:     %8 = arith.sitofp %6 : i64 to f64
+// CHECK-NEXT:     %0 = memref.alloca() : memref<1x2xi64>
+// CHECK-NEXT:     %1 = "polygeist.memref2pointer"(%0) : (memref<1x2xi64>) -> !llvm.ptr<struct<(i64, i64)>>
+// CHECK-NEXT:     %2 = llvm.mlir.null : !llvm.ptr<i8>
+// CHECK-NEXT:     %3 = llvm.bitcast %2 : !llvm.ptr<i8> to !llvm.ptr<struct<(i32, i32)>>
+// CHECK-NEXT:     %4 = llvm.call @gettimeofday(%1, %3) : (!llvm.ptr<struct<(i64, i64)>>, !llvm.ptr<struct<(i32, i32)>>) -> i32
+// CHECK-NEXT:     %5 = affine.load %0[0, 0] : memref<1x2xi64>
+// CHECK-NEXT:     %6 = arith.sitofp %5 : i64 to f64
+// CHECK-NEXT:     %7 = affine.load %0[0, 1] : memref<1x2xi64>
+// CHECK-NEXT:     %8 = arith.sitofp %7 : i64 to f64
 // CHECK-NEXT:     %9 = arith.mulf %8, %cst : f64
-// CHECK-NEXT:     %10 = arith.addf %7, %9 : f64
+// CHECK-NEXT:     %10 = arith.addf %6, %9 : f64
 // CHECK-NEXT:     return %10 : f64
 // CHECK-NEXT:   }
diff --git a/tools/mlir-clang/Test/Verification/memcpystruct.c b/tools/mlir-clang/Test/Verification/memcpystruct.c
new file mode 100644
index 0000000..8f18da9
--- /dev/null
+++ b/tools/mlir-clang/Test/Verification/memcpystruct.c
@@ -0,0 +1,19 @@
+// RUN: mlir-clang %s --function=* -S | FileCheck %s
+
+struct N {
+    int a;
+    int b;
+};
+
+void copy(struct N* dst, void* src) {
+    __builtin_memcpy(dst, src, sizeof(struct N));
+}
+
+// CHECK:   func @copy(%arg0: memref<?x2xi32>, %arg1: !llvm.ptr<i8>)
+// CHECK-NEXT:     %false = arith.constant false
+// CHECK-NEXT:     %c8_i64 = arith.constant 8 : i64
+// CHECK-NEXT:     %0 = "polygeist.memref2pointer"(%arg0) : (memref<?x2xi32>) -> !llvm.ptr<i8>
+// CHECK-NEXT:     "llvm.intr.memcpy"(%0, %arg1, %c8_i64, %false) : (!llvm.ptr<i8>, !llvm.ptr<i8>, i64, i1) -> ()
+// CHECK-NEXT:     return
+// CHECK-NEXT:   }
+
diff --git a/tools/mlir-clang/Test/Verification/redstore.c b/tools/mlir-clang/Test/Verification/redstore.c
index fe0c4d9..f1ff783 100644
--- a/tools/mlir-clang/Test/Verification/redstore.c
+++ b/tools/mlir-clang/Test/Verification/redstore.c
@@ -15,20 +15,19 @@ void sum(double *result, double* array, int N) {
 }
 
 // CHECK:  func @sum(%arg0: memref<?xf64>, %arg1: memref<?xf64>, %arg2: i32)
-// CHECK-NEXT:     %c0_i32 = arith.constant 0 : i32
+// CHECK-NEXT:     %cst = arith.constant 0.000000e+00 : f64
 // CHECK-NEXT:     %0 = arith.index_cast %arg2 : i32 to index
-// CHECK-NEXT:     %1 = arith.sitofp %c0_i32 : i32 to f64
 // CHECK-NEXT:     affine.for %arg3 = 0 to %0 {
-// CHECK-NEXT:       affine.store %1, %arg0[0] : memref<?xf64>
-// CHECK-NEXT:       %2 = affine.load %arg0[0] : memref<?xf64>
-// CHECK-NEXT:       %3 = affine.for %arg4 = 0 to 10 iter_args(%arg5 = %2) -> (f64) {
-// CHECK-NEXT:         %6 = affine.load %arg1[%arg4] : memref<?xf64>
-// CHECK-NEXT:         %7 = arith.addf %arg5, %6 : f64
-// CHECK-NEXT:         affine.yield %7 : f64
+// CHECK-NEXT:       affine.store %cst, %arg0[0] : memref<?xf64>
+// CHECK-NEXT:       %[[i2:.+]] = affine.load %arg0[0] : memref<?xf64>
+// CHECK-NEXT:       %[[i3:.+]] = affine.for %arg4 = 0 to 10 iter_args(%arg5 = %[[i2]]) -> (f64) {
+// CHECK-NEXT:         %[[i6:.+]] = affine.load %arg1[%arg4] : memref<?xf64>
+// CHECK-NEXT:         %[[i7:.+]] = arith.addf %arg5, %[[i6]] : f64
+// CHECK-NEXT:         affine.yield %[[i7]] : f64
 // CHECK-NEXT:       }
-// CHECK-NEXT:       affine.store %3, %arg0[0] : memref<?xf64>
-// CHECK-NEXT:       %4 = affine.load %arg0[0] : memref<?xf64>
-// CHECK-NEXT:       %5 = call @print(%4) : (f64) -> i32
+// CHECK-NEXT:       affine.store %[[i3]], %arg0[0] : memref<?xf64>
+// CHECK-NEXT:       %[[i4:.+]] = affine.load %arg0[0] : memref<?xf64>
+// CHECK-NEXT:       %{{.*}} = call @print(%[[i4]]) : (f64) -> i32
 // CHECK-NEXT:     }
 // CHECK-NEXT:     return
 // CHECK-NEXT:   }
diff --git a/tools/mlir-clang/Test/Verification/redstore2.c b/tools/mlir-clang/Test/Verification/redstore2.c
index 14b3b88..8b73ab0 100644
--- a/tools/mlir-clang/Test/Verification/redstore2.c
+++ b/tools/mlir-clang/Test/Verification/redstore2.c
@@ -10,15 +10,14 @@ void sum(double *result, double* array) {
 }
 
 // CHECK:  func @sum(%arg0: memref<?xf64>, %arg1: memref<?xf64>)
-// CHECK-NEXT:    %c0_i32 = arith.constant 0 : i32
-// CHECK-NEXT:    %0 = arith.sitofp %c0_i32 : i32 to f64
-// CHECK-NEXT:    affine.store %0, %arg0[0] : memref<?xf64>
-// CHECK-NEXT:    %1 = affine.load %arg0[0] : memref<?xf64>
-// CHECK-NEXT:    %2 = affine.for %arg2 = 0 to 10 iter_args(%arg3 = %1) -> (f64) {
-// CHECK-NEXT:      %3 = affine.load %arg1[%arg2] : memref<?xf64>
-// CHECK-NEXT:      %4 = arith.addf %arg3, %3 : f64
-// CHECK-NEXT:      affine.yield %4 : f64
+// CHECK-NEXT:    %cst = arith.constant 0.000000e+00 : f64
+// CHECK-NEXT:    affine.store %cst, %arg0[0] : memref<?xf64>
+// CHECK-NEXT:    %[[i1:.+]] = affine.load %arg0[0] : memref<?xf64>
+// CHECK-NEXT:    %[[i2:.+]] = affine.for %arg2 = 0 to 10 iter_args(%arg3 = %[[i1]]) -> (f64) {
+// CHECK-NEXT:      %[[i3:.+]] = affine.load %arg1[%arg2] : memref<?xf64>
+// CHECK-NEXT:      %[[i4:.+]] = arith.addf %arg3, %[[i3]] : f64
+// CHECK-NEXT:      affine.yield %[[i4]] : f64
 // CHECK-NEXT:    }
-// CHECK-NEXT:    affine.store %2, %arg0[0] : memref<?xf64>
+// CHECK-NEXT:    affine.store %[[i2]], %arg0[0] : memref<?xf64>
 // CHECK-NEXT:    return
 // CHECK-NEXT:  }
diff --git a/tools/mlir-clang/Test/Verification/virt.cpp b/tools/mlir-clang/Test/Verification/virt.cpp
index 33f371a..dda32b1 100644
--- a/tools/mlir-clang/Test/Verification/virt.cpp
+++ b/tools/mlir-clang/Test/Verification/virt.cpp
@@ -65,11 +65,10 @@ void make() {
 // CHECK-NEXT:   }
 // CHECK:   func @_ZN5FRootC1Ev(%arg0: memref<?x1xf32>) attributes {llvm.linkage = #llvm.linkage<linkonce_odr>} {
 // CHECK-NEXT:     %c0_i32 = arith.constant 0 : i32
-// CHECK-NEXT:     %cst = arith.constant 2.180000e+00 : f64
-// CHECK-NEXT:     %0 = arith.truncf %cst : f64 to f32
-// CHECK-NEXT:     affine.store %0, %arg0[0, 0] : memref<?x1xf32>
-// CHECK-NEXT:     %1 = llvm.mlir.addressof @str2 : !llvm.ptr<array<14 x i8>>
-// CHECK-NEXT:     %2 = llvm.getelementptr %1[%c0_i32, %c0_i32] : (!llvm.ptr<array<14 x i8>>, i32, i32) -> !llvm.ptr<i8>
-// CHECK-NEXT:     call @_Z5printPc(%2) : (!llvm.ptr<i8>) -> ()
+// CHECK-NEXT:     %cst = arith.constant 2.180000e+00 : f32
+// CHECK-NEXT:     affine.store %cst, %arg0[0, 0] : memref<?x1xf32>
+// CHECK-NEXT:     %[[i1:.+]] = llvm.mlir.addressof @str2 : !llvm.ptr<array<14 x i8>>
+// CHECK-NEXT:     %[[i2:.+]] = llvm.getelementptr %[[i1]][%c0_i32, %c0_i32] : (!llvm.ptr<array<14 x i8>>, i32, i32) -> !llvm.ptr<i8>
+// CHECK-NEXT:     call @_Z5printPc(%[[i2]]) : (!llvm.ptr<i8>) -> ()
 // CHECK-NEXT:     return
 // CHECK-NEXT:   }