diff --git a/polygeist/tools/cgeist/Lib/CGExpr.cc b/polygeist/tools/cgeist/Lib/CGExpr.cc index be7d1ec8ca42a..1d0f04e4bcb42 100644 --- a/polygeist/tools/cgeist/Lib/CGExpr.cc +++ b/polygeist/tools/cgeist/Lib/CGExpr.cc @@ -6,6 +6,7 @@ // //===----------------------------------------------------------------------===// +#include "Lib/TypeUtils.h" #include "clang-mlir.h" #include "utils.h" #include "llvm/ADT/TypeSwitch.h" @@ -2185,6 +2186,16 @@ ValueCategory MLIRScanner::EmitPromoted(Expr *E, QualType PromotionType) { return Res; } +ValueCategory MLIRScanner::CastToVoidPtr(ValueCategory Ptr) { + assert(mlirclang::isPointerOrMemRefTy(Ptr.val.getType()) && + "Expecting pointer or memref"); + + const auto DestType = + mlirclang::getPtrTyWithNewType(Ptr.val.getType(), builder.getI8Type()); + + return Ptr.BitCast(builder, loc, DestType); +} + ValueCategory MLIRScanner::EmitPromotedValue(Location Loc, ValueCategory Result, QualType PromotionType) { return Result.FPExt(builder, Loc, Glob.getTypes().getMLIRType(PromotionType)); @@ -2576,6 +2587,12 @@ BinOpInfo MLIRScanner::EmitBinOps(BinaryOperator *E, QualType PromotionType) { return {LHS, RHS, Ty, Opcode, E}; } +static void informNoOverflowCheck(LangOptions::SignedOverflowBehaviorTy SOB, + llvm::StringRef OpName) { + if (SOB != clang::LangOptions::SOB_Defined) + llvm::errs() << "Not emitting overflow-checked " << OpName << "\n"; +} + ValueCategory MLIRScanner::EmitBinMul(const BinOpInfo &Info) { auto lhs_v = Info.getLHS().getValue(builder); auto rhs_v = Info.getRHS().getValue(builder); @@ -2620,91 +2637,215 @@ ValueCategory MLIRScanner::EmitBinRem(const BinOpInfo &Info) { } } +/// Casts index of subindex operation conditionally. +static Optional castSubIndexOpIndex(OpBuilder &Builder, Location Loc, + ValueCategory Pointer, + ValueRange IdxList, bool IsSigned) { + if (Pointer.val.getType().isa()) { + assert(IdxList.size() == 1 && "SubIndexOp accepts just an index"); + return ValueCategory(IdxList.front(), false) + .IntCast(Builder, Loc, Builder.getIndexType(), IsSigned) + .val; + } + return llvm::None; +} + +ValueCategory MLIRScanner::EmitCheckedInBoundsPtrOffsetOp(mlir::Type ElemTy, + ValueCategory Pointer, + ValueRange IdxList, + bool IsSigned, bool) { + assert(mlirclang::isPointerOrMemRefTy(Pointer.val.getType()) && + "Expecting pointer or MemRef"); + assert(std::all_of(IdxList.begin(), IdxList.end(), + [](mlir::Value Val) { + return Val.getType().isa(); + }) && + "Expecting indices list"); + + if (Optional NewValue = + castSubIndexOpIndex(builder, loc, Pointer, IdxList, IsSigned)) + IdxList = *NewValue; + + return Pointer.InBoundsGEPOrSubIndex(builder, loc, ElemTy, IdxList); +} + +ValueCategory MLIRScanner::EmitPointerArithmetic(const BinOpInfo &Info) { + const auto *Expr = cast(Info.getExpr()); + + ValueCategory Pointer = Info.getLHS(); + auto *PointerOperand = Expr->getLHS(); + ValueCategory Index = Info.getRHS(); + auto *IndexOperand = Expr->getRHS(); + + const auto Opcode = Info.getOpcode(); + const auto IsSubtraction = + Opcode == clang::BO_Sub || Opcode == clang::BO_SubAssign; + + assert((!IsSubtraction || + mlirclang::isPointerOrMemRefTy(Pointer.val.getType())) && + "The LHS is always a pointer in a subtraction"); + + if (!mlirclang::isPointerOrMemRefTy(Pointer.val.getType())) { + std::swap(Pointer, Index); + std::swap(PointerOperand, IndexOperand); + } + + assert(Index.val.getType().isa() && "Expecting integer type"); + + auto PtrTy = Pointer.val.getType(); + + assert(mlirclang::isPointerOrMemRefTy(PtrTy) && "Expecting pointer type"); + + auto &CGM = Glob.getCGM(); + + // Some versions of glibc and gcc use idioms (particularly in their malloc + // routines) that add a pointer-sized integer (known to be a pointer + // value) to a null pointer in order to cast the value back to an integer + // or as part of a pointer alignment algorithm. This is undefined + // behavior, but we'd like to be able to compile programs that use it. + // + // Normally, we'd generate a GEP with a null-pointer base here in response + // to that code, but it's also UB to dereference a pointer created that + // way. Instead (as an acknowledged hack to tolerate the idiom) we will + // generate a direct cast of the integer value to a pointer. + // + // The idiom (p = nullptr + N) is not met if any of the following are + // true: + // + // The operation is subtraction. + // The index is not pointer-sized. + // The pointer type is not byte-sized. + // + if (BinaryOperator::isNullPointerArithmeticExtension( + CGM.getContext(), Opcode, PointerOperand, IndexOperand)) { + return Index.IntToPtr(builder, loc, PtrTy); + } + + auto &DL = CGM.getDataLayout(); + const unsigned IndexTypeSize = DL.getIndexTypeSizeInBits( + CGM.getTypes().ConvertType(PointerOperand->getType())); + const auto IsSigned = + IndexOperand->getType()->isSignedIntegerOrEnumerationType(); + const unsigned Width = Index.val.getType().getIntOrFloatBitWidth(); + if (Width != IndexTypeSize) { + // Zero-extend or sign-extend the pointer value according to + // whether the index is signed or not. + Index = Index.IntCast(builder, loc, builder.getIntegerType(IndexTypeSize), + IsSigned); + } + + // If this is subtraction, negate the index. + if (IsSubtraction) + Index = Index.Neg(builder, loc); + + const auto *PointerType = + PointerOperand->getType()->getAs(); + + assert(PointerType && "Not pointer type"); + + QualType ElementType = PointerType->getPointeeType(); + assert(!CGM.getContext().getAsVariableArrayType(ElementType) && + "Not implemented yet"); + + // Explicitly handle GNU void* and function pointer arithmetic extensions. + // The GNU void* casts amount to no-ops since our void* type is i8*, but + // this is future proof. + if (ElementType->isVoidType() || ElementType->isFunctionType()) { + assert(PtrTy.isa() && "Expecting pointer type"); + auto Result = CastToVoidPtr(Pointer); + Result = Result.GEP(builder, loc, builder.getI8Type(), Index.val); + return Result.BitCast(builder, loc, Pointer.val.getType()); + } + + auto ElemTy = Glob.getTypes().getMLIRType(ElementType); + if (CGM.getLangOpts().isSignedOverflowDefined()) { + if (Optional NewIndex = + castSubIndexOpIndex(builder, loc, Pointer, Index.val, IsSigned)) + Index.val = *NewIndex; + return Pointer.GEPOrSubIndex(builder, loc, ElemTy, Index.val); + } + + return EmitCheckedInBoundsPtrOffsetOp(ElemTy, Pointer, Index.val, IsSigned, + IsSubtraction); +} + ValueCategory MLIRScanner::EmitBinAdd(const BinOpInfo &Info) { - auto lhs_v = Info.getLHS().getValue(builder); - auto rhs_v = Info.getRHS().getValue(builder); - if (lhs_v.getType().isa()) { - return ValueCategory(builder.create(loc, lhs_v, rhs_v), - /*isReference*/ false); - } else if (auto mt = lhs_v.getType().dyn_cast()) { - auto shape = std::vector(mt.getShape()); - shape[0] = -1; - auto mt0 = - mlir::MemRefType::get(shape, mt.getElementType(), - MemRefLayoutAttrInterface(), mt.getMemorySpace()); - auto ptradd = rhs_v; - ptradd = castToIndex(loc, ptradd); - return ValueCategory( - builder.create(loc, mt0, lhs_v, ptradd), - /*isReference*/ false); - } else if (auto pt = - lhs_v.getType().dyn_cast()) { - return ValueCategory(builder.create( - loc, pt, lhs_v, std::vector({rhs_v})), - /*isReference*/ false); - } else { - if (auto lhs_c = lhs_v.getDefiningOp()) { - if (auto rhs_c = rhs_v.getDefiningOp()) { - return ValueCategory( - builder.create( - loc, lhs_c.value() + rhs_c.value(), lhs_c.getType()), - false); - } - } - return ValueCategory(builder.create(loc, lhs_v, rhs_v), - /*isReference*/ false); + const auto Loc = getMLIRLocation(Info.getExpr()->getExprLoc()); + const auto LHS = Info.getLHS(); + const auto RHS = Info.getRHS().val; + + if (mlirclang::isPointerOrMemRefTy(LHS.val.getType()) || + mlirclang::isPointerOrMemRefTy(RHS.getType())) { + loc = Loc; + return EmitPointerArithmetic(Info); + } + + if (Info.getType()->isSignedIntegerOrEnumerationType()) { + informNoOverflowCheck( + Glob.getCGM().getLangOpts().getSignedOverflowBehavior(), "add"); + return LHS.Add(builder, Loc, RHS); } + + assert(!Info.getType()->isConstantMatrixType() && "Not yet implemented"); + + if (mlirclang::isFPOrFPVectorTy(LHS.val.getType())) + return LHS.FAdd(builder, Loc, RHS); + return LHS.Add(builder, Loc, RHS); } ValueCategory MLIRScanner::EmitBinSub(const BinOpInfo &Info) { - auto lhs_v = Info.getLHS().getValue(builder); - auto rhs_v = Info.getRHS().getValue(builder); - if (auto mt = lhs_v.getType().dyn_cast()) { - lhs_v = builder.create( - loc, - LLVM::LLVMPointerType::get(mt.getElementType(), - mt.getMemorySpaceAsInt()), - lhs_v); - } - if (auto mt = rhs_v.getType().dyn_cast()) { - rhs_v = builder.create( - loc, - LLVM::LLVMPointerType::get(mt.getElementType(), - mt.getMemorySpaceAsInt()), - rhs_v); - } - if (lhs_v.getType().isa()) { - assert(rhs_v.getType() == lhs_v.getType()); - return ValueCategory(builder.create(loc, lhs_v, rhs_v), - /*isReference*/ false); - } else if (auto pt = - lhs_v.getType().dyn_cast()) { - if (auto IT = rhs_v.getType().dyn_cast()) { - mlir::Value vals[1] = {builder.create( - loc, builder.create(loc, 0, IT.getWidth()), rhs_v)}; - return ValueCategory( - builder.create(loc, lhs_v.getType(), lhs_v, - ArrayRef(vals)), - false); + const auto Loc = getMLIRLocation(Info.getExpr()->getExprLoc()); + auto LHS = Info.getLHS(); + auto RHS = Info.getRHS(); + + // The LHS is always a pointer if either side is. + if (!mlirclang::isPointerOrMemRefTy(LHS.val.getType())) { + if (Info.getType()->isSignedIntegerOrEnumerationType()) { + informNoOverflowCheck( + Glob.getCGM().getLangOpts().getSignedOverflowBehavior(), "sub"); + return LHS.Sub(builder, Loc, RHS.val); } - mlir::Value val = builder.create( - loc, - builder.create( - loc, Glob.getTypes().getMLIRType(Info.getType()), lhs_v), - builder.create( - loc, Glob.getTypes().getMLIRType(Info.getType()), rhs_v)); - val = builder.create( - loc, val, - builder.create( - loc, val.getType(), - builder.create( - loc, builder.getIndexType(), - mlir::TypeAttr::get(pt.getElementType())))); - return ValueCategory(val, /*isReference*/ false); - } else { - return ValueCategory(builder.create(loc, lhs_v, rhs_v), - /*isReference*/ false); + assert(!Info.getType()->isConstantMatrixType() && "Not yet implemented"); + if (mlirclang::isFPOrFPVectorTy(LHS.val.getType())) + return LHS.FSub(builder, Loc, RHS.val); + return LHS.Sub(builder, Loc, RHS.val); + } + + // If the RHS is not a pointer, then we have normal pointer + // arithmetic. + if (!mlirclang::isPointerOrMemRefTy(RHS.val.getType())) { + loc = Loc; + return EmitPointerArithmetic(Info); } + + // Otherwise, this is a pointer subtraction. + + // Do the raw subtraction part. + const auto PtrDiffTy = builder.getIntegerType( + Glob.getCGM().getDataLayout().getPointerSizeInBits()); + LHS = LHS.MemRef2Ptr(builder, Loc).PtrToInt(builder, Loc, PtrDiffTy); + RHS = RHS.MemRef2Ptr(builder, Loc).PtrToInt(builder, Loc, PtrDiffTy); + const auto DiffInChars = LHS.Sub(builder, Loc, RHS.val); + + // Okay, figure out the element size. + const QualType ElementType = + Info.getExpr()->getLHS()->getType()->getPointeeType(); + + assert(!Glob.getCGM().getContext().getAsVariableArrayType(ElementType) && + "Not implemented yet"); + + const CharUnits ElementSize = + (ElementType->isVoidType() || ElementType->isFunctionType()) + ? CharUnits::One() + : Glob.getCGM().getContext().getTypeSizeInChars(ElementType); + + if (ElementSize.isOne()) + return DiffInChars; + + const auto Divisor = builder.createOrFold( + Loc, ElementSize.getQuantity(), PtrDiffTy); + + return DiffInChars.ExactSDiv(builder, Loc, Divisor); } ValueCategory MLIRScanner::EmitBinShl(const BinOpInfo &Info) { diff --git a/polygeist/tools/cgeist/Lib/TypeUtils.cc b/polygeist/tools/cgeist/Lib/TypeUtils.cc index 63d4108cdd833..392a7d4c9487d 100644 --- a/polygeist/tools/cgeist/Lib/TypeUtils.cc +++ b/polygeist/tools/cgeist/Lib/TypeUtils.cc @@ -17,6 +17,7 @@ #include "mlir/Dialect/SYCL/IR/SYCLOps.h" #include "mlir/IR/Types.h" +#include "llvm/ADT/TypeSwitch.h" #include "llvm/IR/DerivedTypes.h" #include "llvm/Support/Casting.h" @@ -76,6 +77,28 @@ mlir::IntegerAttr wrapIntegerMemorySpace(unsigned MemorySpace, : nullptr; } +unsigned getAddressSpace(mlir::Type Ty) { + return llvm::TypeSwitch(Ty) + .Case( + [](auto MemRefTy) { return MemRefTy.getMemorySpaceAsInt(); }) + .Case( + [](auto PtrTy) { return PtrTy.getAddressSpace(); }) + .Default([](auto) -> unsigned { llvm_unreachable("Invalid type"); }); +} + +mlir::Type getPtrTyWithNewType(mlir::Type Orig, mlir::Type NewElementType) { + return llvm::TypeSwitch(Orig) + .Case([NewElementType](auto Ty) { + return mlir::MemRefType::get(Ty.getShape(), NewElementType, + Ty.getLayout(), Ty.getMemorySpace()); + }) + .Case([NewElementType](auto Ty) { + return mlir::LLVM::LLVMPointerType::get(NewElementType, + Ty.getAddressSpace()); + }) + .Default([](auto) -> mlir::Type { llvm_unreachable("Invalid type"); }); +} + mlir::Type getSYCLType(const clang::RecordType *RT, mlirclang::CodeGen::CodeGenTypes &CGT) { const auto *RD = RT->getAsRecordDecl(); @@ -177,4 +200,33 @@ llvm::Type *getLLVMType(const clang::QualType QT, return CGM.getTypes().ConvertType(QT); } +template static bool isTyOrTyVectorTy(mlir::Type Ty) { + if (Ty.isa()) + return true; + const auto VecTy = Ty.dyn_cast(); + return VecTy && VecTy.getElementType().isa(); +} + +bool isFPOrFPVectorTy(mlir::Type Ty) { + return isTyOrTyVectorTy(Ty); +} + +bool isIntOrIntVectorTy(mlir::Type Ty) { + return isTyOrTyVectorTy(Ty); +} + +unsigned getPrimitiveSizeInBits(mlir::Type Ty) { + return llvm::TypeSwitch(Ty) + .Case([](auto IntTy) { return IntTy.getWidth(); }) + .Case([](auto FloatTy) { return FloatTy.getWidth(); }) + .Case( + [](auto) { return mlir::IndexType::kInternalStorageBitWidth; }) + .Case([](auto VecTy) { + return VecTy.getNumElements() * + getPrimitiveSizeInBits(VecTy.getElementType()); + }) + .Default( + [](auto) -> unsigned { llvm_unreachable("Invalid primitive type"); }); +} + } // namespace mlirclang diff --git a/polygeist/tools/cgeist/Lib/TypeUtils.h b/polygeist/tools/cgeist/Lib/TypeUtils.h index 1f9b648e335a3..b15dae864bd89 100644 --- a/polygeist/tools/cgeist/Lib/TypeUtils.h +++ b/polygeist/tools/cgeist/Lib/TypeUtils.h @@ -9,6 +9,9 @@ #ifndef MLIR_TOOLS_MLIRCLANG_TYPE_UTILS_H #define MLIR_TOOLS_MLIRCLANG_TYPE_UTILS_H +#include "mlir/Dialect/LLVMIR/LLVMTypes.h" +#include "mlir/Dialect/SYCL/IR/SYCLOpsTypes.h" +#include "mlir/IR/BuiltinTypes.h" #include "llvm/ADT/SmallPtrSet.h" namespace clang { @@ -46,11 +49,37 @@ bool isRecursiveStruct(llvm::Type *T, llvm::Type *Meta, mlir::IntegerAttr wrapIntegerMemorySpace(unsigned MemorySpace, mlir::MLIRContext *Ctx); +unsigned getAddressSpace(mlir::Type Ty); + +/// Given a MemRefType or LLVMPointerType, change the element type, keeping the +/// rest of the parameters. +mlir::Type getPtrTyWithNewType(mlir::Type Orig, mlir::Type NewElementType); + mlir::Type getSYCLType(const clang::RecordType *RT, mlirclang::CodeGen::CodeGenTypes &CGT); llvm::Type *getLLVMType(clang::QualType QT, clang::CodeGen::CodeGenModule &CGM); +bool isFPOrFPVectorTy(mlir::Type Ty); +bool isIntOrIntVectorTy(mlir::Type Ty); + +inline bool isPointerOrMemRefTy(mlir::Type Ty) { + return Ty.isa(); +} + +inline bool isFirstClassType(mlir::Type Ty) { + return Ty.isa() || + mlir::sycl::isSYCLType(Ty); +} + +inline bool isAggregateType(mlir::Type Ty) { + return Ty.isa() || mlir::sycl::isSYCLType(Ty); +} + +unsigned getPrimitiveSizeInBits(mlir::Type Ty); + } // namespace mlirclang #endif // MLIR_TOOLS_MLIRCLANG_TYPE_UTILS_H diff --git a/polygeist/tools/cgeist/Lib/ValueCategory.cc b/polygeist/tools/cgeist/Lib/ValueCategory.cc index 7b417e84a5112..1b90c54d96e50 100644 --- a/polygeist/tools/cgeist/Lib/ValueCategory.cc +++ b/polygeist/tools/cgeist/Lib/ValueCategory.cc @@ -9,6 +9,7 @@ //===----------------------------------------------------------------------===// #include "ValueCategory.h" +#include "Lib/TypeUtils.h" #include "mlir/Dialect/Arith/IR/Arith.h" #include "mlir/Dialect/Func/IR/FuncOps.h" #include "mlir/Dialect/LLVMIR/LLVMDialect.h" @@ -287,7 +288,7 @@ void ValueCategory::store(mlir::OpBuilder &builder, ValueCategory toStore, } } -template inline void warnUnconstrainedCast() { +template inline void warnUnconstrainedOp() { llvm::WithColor::warning() << "Creating unconstrained " << OpTy::getOperationName() << "\n"; } @@ -302,7 +303,7 @@ ValueCategory ValueCategory::FPTrunc(OpBuilder &Builder, Location Loc, PromotionType.getIntOrFloatBitWidth() && "Source type must be wider than promotion type"); - warnUnconstrainedCast(); + warnUnconstrainedOp(); return Cast(Builder, Loc, PromotionType); } @@ -316,7 +317,7 @@ ValueCategory ValueCategory::FPExt(OpBuilder &Builder, Location Loc, PromotionType.getIntOrFloatBitWidth() && "Source type must be narrower than promotion type"); - warnUnconstrainedCast(); + warnUnconstrainedOp(); return Cast(Builder, Loc, PromotionType); } @@ -326,7 +327,7 @@ ValueCategory ValueCategory::SIToFP(OpBuilder &Builder, Location Loc, assert(PromotionType.isa() && "Expecting floating point promotion type"); - warnUnconstrainedCast(); + warnUnconstrainedOp(); return Cast(Builder, Loc, PromotionType); } @@ -336,7 +337,7 @@ ValueCategory ValueCategory::UIToFP(OpBuilder &Builder, Location Loc, assert(PromotionType.isa() && "Expecting floating point promotion type"); - warnUnconstrainedCast(); + warnUnconstrainedOp(); return Cast(Builder, Loc, PromotionType); } @@ -347,7 +348,7 @@ ValueCategory ValueCategory::FPToUI(OpBuilder &Builder, Location Loc, assert(PromotionType.isa() && "Expecting integer promotion type"); - warnUnconstrainedCast(); + warnUnconstrainedOp(); return Cast(Builder, Loc, PromotionType); } @@ -358,26 +359,35 @@ ValueCategory ValueCategory::FPToSI(OpBuilder &Builder, Location Loc, assert(PromotionType.isa() && "Expecting integer promotion type"); - warnUnconstrainedCast(); + warnUnconstrainedOp(); return Cast(Builder, Loc, PromotionType); } ValueCategory ValueCategory::IntCast(OpBuilder &Builder, Location Loc, Type PromotionType, bool IsSigned) const { - assert(val.getType().isa() && "Expecting integer source type"); - assert(PromotionType.isa() && - "Expecting integer promotion type"); + assert((val.getType().isa()) && + "Expecting integer or index source type"); + assert((PromotionType.isa()) && + "Expecting integer or index promotion type"); if (val.getType() == PromotionType) return *this; - auto SrcIntTy = val.getType().cast(); - auto DstIntTy = PromotionType.cast(); + auto Res = [&]() -> Value { + if (val.getType().isa() || PromotionType.isa()) { + // Special indexcast case + if (IsSigned) + return Builder.createOrFold(Loc, PromotionType, + val); + return Builder.createOrFold(Loc, PromotionType, + val); + } - const unsigned SrcBits = SrcIntTy.getWidth(); - const unsigned DstBits = DstIntTy.getWidth(); + auto SrcIntTy = val.getType().cast(); + auto DstIntTy = PromotionType.cast(); - auto Res = [&]() -> Value { + const unsigned SrcBits = SrcIntTy.getWidth(); + const unsigned DstBits = DstIntTy.getWidth(); if (SrcBits == DstBits) return Builder.createOrFold(Loc, PromotionType, val); if (SrcBits > DstBits) @@ -390,6 +400,74 @@ ValueCategory ValueCategory::IntCast(OpBuilder &Builder, Location Loc, return {Res, /*IsReference*/ false}; } +ValueCategory ValueCategory::PtrToInt(OpBuilder &Builder, Location Loc, + Type DestTy) const { + assert(val.getType().isa() && + "Expecting pointer source type"); + assert(DestTy.isa() && + "Expecting floating point promotion type"); + + return Cast(Builder, Loc, DestTy); +} + +ValueCategory ValueCategory::IntToPtr(OpBuilder &Builder, Location Loc, + Type DestTy) const { + assert(val.getType().isa() && "Expecting pointer source type"); + assert(DestTy.isa() && + "Expecting floating point promotion type"); + + return Cast(Builder, Loc, DestTy); +} + +ValueCategory ValueCategory::BitCast(OpBuilder &Builder, Location Loc, + Type DestTy) const { + assert(mlirclang::isFirstClassType(val.getType()) && + "Expecting first class type"); + assert(mlirclang::isFirstClassType(DestTy) && "Expecting first class type"); + assert(!mlirclang::isAggregateType(val.getType()) && + "Not expecting aggregate type"); + assert(!mlirclang::isAggregateType(DestTy) && "Not expecting aggregate type"); + assert((!mlirclang::isPointerOrMemRefTy(val.getType()) || + mlirclang::isPointerOrMemRefTy(DestTy)) && + "Cannot cast pointers to anything but pointers"); + assert((mlirclang::isPointerOrMemRefTy(val.getType()) || + mlirclang::getPrimitiveSizeInBits(val.getType()) == + mlirclang::getPrimitiveSizeInBits(DestTy)) && + "Expecting equal bitwidth"); + assert((!mlirclang::isPointerOrMemRefTy(val.getType()) || + mlirclang::getAddressSpace(val.getType()) == + mlirclang::getAddressSpace(DestTy)) && + "Expecting equal address spaces"); + assert((!(val.getType().isa() && + DestTy.isa()) || + val.getType().cast().getNumElements() == + DestTy.cast().getNumElements()) && + "Expecting same number of elements"); + assert((!val.getType().isa() || + val.getType().cast().getNumElements() == 1) && + "Expecting single-element vector"); + assert((!DestTy.isa() || + DestTy.cast().getNumElements() == 1) && + "Expecting single-element vector"); + + return Cast(Builder, Loc, DestTy); +} + +ValueCategory ValueCategory::MemRef2Ptr(OpBuilder &Builder, + Location Loc) const { + const auto Ty = val.getType().dyn_cast(); + if (!Ty) { + assert(val.getType().isa() && + "Expecting pointer type"); + return *this; + } + + auto DestTy = + LLVM::LLVMPointerType::get(Ty.getElementType(), Ty.getMemorySpaceAsInt()); + return {Builder.createOrFold(Loc, DestTy, val), + isReference}; +} + ValueCategory ValueCategory::ICmpNE(mlir::OpBuilder &builder, Location Loc, mlir::Value RHS) const { return ICmp(builder, Loc, arith::CmpIPredicate::ne, RHS); @@ -414,6 +492,143 @@ ValueCategory ValueCategory::FCmp(OpBuilder &builder, Location Loc, mlir::Value RHS) const { assert(val.getType() == RHS.getType() && "Cannot compare values of different types"); - assert(val.getType().isa() && "Expecting floatint point inputs"); + assert(val.getType().isa() && "Expecting floating point inputs"); return {builder.createOrFold(Loc, predicate, val, RHS), false}; } + +template +static ValueCategory IntBinOp(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value LHS, mlir::Value RHS) { + assert(LHS.getType() == RHS.getType() && + "Cannot operate on values of different types"); + assert(mlirclang::isIntOrIntVectorTy(LHS.getType()) && + "Expecting integers or integer vectors as inputs"); + return {Builder.createOrFold(Loc, LHS, RHS), false}; +} + +template +static ValueCategory FPBinOp(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value LHS, mlir::Value RHS) { + assert(LHS.getType() == RHS.getType() && + "Cannot operate on values of different types"); + assert(mlirclang::isFPOrFPVectorTy(LHS.getType()) && + "Expecting integers or integer vectors as inputs"); + + warnUnconstrainedOp(); + + return {Builder.createOrFold(Loc, LHS, RHS), false}; +} + +template +static ValueCategory NUWNSWBinOp(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value LHS, mlir::Value RHS, bool HasNUW, + bool HasNSW) { + // No way of adding these flags to MLIR. + if (HasNUW) + llvm::WithColor::warning() << "Not adding NUW flag.\n"; + if (HasNSW) + llvm::WithColor::warning() << "Not adding NSW flag.\n"; + return IntBinOp(Builder, Loc, LHS, RHS); +} + +ValueCategory ValueCategory::SDiv(OpBuilder &Builder, Location Loc, Value RHS, + bool IsExact) const { + if (IsExact) + llvm::WithColor::warning() << "Creating exact division is not supported\n"; + return IntBinOp(Builder, Loc, val, RHS); +} + +ValueCategory ValueCategory::ExactSDiv(OpBuilder &Builder, Location Loc, + Value RHS) const { + return SDiv(Builder, Loc, RHS, /*IsExact*/ true); +} + +ValueCategory ValueCategory::Neg(OpBuilder &Builder, Location Loc, bool HasNUW, + bool HasNSW) const { + ValueCategory Zero(Builder.createOrFold(Loc, 0, val.getType()), + /*IsReference*/ false); + return Zero.Sub(Builder, Loc, val, HasNUW, HasNSW); +} + +ValueCategory ValueCategory::Add(OpBuilder &Builder, Location Loc, Value RHS, + bool HasNUW, bool HasNSW) const { + return NUWNSWBinOp(Builder, Loc, val, RHS, HasNUW, HasNSW); +} + +ValueCategory ValueCategory::FAdd(OpBuilder &Builder, Location Loc, + Value RHS) const { + return FPBinOp(Builder, Loc, val, RHS); +} + +ValueCategory ValueCategory::Sub(OpBuilder &Builder, Location Loc, Value RHS, + bool HasNUW, bool HasNSW) const { + return NUWNSWBinOp(Builder, Loc, val, RHS, HasNUW, HasNSW); +} + +ValueCategory ValueCategory::FSub(OpBuilder &Builder, Location Loc, + Value RHS) const { + return FPBinOp(Builder, Loc, val, RHS); +} + +ValueCategory ValueCategory::SubIndex(OpBuilder &Builder, Location Loc, + Type Type, Value Index, + bool IsInBounds) const { + assert(val.getType().isa() && "Expecting a pointer as operand"); + assert(Index.getType().isa() && "Expecting an index type index"); + + if (IsInBounds) { + llvm::WithColor::warning() + << "Cannot create an inbounds SubIndex operation\n"; + } + auto PtrTy = mlirclang::getPtrTyWithNewType(val.getType(), Type); + return {Builder.createOrFold(Loc, PtrTy, val, Index), + isReference}; +} + +ValueCategory ValueCategory::InBoundsSubIndex(OpBuilder &Builder, Location Loc, + Type Type, Value Index) const { + return SubIndex(Builder, Loc, Type, Index, /*IsInBounds*/ true); +} + +ValueCategory ValueCategory::GEP(OpBuilder &Builder, Location Loc, Type Type, + ValueRange IdxList, bool IsInBounds) const { + assert(val.getType().isa() && + "Expecting a pointer as operand"); + assert(std::all_of(IdxList.getType().begin(), IdxList.getType().end(), + [](mlir::Type Ty) { return Ty.isa(); }) && + "Expecting integer indices"); + + if (IsInBounds) + llvm::WithColor::warning() << "Cannot create an inbounds GEP operation\n"; + auto PtrTy = mlirclang::getPtrTyWithNewType(val.getType(), Type); + return {Builder.createOrFold(Loc, PtrTy, val, IdxList), + isReference}; +} + +ValueCategory ValueCategory::InBoundsGEP(OpBuilder &Builder, Location Loc, + Type Type, ValueRange IdxList) const { + return GEP(Builder, Loc, Type, IdxList, /*IsInBounds*/ true); +} + +ValueCategory ValueCategory::GEPOrSubIndex(OpBuilder &Builder, Location Loc, + Type Type, ValueRange IdxList, + bool IsInBounds) const { + const auto ValType = val.getType(); + assert((ValType.isa()) && + "Expecting an LLVMPointer or MemRefType input"); + + return llvm::TypeSwitch(ValType) + .Case([&](auto) { + assert(IdxList.size() == 1 && "SubIndexOp expects a single index"); + return SubIndex(Builder, Loc, Type, IdxList[0], IsInBounds); + }) + .Case( + [&](auto) { return GEP(Builder, Loc, Type, IdxList, IsInBounds); }) + .Default([](auto) -> ValueCategory { llvm_unreachable("Invalid type"); }); +} + +ValueCategory ValueCategory::InBoundsGEPOrSubIndex(OpBuilder &Builder, + Location Loc, Type Type, + ValueRange IdxList) const { + return GEPOrSubIndex(Builder, Loc, Type, IdxList, /*IsInBounds*/ true); +} diff --git a/polygeist/tools/cgeist/Lib/ValueCategory.h b/polygeist/tools/cgeist/Lib/ValueCategory.h index 0ffff1017e7ac..7da481662c1c0 100644 --- a/polygeist/tools/cgeist/Lib/ValueCategory.h +++ b/polygeist/tools/cgeist/Lib/ValueCategory.h @@ -48,6 +48,25 @@ class ValueCategory { void store(mlir::OpBuilder &Builder, mlir::Value toStore) const; ValueCategory dereference(mlir::OpBuilder &Builder) const; + ValueCategory SubIndex(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type Type, mlir::Value Index, + bool IsInBounds = false) const; + ValueCategory InBoundsSubIndex(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type Type, mlir::Value Index) const; + + ValueCategory GEP(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type Type, mlir::ValueRange IdxList, + bool IsInBounds = false) const; + ValueCategory InBoundsGEP(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type Type, mlir::ValueRange IdxList) const; + + ValueCategory GEPOrSubIndex(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type Type, mlir::ValueRange IdxList, + bool IsInBounds = false) const; + ValueCategory InBoundsGEPOrSubIndex(mlir::OpBuilder &Builder, + mlir::Location Loc, mlir::Type Type, + mlir::ValueRange IdxList) const; + ValueCategory FPTrunc(mlir::OpBuilder &Builder, mlir::Location Loc, mlir::Type PromotionType) const; @@ -63,11 +82,35 @@ class ValueCategory { mlir::Type PromotionType) const; ValueCategory FPToSI(mlir::OpBuilder &Builder, mlir::Location Loc, mlir::Type PromotionType) const; + ValueCategory PtrToInt(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type DestTy) const; + ValueCategory IntToPtr(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type DestTy) const; + ValueCategory BitCast(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Type DestTy) const; + ValueCategory MemRef2Ptr(mlir::OpBuilder &Builder, mlir::Location Loc) const; ValueCategory ICmpNE(mlir::OpBuilder &builder, mlir::Location Loc, mlir::Value RHS) const; ValueCategory FCmpUNE(mlir::OpBuilder &builder, mlir::Location Loc, mlir::Value RHS) const; + + ValueCategory SDiv(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value RHS, bool IsExact = false) const; + ValueCategory ExactSDiv(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value RHS) const; + ValueCategory Neg(mlir::OpBuilder &Builder, mlir::Location Loc, + bool HasNUW = false, bool HasNSW = false) const; + ValueCategory Add(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value RHS, bool HasNUW = false, + bool HasNSW = false) const; + ValueCategory FAdd(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value RHS) const; + ValueCategory Sub(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value RHS, bool HasNUW = false, + bool HasNSW = false) const; + ValueCategory FSub(mlir::OpBuilder &Builder, mlir::Location Loc, + mlir::Value RHS) const; }; #endif /* CLANG_MLIR_VALUE_CATEGORY */ diff --git a/polygeist/tools/cgeist/Lib/clang-mlir.h b/polygeist/tools/cgeist/Lib/clang-mlir.h index 8f7acf8e71915..7b05026403463 100644 --- a/polygeist/tools/cgeist/Lib/clang-mlir.h +++ b/polygeist/tools/cgeist/Lib/clang-mlir.h @@ -353,6 +353,8 @@ class MLIRScanner : public clang::StmtVisitor { // Reshape memref to memref<1 x elemTy>. mlir::Value reshapeRanklessGlobal(mlir::memref::GetGlobalOp GV); + ValueCategory CastToVoidPtr(ValueCategory Ptr); + /// TODO: Add ScalarConversion options ValueCategory EmitScalarCast(mlir::Location Loc, ValueCategory Src, clang::QualType SrcType, clang::QualType DstType, @@ -505,6 +507,13 @@ class MLIRScanner : public clang::StmtVisitor { EmitCompoundAssign(clang::CompoundAssignOperator *E, ValueCategory (MLIRScanner::*F)(const BinOpInfo &)); + ValueCategory EmitCheckedInBoundsPtrOffsetOp(mlir::Type ElemTy, + ValueCategory Pointer, + mlir::ValueRange IdxList, + bool IsSigned, + bool IsSubtraction); + ValueCategory EmitPointerArithmetic(const BinOpInfo &Info); + BinOpInfo EmitBinOps(clang::BinaryOperator *E, clang::QualType PromotionTy = clang::QualType()); #define HANDLEBINOP(OP) \ diff --git a/polygeist/tools/cgeist/Test/Verification/add.c b/polygeist/tools/cgeist/Test/Verification/add.c new file mode 100644 index 0000000000000..b4a4a6aecbe41 --- /dev/null +++ b/polygeist/tools/cgeist/Test/Verification/add.c @@ -0,0 +1,150 @@ +// RUN: cgeist -O0 %s --function=* -S | FileCheck %s + +typedef char char_vec __attribute__((ext_vector_type(3))); +typedef short short_vec __attribute__((ext_vector_type(3))); +typedef int int_vec __attribute__((ext_vector_type(3))); +typedef long long_vec __attribute__((ext_vector_type(3))); +typedef float float_vec __attribute__((ext_vector_type(3))); +typedef double double_vec __attribute__((ext_vector_type(3))); + +// CHECK-LABEL: func.func @add_i8( +// CHECK-SAME: %[[VAL_0:.*]]: i8, +// CHECK-SAME: %[[VAL_1:.*]]: i8) -> i8 +// CHECK: %[[VAL_2:.*]] = arith.extsi %[[VAL_0]] : i8 to i32 +// CHECK: %[[VAL_3:.*]] = arith.extsi %[[VAL_1]] : i8 to i32 +// CHECK: %[[VAL_4:.*]] = arith.addi %[[VAL_2]], %[[VAL_3]] : i32 +// CHECK: %[[VAL_5:.*]] = arith.trunci %[[VAL_4]] : i32 to i8 +// CHECK: return %[[VAL_5]] : i8 +// CHECK: } + +char add_i8(char a, char b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_i16( +// CHECK-SAME: %[[VAL_0:.*]]: i16, +// CHECK-SAME: %[[VAL_1:.*]]: i16) -> i16 +// CHECK: %[[VAL_2:.*]] = arith.extsi %[[VAL_0]] : i16 to i32 +// CHECK: %[[VAL_3:.*]] = arith.extsi %[[VAL_1]] : i16 to i32 +// CHECK: %[[VAL_4:.*]] = arith.addi %[[VAL_2]], %[[VAL_3]] : i32 +// CHECK: %[[VAL_5:.*]] = arith.trunci %[[VAL_4]] : i32 to i16 +// CHECK: return %[[VAL_5]] : i16 +// CHECK: } + +short add_i16(short a, short b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_i32( +// CHECK-SAME: %[[VAL_0:.*]]: i32, +// CHECK-SAME: %[[VAL_1:.*]]: i32) -> i32 +// CHECK: %[[ADD:.*]] = arith.addi %[[VAL_0]], %[[VAL_1]] : i32 +// CHECK: return %[[ADD]] : i32 +// CHECK: } + +int add_i32(int a, int b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_i64( +// CHECK-SAME: %[[VAL_0:.*]]: i64, +// CHECK-SAME: %[[VAL_1:.*]]: i64) -> i64 +// CHECK: %[[VAL_2:.*]] = arith.addi %[[VAL_0]], %[[VAL_1]] : i64 +// CHECK: return %[[VAL_2]] : i64 +// CHECK: } + +long add_i64(long a, long b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_f32( +// CHECK-SAME: %[[VAL_0:.*]]: f32, +// CHECK-SAME: %[[VAL_1:.*]]: f32) -> f32 +// CHECK: %[[ADD:.*]] = arith.addf %[[VAL_0]], %[[VAL_1]] : f32 +// CHECK: return %[[ADD]] : f32 +// CHECK: } + +float add_f32(float a, float b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_f64( +// CHECK-SAME: %[[VAL_0:.*]]: f32, +// CHECK-SAME: %[[VAL_1:.*]]: f32) -> f32 +// CHECK: %[[VAL_2:.*]] = arith.addf %[[VAL_0]], %[[VAL_1]] : f32 +// CHECK: return %[[VAL_2]] : f32 +// CHECK: } + +float add_f64(float a, float b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_vi8( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xi8>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xi8>) -> vector<3xi8> +// CHECK: %[[ADD:.*]] = arith.addi %[[VAL_0]], %[[VAL_1]] : vector<3xi8> +// CHECK: return %[[ADD]] : vector<3xi8> +// CHECK: } + +char_vec add_vi8(char_vec a, char_vec b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_vi16( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xi16>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xi16>) -> vector<3xi16> +// CHECK: %[[VAL_2:.*]] = arith.addi %[[VAL_0]], %[[VAL_1]] : vector<3xi16> +// CHECK: return %[[VAL_2]] : vector<3xi16> +// CHECK: } + +short_vec add_vi16(short_vec a, short_vec b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_vi32( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xi32>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xi32>) -> vector<3xi32> +// CHECK: %[[ADD:.*]] = arith.addi %[[VAL_0]], %[[VAL_1]] : vector<3xi32> +// CHECK: return %[[ADD]] : vector<3xi32> +// CHECK: } + +int_vec add_vi32(int_vec a, int_vec b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_vi64( +// CHECK-SAME: %[[VAL_0:.*]]: memref>, +// CHECK-SAME: %[[VAL_1:.*]]: memref>) -> vector<3xi64> +// CHECK: %[[VAL_2:.*]] = affine.load %[[VAL_0]][0] : memref> +// CHECK: %[[VAL_3:.*]] = affine.load %[[VAL_1]][0] : memref> +// CHECK: %[[VAL_4:.*]] = arith.addi %[[VAL_2]], %[[VAL_3]] : vector<3xi64> +// CHECK: return %[[VAL_4]] : vector<3xi64> +// CHECK: } + +long_vec add_vi64(long_vec a, long_vec b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_vf32( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xf32>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xf32>) -> vector<3xf32> +// CHECK: %[[ADD:.*]] = arith.addf %[[VAL_0]], %[[VAL_1]] : vector<3xf32> +// CHECK: return %[[ADD]] : vector<3xf32> +// CHECK: } + +float_vec add_vf32(float_vec a, float_vec b) { + return a + b; +} + +// CHECK-LABEL: func.func @add_vf64( +// CHECK-SAME: %[[VAL_0:.*]]: memref>, +// CHECK-SAME: %[[VAL_1:.*]]: memref>) -> vector<3xf64> +// CHECK: %[[VAL_2:.*]] = affine.load %[[VAL_0]][0] : memref> +// CHECK: %[[VAL_3:.*]] = affine.load %[[VAL_1]][0] : memref> +// CHECK: %[[VAL_4:.*]] = arith.addf %[[VAL_2]], %[[VAL_3]] : vector<3xf64> +// CHECK: return %[[VAL_4]] : vector<3xf64> +// CHECK: } + +double_vec add_vf64(double_vec a, double_vec b) { + return a + b; +} diff --git a/polygeist/tools/cgeist/Test/Verification/ptrarith.c b/polygeist/tools/cgeist/Test/Verification/ptrarith.c new file mode 100644 index 0000000000000..3bdc3d8c8fb03 --- /dev/null +++ b/polygeist/tools/cgeist/Test/Verification/ptrarith.c @@ -0,0 +1,76 @@ +// RUN: cgeist %s --function=* -S | FileCheck %s + +#include + +// CHECK-LABEL: func.func @f0( +// CHECK-SAME: %[[VAL_0:.*]]: memref, +// CHECK-SAME: %[[VAL_1:.*]]: i32) -> memref +// CHECK: %[[INDEX:.*]] = arith.index_cast %[[VAL_1]] : i32 to index +// CHECK: %[[ADD:.*]] = "polygeist.subindex"(%[[VAL_0]], %[[INDEX]]) : (memref, index) -> memref +// CHECK: return %[[ADD]] : memref +// CHECK: } + +int *f0(int *ptr, int index) { + return ptr + index; +} + +// CHECK-LABEL: func.func @f1( +// CHECK-SAME: %[[VAL_0:.*]]: i64, +// CHECK-SAME: %[[VAL_1:.*]]: memref) -> memref +// CHECK: %[[INDEX:.*]] = arith.index_castui %[[VAL_0]] : i64 to index +// CHECK: %[[ADD:.*]] = "polygeist.subindex"(%[[VAL_1]], %[[INDEX]]) : (memref, index) -> memref +// CHECK: return %[[ADD]] : memref +// CHECK: } + +int *f1(size_t index, int *ptr) { + return index + ptr; +} + +// CHECK-LABEL: func.func @f2( +// CHECK-SAME: %[[VAL_0:.*]]: i64) -> !llvm.ptr +// CHECK: %[[PTR_0:.*]] = llvm.inttoptr %[[VAL_0]] : i64 to !llvm.ptr +// CHECK: return %[[PTR_0]] : !llvm.ptr +// CHECK: } + +void *f2(size_t index) { + return ((char*) NULL) + index; +} + +// CHECK-LABEL: func.func @f3( +// CHECK-SAME: %[[VAL_0:.*]]: memref, +// CHECK-SAME: %[[VAL_1:.*]]: i64) -> memref +// CHECK: %[[I64_0:.*]] = arith.constant 0 : i64 +// CHECK: %[[NEG:.*]] = arith.subi %[[I64_0]], %[[VAL_1]] : i64 +// CHECK: %[[INDEX:.*]] = arith.index_castui %[[NEG]] : i64 to index +// CHECK: %[[ADD:.*]] = "polygeist.subindex"(%[[VAL_0]], %[[INDEX]]) : (memref, index) -> memref +// CHECK: return %[[ADD]] : memref +// CHECK: } + +int *f3(int *ptr, size_t index) { + return ptr - index; +} + +// CHECK-LABEL: func.func @f4( +// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr, +// CHECK-SAME: %[[VAL_1:.*]]: i64) -> !llvm.ptr +// CHECK: %[[ADD:.*]] = llvm.getelementptr %[[VAL_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr +// CHECK: return %[[ADD]] : !llvm.ptr +// CHECK: } + +void *f4(void *ptr, size_t index) { + return ptr + index; +} + +// CHECK-LABEL: func.func @f5( +// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr>, +// CHECK-SAME: %[[VAL_1:.*]]: i64) -> i32 +// CHECK: %[[VOIDPTR_0:.*]] = llvm.bitcast %[[VAL_0]] : !llvm.ptr> to !llvm.ptr +// CHECK: %[[PTR:.*]] = llvm.getelementptr %[[VOIDPTR_0]]{{\[}}%[[VAL_1]]] : (!llvm.ptr, i64) -> !llvm.ptr +// CHECK: %[[FUNCPTR:.*]] = llvm.bitcast %[[PTR]] : !llvm.ptr to !llvm.ptr> +// CHECK: %[[RES:.*]] = llvm.call %[[FUNCPTR]]() : () -> i32 +// CHECK: return %[[RES]] : i32 +// CHECK: } + +int f5(int (*ptr)(void), size_t index) { + return (ptr + index)(); +} diff --git a/polygeist/tools/cgeist/Test/Verification/sub.c b/polygeist/tools/cgeist/Test/Verification/sub.c new file mode 100644 index 0000000000000..d3a38191492c6 --- /dev/null +++ b/polygeist/tools/cgeist/Test/Verification/sub.c @@ -0,0 +1,182 @@ +// RUN: cgeist -O0 %s --function=* -S | FileCheck %s + +#include + +typedef char char_vec __attribute__((ext_vector_type(3))); +typedef short short_vec __attribute__((ext_vector_type(3))); +typedef int int_vec __attribute__((ext_vector_type(3))); +typedef long long_vec __attribute__((ext_vector_type(3))); +typedef float float_vec __attribute__((ext_vector_type(3))); +typedef double double_vec __attribute__((ext_vector_type(3))); + +// CHECK-LABEL: func.func @sub_i8( +// CHECK-SAME: %[[VAL_0:.*]]: i8, +// CHECK-SAME: %[[VAL_1:.*]]: i8) -> i8 +// CHECK: %[[VAL_2:.*]] = arith.extsi %[[VAL_0]] : i8 to i32 +// CHECK: %[[VAL_3:.*]] = arith.extsi %[[VAL_1]] : i8 to i32 +// CHECK: %[[VAL_4:.*]] = arith.subi %[[VAL_2]], %[[VAL_3]] : i32 +// CHECK: %[[VAL_5:.*]] = arith.trunci %[[VAL_4]] : i32 to i8 +// CHECK: return %[[VAL_5]] : i8 +// CHECK: } + +char sub_i8(char a, char b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_i16( +// CHECK-SAME: %[[VAL_0:.*]]: i16, +// CHECK-SAME: %[[VAL_1:.*]]: i16) -> i16 +// CHECK: %[[VAL_2:.*]] = arith.extsi %[[VAL_0]] : i16 to i32 +// CHECK: %[[VAL_3:.*]] = arith.extsi %[[VAL_1]] : i16 to i32 +// CHECK: %[[VAL_4:.*]] = arith.subi %[[VAL_2]], %[[VAL_3]] : i32 +// CHECK: %[[VAL_5:.*]] = arith.trunci %[[VAL_4]] : i32 to i16 +// CHECK: return %[[VAL_5]] : i16 +// CHECK: } + +short sub_i16(short a, short b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_i32( +// CHECK-SAME: %[[VAL_0:.*]]: i32, +// CHECK-SAME: %[[VAL_1:.*]]: i32) -> i32 +// CHECK: %[[VAL_2:.*]] = arith.subi %[[VAL_0]], %[[VAL_1]] : i32 +// CHECK: return %[[VAL_2]] : i32 +// CHECK: } + +int sub_i32(int a, int b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_i64( +// CHECK-SAME: %[[VAL_0:.*]]: i64, +// CHECK-SAME: %[[VAL_1:.*]]: i64) -> i64 +// CHECK: %[[VAL_2:.*]] = arith.subi %[[VAL_0]], %[[VAL_1]] : i64 +// CHECK: return %[[VAL_2]] : i64 +// CHECK: } + +long sub_i64(long a, long b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_f32( +// CHECK-SAME: %[[VAL_0:.*]]: f32, +// CHECK-SAME: %[[VAL_1:.*]]: f32) -> f32 +// CHECK: %[[SUB:.*]] = arith.subf %[[VAL_0]], %[[VAL_1]] : f32 +// CHECK: return %[[SUB]] : f32 +// CHECK: } + +float sub_f32(float a, float b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_f64( +// CHECK-SAME: %[[VAL_0:.*]]: f32, +// CHECK-SAME: %[[VAL_1:.*]]: f32) -> f32 +// CHECK: %[[VAL_2:.*]] = arith.subf %[[VAL_0]], %[[VAL_1]] : f32 +// CHECK: return %[[VAL_2]] : f32 +// CHECK: } + +float sub_f64(float a, float b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_vi8( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xi8>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xi8>) -> vector<3xi8> +// CHECK: %[[VAL_2:.*]] = arith.subi %[[VAL_0]], %[[VAL_1]] : vector<3xi8> +// CHECK: return %[[VAL_2]] : vector<3xi8> +// CHECK: } + +char_vec sub_vi8(char_vec a, char_vec b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_vi16( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xi16>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xi16>) -> vector<3xi16> +// CHECK: %[[VAL_2:.*]] = arith.subi %[[VAL_0]], %[[VAL_1]] : vector<3xi16> +// CHECK: return %[[VAL_2]] : vector<3xi16> +// CHECK: } + +short_vec sub_vi16(short_vec a, short_vec b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_vi32( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xi32>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xi32>) -> vector<3xi32> +// CHECK: %[[VAL_2:.*]] = arith.subi %[[VAL_0]], %[[VAL_1]] : vector<3xi32> +// CHECK: return %[[VAL_2]] : vector<3xi32> +// CHECK: } + +int_vec sub_vi32(int_vec a, int_vec b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_vi64( +// CHECK-SAME: %[[VAL_0:.*]]: memref>, +// CHECK-SAME: %[[VAL_1:.*]]: memref>) -> vector<3xi64> +// CHECK: %[[VAL_2:.*]] = affine.load %[[VAL_0]][0] : memref> +// CHECK: %[[VAL_3:.*]] = affine.load %[[VAL_1]][0] : memref> +// CHECK: %[[VAL_4:.*]] = arith.subi %[[VAL_2]], %[[VAL_3]] : vector<3xi64> +// CHECK: return %[[VAL_4]] : vector<3xi64> +// CHECK: } + +long_vec sub_vi64(long_vec a, long_vec b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_vf32( +// CHECK-SAME: %[[VAL_0:.*]]: vector<3xf32>, +// CHECK-SAME: %[[VAL_1:.*]]: vector<3xf32>) -> vector<3xf32> +// CHECK: %[[SUB:.*]] = arith.subf %[[VAL_0]], %[[VAL_1]] : vector<3xf32> +// CHECK: return %[[SUB]] : vector<3xf32> +// CHECK: } + +float_vec sub_vf32(float_vec a, float_vec b) { + return a - b; +} + +// CHECK-LABEL: func.func @sub_vf64( +// CHECK-SAME: %[[VAL_0:.*]]: memref>, +// CHECK-SAME: %[[VAL_1:.*]]: memref>) -> vector<3xf64> +// CHECK: %[[VAL_2:.*]] = affine.load %[[VAL_0]][0] : memref> +// CHECK: %[[VAL_3:.*]] = affine.load %[[VAL_1]][0] : memref> +// CHECK: %[[VAL_4:.*]] = arith.subf %[[VAL_2]], %[[VAL_3]] : vector<3xf64> +// CHECK: return %[[VAL_4]] : vector<3xf64> +// CHECK: } + +double_vec sub_vf64(double_vec a, double_vec b) { + return a - b; +} + +// CHECK-LABEL: func.func @ptr_diff_i8( +// CHECK-SAME: %[[VAL_0:.*]]: !llvm.ptr, +// CHECK-SAME: %[[VAL_1:.*]]: !llvm.ptr) -> i64 +// CHECK: %[[VAL_2:.*]] = llvm.ptrtoint %[[VAL_0]] : !llvm.ptr to i64 +// CHECK: %[[VAL_3:.*]] = llvm.ptrtoint %[[VAL_1]] : !llvm.ptr to i64 +// CHECK: %[[VAL_4:.*]] = arith.subi %[[VAL_2]], %[[VAL_3]] : i64 +// CHECK: return %[[VAL_4]] : i64 +// CHECK: } + +size_t ptr_diff_i8(char *a, char *b) { + return a - b; +} + +// CHECK-LABEL: func.func @ptr_diff_f32( +// CHECK-SAME: %[[VAL_0:.*]]: memref, +// CHECK-SAME: %[[VAL_1:.*]]: memref) -> i64 +// CHECK: %[[I64_0:.*]] = arith.constant 4 : i64 +// CHECK: %[[PTR_0:.*]] = "polygeist.memref2pointer"(%[[VAL_0]]) : (memref) -> !llvm.ptr +// CHECK: %[[INT_0:.*]] = llvm.ptrtoint %[[PTR_0]] : !llvm.ptr to i64 +// CHECK: %[[PTR_1:.*]] = "polygeist.memref2pointer"(%[[VAL_1]]) : (memref) -> !llvm.ptr +// CHECK: %[[INT_1:.*]] = llvm.ptrtoint %[[PTR_1]] : !llvm.ptr to i64 +// CHECK: %[[DIFF:.*]] = arith.subi %[[INT_0]], %[[INT_1]] : i64 +// CHECK: %[[SUB:.*]] = arith.divsi %[[DIFF]], %[[I64_0]] : i64 +// CHECK: return %[[SUB]] : i64 +// CHECK: } + +size_t ptr_diff_f32(float *a, float *b) { + return a - b; +}