replica
/
hanchenye-Polygeist
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

Fix cast error

This commit is contained in:
William S. Moses 2021-12-29 18:54:19 -05:00 committed by William Moses
parent c689e8b21b
commit 3bfbd54ee5
12 changed files with 249 additions and 156 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
include/polygeist/PolygeistOps.td
View File

@ -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 = [{

98
lib/polygeist/Ops.cpp
View File

@ -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;
}

33
lib/polygeist/Passes/ParallelLower.cpp
View File

@ -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())

2
llvm-project

@ -1 +1 @@
Subproject commit 423ba12971bac8397c87fcf975ba6a4b7530ed28
Subproject commit 4ce3bb1143f2b6dd72bee09f045ca1bce509adb9

24
test/polygeist-opt/cudalower.mlir Normal file
View File

@ -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: }

129
tools/mlir-clang/Lib/clang-mlir.cc
View File

@ -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));
}
if (args[1].getType().isa<mlir::IntegerType>())
return ValueCategory(
builder.create<mlir::math::PowFOp>(loc, mlirType, args[0], args[1]),
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" ||
@ -2245,6 +2271,17 @@ ValueCategory MLIRScanner::VisitCallExpr(clang::CallExpr *expr) {
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,13 +4383,10 @@ 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>()) {
E->dump();
E->getType()->dump();
llvm::errs() << " scalar: " << scalar << " mlirty: " << mlirty << "\n";
}
auto mt = mlirty.cast<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());
@ -4389,6 +4397,13 @@ ValueCategory MLIRScanner::VisitCastExpr(CastExpr *E) {
}
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");
}
}
case clang::CastKind::CK_LValueToRValue: {
if (auto dr = dyn_cast<DeclRefExpr>(E->getSubExpr())) {

9
tools/mlir-clang/Test/Verification/dynalloc.c
View File

@ -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: }

23
tools/mlir-clang/Test/Verification/gettimeofday.c
View File

@ -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: }

19
tools/mlir-clang/Test/Verification/memcpystruct.c Normal file
View File

@ -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: }

21
tools/mlir-clang/Test/Verification/redstore.c
View File

@ -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: }

17
tools/mlir-clang/Test/Verification/redstore2.c
View File

@ -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: }

11
tools/mlir-clang/Test/Verification/virt.cpp
View File

@ -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: }