llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-flang-fir-hlfir Author: Ming Yan (NexMing) <details> <summary>Changes</summary> This patch begins a long-term effort to establish an incremental FIR to MLIR lowering path while preserving the existing FIR to LLVM pipeline. It introduces the core conversion infrastructure and demonstrates the approach with example conversions such as convert `fir.load` to `memref.load`. The new lowering can coexist with the current pipeline, enabling mixed-mode lowering without affecting existing behavior. Future patches will extend the coverage to more operations and types. --- Patch is 34.87 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/168703.diff 16 Files Affected: - (modified) clang/include/clang/Options/Options.td (+5) - (modified) clang/lib/Driver/ToolChains/Flang.cpp (+1) - (modified) flang/include/flang/Lower/LoweringOptions.def (+3) - (modified) flang/include/flang/Optimizer/Passes/Pipelines.h (+3) - (modified) flang/include/flang/Optimizer/Support/InitFIR.h (+8-1) - (modified) flang/include/flang/Optimizer/Transforms/Passes.td (+11) - (modified) flang/include/flang/Tools/CrossToolHelpers.h (+1) - (modified) flang/lib/Frontend/CompilerInvocation.cpp (+4) - (modified) flang/lib/Frontend/FrontendActions.cpp (+1) - (modified) flang/lib/Optimizer/CodeGen/CodeGen.cpp (+35-2) - (modified) flang/lib/Optimizer/Passes/Pipelines.cpp (+18-5) - (modified) flang/lib/Optimizer/Transforms/CMakeLists.txt (+1) - (added) flang/lib/Optimizer/Transforms/ConvertFIRToMLIR.cpp (+238) - (modified) flang/test/Driver/frontend-forwarding.f90 (+2) - (modified) flang/test/Fir/convert-to-llvm.fir (+75) - (added) flang/test/Fir/convert-to-mlir.fir (+135) ``````````diff diff --git a/clang/include/clang/Options/Options.td b/clang/include/clang/Options/Options.td index cda11fdc94230..cd8409de8c5a9 100644 --- a/clang/include/clang/Options/Options.td +++ b/clang/include/clang/Options/Options.td @@ -7217,6 +7217,11 @@ def flang_deprecated_no_hlfir : Flag<["-"], "flang-deprecated-no-hlfir">, Flags<[HelpHidden]>, Visibility<[FlangOption, FC1Option]>, HelpText<"Do not use HLFIR lowering (deprecated)">; +def flang_experimental_lower_through_mlir + : Flag<["-"], "flang-experimental-lower-through-mlir">, + Flags<[HelpHidden]>, Visibility<[FlangOption, FC1Option]>, + HelpText<"Lower form FIR through MLIR to LLVM (experimental)">; + //===----------------------------------------------------------------------===// // FLangOption + CoreOption + NoXarchOption //===----------------------------------------------------------------------===// diff --git a/clang/lib/Driver/ToolChains/Flang.cpp b/clang/lib/Driver/ToolChains/Flang.cpp index 270904de544d6..e294ac59af73d 100644 --- a/clang/lib/Driver/ToolChains/Flang.cpp +++ b/clang/lib/Driver/ToolChains/Flang.cpp @@ -222,6 +222,7 @@ void Flang::addCodegenOptions(const ArgList &Args, {options::OPT_fdo_concurrent_to_openmp_EQ, options::OPT_flang_experimental_hlfir, options::OPT_flang_deprecated_no_hlfir, + options::OPT_flang_experimental_lower_through_mlir, options::OPT_fno_ppc_native_vec_elem_order, options::OPT_fppc_native_vec_elem_order, options::OPT_finit_global_zero, options::OPT_fno_init_global_zero, options::OPT_frepack_arrays, diff --git a/flang/include/flang/Lower/LoweringOptions.def b/flang/include/flang/Lower/LoweringOptions.def index 39f197d8d35c8..01fc96b78df50 100644 --- a/flang/include/flang/Lower/LoweringOptions.def +++ b/flang/include/flang/Lower/LoweringOptions.def @@ -38,6 +38,9 @@ ENUM_LOWERINGOPT(Underscoring, unsigned, 1, 1) /// (i.e. wraps around as two's complement). Off by default. ENUM_LOWERINGOPT(IntegerWrapAround, unsigned, 1, 0) +/// If true, lower form FIR through MLIR to LLVM +ENUM_LOWERINGOPT(LowerThroughMLIR, unsigned, 1, 0) + /// If true (default), follow Fortran 2003 rules for (re)allocating /// the allocatable on the left side of the intrinsic assignment, /// if LHS and RHS have mismatching shapes/types. diff --git a/flang/include/flang/Optimizer/Passes/Pipelines.h b/flang/include/flang/Optimizer/Passes/Pipelines.h index 70b9341347244..e1e990bfd1c01 100644 --- a/flang/include/flang/Optimizer/Passes/Pipelines.h +++ b/flang/include/flang/Optimizer/Passes/Pipelines.h @@ -18,10 +18,13 @@ #include "flang/Optimizer/Passes/CommandLineOpts.h" #include "flang/Optimizer/Transforms/Passes.h" #include "flang/Tools/CrossToolHelpers.h" +#include "mlir/Conversion/AffineToStandard/AffineToStandard.h" +#include "mlir/Conversion/MemRefToLLVM/MemRefToLLVM.h" #include "mlir/Conversion/ReconcileUnrealizedCasts/ReconcileUnrealizedCasts.h" #include "mlir/Conversion/SCFToControlFlow/SCFToControlFlow.h" #include "mlir/Dialect/GPU/IR/GPUDialect.h" #include "mlir/Dialect/LLVMIR/LLVMAttrs.h" +#include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Dialect/OpenMP/Transforms/Passes.h" #include "mlir/Pass/PassManager.h" #include "mlir/Transforms/GreedyPatternRewriteDriver.h" diff --git a/flang/include/flang/Optimizer/Support/InitFIR.h b/flang/include/flang/Optimizer/Support/InitFIR.h index 67e9287ddad4f..184e2aa141d85 100644 --- a/flang/include/flang/Optimizer/Support/InitFIR.h +++ b/flang/include/flang/Optimizer/Support/InitFIR.h @@ -32,6 +32,8 @@ #include "mlir/Dialect/LLVMIR/NVVMDialect.h" #include "mlir/Dialect/LLVMIR/Transforms/InlinerInterfaceImpl.h" #include "mlir/Dialect/Math/IR/Math.h" +#include "mlir/Dialect/MemRef/IR/MemRef.h" +#include "mlir/Dialect/MemRef/Transforms/Passes.h" #include "mlir/Dialect/OpenACC/OpenACC.h" #include "mlir/Dialect/OpenACC/Transforms/Passes.h" #include "mlir/Dialect/SCF/IR/SCF.h" @@ -54,7 +56,8 @@ namespace fir::support { mlir::NVVM::NVVMDialect, mlir::gpu::GPUDialect, \ mlir::index::IndexDialect, mif::MIFDialect -#define FLANG_CODEGEN_DIALECT_LIST FIRCodeGenDialect, mlir::LLVM::LLVMDialect +#define FLANG_CODEGEN_DIALECT_LIST \ + FIRCodeGenDialect, mlir::memref::MemRefDialect, mlir::LLVM::LLVMDialect // The definitive list of dialects used by flang. #define FLANG_DIALECT_LIST \ @@ -129,7 +132,11 @@ inline void registerMLIRPassesForFortranTools() { mlir::affine::registerAffineLoopTilingPass(); mlir::affine::registerAffineDataCopyGenerationPass(); + mlir::registerMem2RegPass(); + mlir::memref::registerMemRefPasses(); + mlir::registerLowerAffinePass(); + mlir::registerReconcileUnrealizedCastsPass(); } /// Register the interfaces needed to lower to LLVM IR. diff --git a/flang/include/flang/Optimizer/Transforms/Passes.td b/flang/include/flang/Optimizer/Transforms/Passes.td index bb2509b1747d5..0bf1537b2215c 100644 --- a/flang/include/flang/Optimizer/Transforms/Passes.td +++ b/flang/include/flang/Optimizer/Transforms/Passes.td @@ -87,6 +87,17 @@ def FIRToSCFPass : Pass<"fir-to-scf"> { ]; } +def ConvertFIRToMLIRPass : Pass<"fir-to-mlir", "mlir::ModuleOp"> { + let summary = "Convert the FIR dialect module to MLIR standard dialects."; + let description = [{ + Convert the FIR dialect module to MLIR standard dialects. + }]; + let dependentDialects = [ + "fir::FIROpsDialect", "fir::FIRCodeGenDialect", "mlir::scf::SCFDialect", + "mlir::memref::MemRefDialect", "mlir::affine::AffineDialect" + ]; +} + def AnnotateConstantOperands : Pass<"annotate-constant"> { let summary = "Annotate constant operands to all FIR operations"; let description = [{ diff --git a/flang/include/flang/Tools/CrossToolHelpers.h b/flang/include/flang/Tools/CrossToolHelpers.h index e964882ef6dac..0dcb99e1eb5b1 100644 --- a/flang/include/flang/Tools/CrossToolHelpers.h +++ b/flang/include/flang/Tools/CrossToolHelpers.h @@ -137,6 +137,7 @@ struct MLIRToLLVMPassPipelineConfig : public FlangEPCallBacks { bool EnableOpenMP = false; ///< Enable OpenMP lowering. bool EnableOpenMPSimd = false; ///< Enable OpenMP simd-only mode. bool SkipConvertComplexPow = false; ///< Do not run complex pow conversion. + bool LowerThroughMLIR = false; ///< Lower form FIR through MLIR to LLVM std::string InstrumentFunctionEntry = ""; ///< Name of the instrument-function that is called on each ///< function-entry diff --git a/flang/lib/Frontend/CompilerInvocation.cpp b/flang/lib/Frontend/CompilerInvocation.cpp index 893121fe01f27..8c3fde0a27153 100644 --- a/flang/lib/Frontend/CompilerInvocation.cpp +++ b/flang/lib/Frontend/CompilerInvocation.cpp @@ -1580,6 +1580,10 @@ bool CompilerInvocation::createFromArgs( invoc.loweringOpts.setLowerToHighLevelFIR(false); } + // -flang-experimental-lower-through-mlir + invoc.loweringOpts.setLowerThroughMLIR( + args.hasArg(clang::options::OPT_flang_experimental_lower_through_mlir)); + // -fno-ppc-native-vector-element-order if (args.hasArg(clang::options::OPT_fno_ppc_native_vec_elem_order)) { invoc.loweringOpts.setNoPPCNativeVecElemOrder(true); diff --git a/flang/lib/Frontend/FrontendActions.cpp b/flang/lib/Frontend/FrontendActions.cpp index 159d08a2797b3..0cb241f209522 100644 --- a/flang/lib/Frontend/FrontendActions.cpp +++ b/flang/lib/Frontend/FrontendActions.cpp @@ -769,6 +769,7 @@ void CodeGenAction::generateLLVMIR() { config.NSWOnLoopVarInc = false; config.ComplexRange = opts.getComplexRange(); + config.LowerThroughMLIR = invoc.getLoweringOpts().getLowerThroughMLIR(); // Create the pass pipeline fir::createMLIRToLLVMPassPipeline(pm, config, getCurrentFile()); diff --git a/flang/lib/Optimizer/CodeGen/CodeGen.cpp b/flang/lib/Optimizer/CodeGen/CodeGen.cpp index f96d45d3f6b66..1f52cd4011b7a 100644 --- a/flang/lib/Optimizer/CodeGen/CodeGen.cpp +++ b/flang/lib/Optimizer/CodeGen/CodeGen.cpp @@ -1008,6 +1008,29 @@ struct ConvertOpConversion : public fir::FIROpConversion<fir::ConvertOp> { rewriter.replaceOpWithNewOp<mlir::LLVM::BitcastOp>(convert, toTy, op0); return mlir::success(); } + // Pointer to MemRef conversion. + if (mlir::isa<mlir::MemRefType>(toFirTy)) { + auto dstMemRef = mlir::MemRefDescriptor::poison(rewriter, loc, toTy); + dstMemRef.setAlignedPtr(rewriter, loc, op0); + dstMemRef.setOffset( + rewriter, loc, + createIndexAttrConstant(rewriter, loc, getIndexType(), 0)); + rewriter.replaceOp(convert, {dstMemRef}); + return mlir::success(); + } + } else if (mlir::isa<mlir::MemRefType>(fromFirTy) && + mlir::isa<mlir::LLVM::LLVMPointerType>(toTy)) { + // MemRef to pointer conversion. + auto srcMemRef = mlir::MemRefDescriptor(op0); + mlir::Type elementType = typeConverter->convertType( + mlir::cast<mlir::MemRefType>(fromFirTy).getElementType()); + mlir::Value srcBasePtr = srcMemRef.alignedPtr(rewriter, loc); + mlir::Value srcOffset = srcMemRef.offset(rewriter, loc); + mlir::Value srcPtr = + mlir::LLVM::GEPOp::create(rewriter, loc, srcBasePtr.getType(), + elementType, srcBasePtr, srcOffset); + rewriter.replaceOp(convert, srcPtr); + return mlir::success(); } return emitError(loc) << "cannot convert " << fromTy << " to " << toTy; } @@ -3521,6 +3544,13 @@ struct SelectCaseOpConversion : public fir::FIROpConversion<fir::SelectCaseOp> { auto loc = caseOp.getLoc(); for (unsigned t = 0; t != conds; ++t) { mlir::Block *dest = caseOp.getSuccessor(t); + std::optional<mlir::TypeConverter::SignatureConversion> conversion = + getTypeConverter()->convertBlockSignature(dest); + if (!conversion) { + return rewriter.notifyMatchFailure(caseOp, + "could not compute block signature"); + } + dest = rewriter.applySignatureConversion(dest, *conversion); std::optional<mlir::ValueRange> destOps = caseOp.getSuccessorOperands(adaptor.getOperands(), t); std::optional<mlir::ValueRange> cmpOps = @@ -4326,7 +4356,7 @@ class FIRToLLVMLowering target.addLegalDialect<mlir::gpu::GPUDialect>(); // required NOPs for applying a full conversion - target.addLegalOp<mlir::ModuleOp>(); + target.addLegalOp<mlir::ModuleOp, mlir::UnrealizedConversionCastOp>(); // If we're on Windows, we might need to rename some libm calls. bool isMSVC = fir::getTargetTriple(mod).isOSMSVCRT(); @@ -4346,9 +4376,12 @@ class FIRToLLVMLowering }); } + mlir::ConversionConfig config; + config.buildMaterializations = false; + // apply the patterns if (mlir::failed(mlir::applyFullConversion(getModule(), target, - std::move(pattern)))) { + std::move(pattern), config))) { signalPassFailure(); } diff --git a/flang/lib/Optimizer/Passes/Pipelines.cpp b/flang/lib/Optimizer/Passes/Pipelines.cpp index 103e736accca0..3e233e07228ed 100644 --- a/flang/lib/Optimizer/Passes/Pipelines.cpp +++ b/flang/lib/Optimizer/Passes/Pipelines.cpp @@ -109,6 +109,21 @@ void addDebugInfoPass(mlir::PassManager &pm, void addFIRToLLVMPass(mlir::PassManager &pm, const MLIRToLLVMPassPipelineConfig &config) { + if (disableFirToLlvmIr) + return; + + if (config.LowerThroughMLIR) { + pm.addPass(createConvertFIRToMLIRPass()); + pm.addPass(mlir::memref::createFoldMemRefAliasOpsPass()); + pm.addPass(mlir::createMem2Reg()); + pm.addPass(mlir::createCSEPass()); + pm.addPass(mlir::createCanonicalizerPass()); + pm.addPass(mlir::memref::createExpandOpsPass()); + pm.addPass(mlir::memref::createExpandStridedMetadataPass()); + pm.addPass(mlir::createLowerAffinePass()); + pm.addPass(mlir::createFinalizeMemRefToLLVMConversionPass()); + } + fir::FIRToLLVMPassOptions options; options.ignoreMissingTypeDescriptors = ignoreMissingTypeDescriptors; options.skipExternalRttiDefinition = skipExternalRttiDefinition; @@ -117,13 +132,11 @@ void addFIRToLLVMPass(mlir::PassManager &pm, options.typeDescriptorsRenamedForAssembly = !disableCompilerGeneratedNamesConversion; options.ComplexRange = config.ComplexRange; - addPassConditionally(pm, disableFirToLlvmIr, - [&]() { return fir::createFIRToLLVMPass(options); }); + pm.addPass(fir::createFIRToLLVMPass(options)); + // The dialect conversion framework may leave dead unrealized_conversion_cast // ops behind, so run reconcile-unrealized-casts to clean them up. - addPassConditionally(pm, disableFirToLlvmIr, [&]() { - return mlir::createReconcileUnrealizedCastsPass(); - }); + pm.addPass(mlir::createReconcileUnrealizedCastsPass()); } void addLLVMDialectToLLVMPass(mlir::PassManager &pm, diff --git a/flang/lib/Optimizer/Transforms/CMakeLists.txt b/flang/lib/Optimizer/Transforms/CMakeLists.txt index 0388439f89a54..a6423b3dea5a9 100644 --- a/flang/lib/Optimizer/Transforms/CMakeLists.txt +++ b/flang/lib/Optimizer/Transforms/CMakeLists.txt @@ -36,6 +36,7 @@ add_flang_library(FIRTransforms SimplifyFIROperations.cpp OptimizeArrayRepacking.cpp ConvertComplexPow.cpp + ConvertFIRToMLIR.cpp MIFOpConversion.cpp DEPENDS diff --git a/flang/lib/Optimizer/Transforms/ConvertFIRToMLIR.cpp b/flang/lib/Optimizer/Transforms/ConvertFIRToMLIR.cpp new file mode 100644 index 0000000000000..1470f78dc4e9c --- /dev/null +++ b/flang/lib/Optimizer/Transforms/ConvertFIRToMLIR.cpp @@ -0,0 +1,238 @@ +//===-- ConvertFIRToMLIR.cpp ----------------------------------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "flang/Optimizer/Dialect/FIRCG/CGOps.h" +#include "flang/Optimizer/Dialect/FIRDialect.h" +#include "flang/Optimizer/Transforms/Passes.h" +#include "mlir/Dialect/Affine/IR/AffineOps.h" +#include "mlir/Dialect/MemRef/IR/MemRef.h" +#include "mlir/Dialect/SCF/IR/SCF.h" +#include "mlir/Transforms/DialectConversion.h" + +namespace fir { +#define GEN_PASS_DEF_CONVERTFIRTOMLIRPASS +#include "flang/Optimizer/Transforms/Passes.h.inc" +} // namespace fir + +namespace { +class ConvertFIRToMLIRPass + : public fir::impl::ConvertFIRToMLIRPassBase<ConvertFIRToMLIRPass> { +public: + void runOnOperation() override; +}; + +class FIRLoadOpLowering : public mlir::OpConversionPattern<fir::LoadOp> { +public: + using mlir::OpConversionPattern<fir::LoadOp>::OpConversionPattern; + + mlir::LogicalResult + matchAndRewrite(fir::LoadOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const override { + if (!getTypeConverter()->convertType(op.getMemref())) + return mlir::failure(); + + rewriter.replaceOpWithNewOp<mlir::memref::LoadOp>(op, adaptor.getMemref(), + mlir::ValueRange{}); + return mlir::success(); + } +}; + +class FIRStoreOpLowering : public mlir::OpConversionPattern<fir::StoreOp> { +public: + using mlir::OpConversionPattern<fir::StoreOp>::OpConversionPattern; + + mlir::LogicalResult + matchAndRewrite(fir::StoreOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const override { + if (!getTypeConverter()->convertType(op.getMemref())) + return mlir::failure(); + + rewriter.replaceOpWithNewOp<mlir::memref::StoreOp>( + op, adaptor.getValue(), adaptor.getMemref(), mlir::ValueRange{}); + return mlir::success(); + } +}; + +class FIRConvertOpLowering : public mlir::OpConversionPattern<fir::ConvertOp> { +public: + using mlir::OpConversionPattern<fir::ConvertOp>::OpConversionPattern; + + mlir::LogicalResult + matchAndRewrite(fir::ConvertOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const override { + auto srcType = getTypeConverter()->convertType(op.getValue().getType()); + auto dstType = getTypeConverter()->convertType(op.getType()); + + if (!srcType || !dstType) + return mlir::failure(); + + auto srcVal = adaptor.getValue(); + + if (srcType == dstType) { + rewriter.replaceOp(op, mlir::ValueRange{srcVal}); + } else if (srcType.isIntOrIndex() && dstType.isIntOrIndex()) { + if (srcType.isIndex() || dstType.isIndex()) { + rewriter.replaceOpWithNewOp<mlir::arith::IndexCastOp>(op, dstType, + srcVal); + } else if (srcType.getIntOrFloatBitWidth() < + dstType.getIntOrFloatBitWidth()) { + rewriter.replaceOpWithNewOp<mlir::arith::ExtSIOp>(op, dstType, srcVal); + } else { + rewriter.replaceOpWithNewOp<mlir::arith::TruncIOp>(op, dstType, srcVal); + } + } else if (srcType.isFloat() && dstType.isFloat()) { + if (srcType.getIntOrFloatBitWidth() < dstType.getIntOrFloatBitWidth()) { + rewriter.replaceOpWithNewOp<mlir::arith::ExtFOp>(op, dstType, srcVal); + } else { + rewriter.replaceOpWithNewOp<mlir::arith::TruncFOp>(op, dstType, srcVal); + } + } else { + return mlir::failure(); + } + + return mlir::success(); + } +}; + +class FIRAllocOpLowering : public mlir::OpConversionPattern<fir::AllocaOp> { +public: + using mlir::OpConversionPattern<fir::AllocaOp>::OpConversionPattern; + + mlir::LogicalResult + matchAndRewrite(fir::AllocaOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const override { + if (!mlir::MemRefType::isValidElementType(op.getAllocatedType()) || + op.hasLenParams()) + return mlir::failure(); + + auto dstType = getTypeConverter()->convertType(op.getType()); + auto allocaOp = mlir::memref::AllocaOp::create( + rewriter, op.getLoc(), + mlir::MemRefType::get({}, op.getAllocatedType())); + allocaOp->setAttrs(op->getAttrs()); + rewriter.replaceOpWithNewOp<mlir::memref::CastOp>(op, dstType, allocaOp); + return mlir::success(); + } +}; + +class FIRXArrayCoorOpLowering + : public mlir::OpConversionPattern<fir::cg::XArrayCoorOp> { +public: + using mlir::OpConversionPattern<fir::cg::XArrayCoorOp>::OpConversionPattern; + + mlir::LogicalResult + matchAndRewrite(fir::cg::XArrayCoorOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const override { + if (!getTypeConverter()->convertType(op.getMemref())) + return mlir::failure(); + + mlir::Location loc = op.getLoc(); + auto metadata = mlir::memref::ExtractStridedMetadataOp::create( + rewriter, loc, adaptor.getMemref()); + auto base = metadata.getBaseBuffer(); + auto offset = metadata.getOffset(); + mlir::ValueRange shape = adaptor.getShape(); + unsigned rank = op.getRank(); + + assert(rank > 0 && "expected rank to be greater than zero"); + + auto sizes = llvm::to_vector_of<mlir::OpFoldResult>(llvm::reverse(shape)); + mlir::SmallVector<mlir::OpFoldResult> strides(rank); + + strides[rank - 1] = rewriter.getIndexAttr(1); + mlir::Value stride = mlir::arith::ConstantIndexOp::create(rewriter, loc, 1); + for (unsigned i = 1; i < rank; ++i) { + stride = mlir::arith::MulIOp::create(rewriter, loc, stride, shape[i - 1]); + strides[rank - 1 - i] = stride; + } + + mlir::Value memref = mlir::memref::ReinterpretCastOp::create( + rewriter, loc, base, offset, sizes, strides); + + mlir::SmallVector<mlir::OpFoldResult> oneAttrs(rank, + rewriter.getIndexAttr(1)); + auto one = mlir::arith::ConstantIndexOp::create(rewriter, loc, 1); + auto offsets = llvm::map_to_vector( + llvm::reverse(adaptor.getIndices()), + [&](mlir::Value idx) -> mlir::OpFoldResult { + if (idx.getType().isInteger()) + idx = mlir::arith::IndexCastOp::create( + rewriter, loc, rewriter.getIndexType(), idx); + + assert(idx.getType().isIndex() && "expected index type"); + idx = mlir::arith::SubIOp... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/168703 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
