Code Monkey home page Code Monkey logo

buddy-mlir's Introduction

BUDDY MLIR

An MLIR-based compiler framework designed for a co-design ecosystem from DSL (domain-specific languages) to DSA (domain-specific architectures). (Project page).

Getting Started

The default build system uses LLVM/MLIR as an external library. We also provide a one-step build strategy for users who only want to use our tools. Please make sure the dependencies are available on your machine.

LLVM/MLIR Dependencies

Before building, please make sure the dependencies are available on your machine.

Clone and Initialize

$ git clone [email protected]:buddy-compiler/buddy-mlir.git
$ cd buddy-mlir
$ git submodule update --init

Build and Test LLVM/MLIR/CLANG

$ cd buddy-mlir
$ mkdir llvm/build
$ cd llvm/build
$ cmake -G Ninja ../llvm \
    -DLLVM_ENABLE_PROJECTS="mlir;clang" \
    -DLLVM_TARGETS_TO_BUILD="host;RISCV" \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DCMAKE_BUILD_TYPE=RELEASE
$ ninja check-mlir check-clang

If your target machine includes a Nvidia GPU, you can use the following configuration:

$ cmake -G Ninja ../llvm \
    -DLLVM_ENABLE_PROJECTS="mlir;clang" \
    -DLLVM_TARGETS_TO_BUILD="host;RISCV;NVPTX" \
    -DMLIR_ENABLE_CUDA_RUNNER=ON \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DCMAKE_BUILD_TYPE=RELEASE

To enable MLIR Python bindings, please use the following configuration:

$ cmake -G Ninja ../llvm \
    -DLLVM_ENABLE_PROJECTS="mlir;clang" \
    -DLLVM_TARGETS_TO_BUILD="host;RISCV" \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DCMAKE_BUILD_TYPE=RELEASE \
    -DMLIR_ENABLE_BINDINGS_PYTHON=ON \
    -DPython3_EXECUTABLE=$(which python3)

If your target machine has lld installed, you can use the following configuration:

$ cmake -G Ninja ../llvm \
    -DLLVM_ENABLE_PROJECTS="mlir;clang" \
    -DLLVM_TARGETS_TO_BUILD="host;RISCV" \
    -DLLVM_USE_LINKER=lld \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DCMAKE_BUILD_TYPE=RELEASE

Build buddy-mlir

If you have previously built the llvm-project, you can replace the $PWD with the path to the directory where you have successfully built the llvm-project.

$ cd buddy-mlir
$ mkdir build
$ cd build
$ cmake -G Ninja .. \
    -DMLIR_DIR=$PWD/../llvm/build/lib/cmake/mlir \
    -DLLVM_DIR=$PWD/../llvm/build/lib/cmake/llvm \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DCMAKE_BUILD_TYPE=RELEASE
$ ninja
$ ninja check-buddy

To utilize the Buddy Compiler Python package, please ensure that the MLIR Python bindings are enabled and use the following configuration:

$ cmake -G Ninja .. \
    -DMLIR_DIR=$PWD/../llvm/build/lib/cmake/mlir \
    -DLLVM_DIR=$PWD/../llvm/build/lib/cmake/llvm \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DCMAKE_BUILD_TYPE=RELEASE \
    -DBUDDY_MLIR_ENABLE_PYTHON_PACKAGES=ON \
    -DPython3_EXECUTABLE=$(which python3)

If you want to add domain-specific framework support, please add the following cmake options:

Framework Enable Option Other Options
OpenCV -DBUDDY_ENABLE_OPENCV=ON Add -DOpenCV_DIR=</PATH/TO/OPENCV/BUILD/> or install OpenCV release version on your local device.

One-step building strategy

If you only want to use our tools and integrate them more easily into your projects, you can choose to use the one-step build strategy.

$ cmake -G Ninja -Bbuild \
    -DCMAKE_BUILD_TYPE=Release \
    -DLLVM_ENABLE_PROJECTS="mlir;clang" \
    -DLLVM_TARGETS_TO_BUILD="host;RISCV" \
    -DLLVM_EXTERNAL_PROJECTS="buddy-mlir" \
    -DLLVM_EXTERNAL_BUDDY_MLIR_SOURCE_DIR="$PWD" \
    -DLLVM_ENABLE_ASSERTIONS=ON \
    -DCMAKE_BUILD_TYPE=RELEASE \
    llvm/llvm
$ cd build
$ ninja check-mlir check-clang
$ ninja
$ ninja check-buddy

Use nix

This repository have nix flake support. You can follow the nix installation instruction and enable the flake features to have nix setup.

  • If you want to contribute to this project:
nix develop .

This will setup a bash shell with clang, clangd, cmake, ninja, and other necessary dependencies to build buddy-mlir from source.

  • If you want to use the buddy-mlir bintools
nix build .#buddy-mlir
./result/bin/buddy-opt --version

Dialects

Bud Dialect

Bud dialect is designed for testing and demonstrating.

DIP Dialect

DIP dialect is designed for digital image processing abstraction.

Tools

buddy-opt

The buddy-opt is the driver for dialects and optimization in buddy-mlir project.

buddy-lsp-server

This program should be a drop-in replacement for mlir-lsp-server, supporting new dialects defined in buddy-mlir. To use it, please directly modify mlir LSP server path in VSCode settings (or similar settings for other editors) to:

{
    "mlir.server_path": "YOUR_BUDDY_MLIR_BUILD/bin/buddy-lsp-server",
}

After modification, your editor should have correct completion and error prompts for new dialects such as rvv and gemmini.

Examples

The purpose of the examples is to give users a better understanding of how to use the passes and the interfaces in buddy-mlir. Currently, we provide three types of examples.

  • IR level conversion and transformation examples.
  • Domain-specific application level examples.
  • Testing and demonstrating examples.

For more details, please see the documentation of the examples.

How to Cite

If you find our project and research useful or refer to it in your own work, please cite our paper as follows:

@article{zhang2023compiler,
  title={Compiler Technologies in Deep Learning Co-Design: A Survey},
  author={Zhang, Hongbin and Xing, Mingjie and Wu, Yanjun and Zhao, Chen},
  journal={Intelligent Computing},
  year={2023},
  publisher={AAAS}
}

For direct access to the paper, please visit Compiler Technologies in Deep Learning Co-Design: A Survey.

buddy-mlir's People

Contributors

amanchhaparia avatar artemskrebkov avatar avimitin avatar axmat avatar bbuf avatar ellislambda avatar flagerlee avatar guan-schoolmate avatar guessmewho123888 avatar inclyc avatar joejiong avatar leiwang1999 avatar lester-1 avatar lhy-24 avatar linuxlonelyeagle avatar meshtag avatar origami404 avatar ris-bali avatar sforekeeper avatar taiqzheng avatar tsworld1314 avatar weilinquan avatar wlfj avatar wuxintong123 avatar xinyu302 avatar xlinsist avatar xtayex avatar zdx0317 avatar zhanghb97 avatar zrr1999 avatar

Stargazers

 avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar

Watchers

 avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar  avatar

buddy-mlir's Issues

Get rid of warnings during the basic build step

Is your feature request related to a problem? Please describe.
Following is the output log for ninja check-buddy

In file included from ../include/Interface/buddy/dap/AudioContainer.h:24,
                 from ../tests/Interface/core/AudioContainerTest.cpp:23:
../thirdparty/include/AudioFile.h: In instantiation of 'bool AudioFile<T>::decodeWaveFile(std::vector<unsigned char>&) [with T = float]':
../thirdparty/include/AudioFile.h:419:12:   required from 'bool AudioFile<T>::loadFromMemory(std::vector<unsigned char>&) [with T = float]'
../thirdparty/include/AudioFile.h:409:10:   required from 'bool AudioFile<T>::load(std::string) [with T = float; std::string = std::__cxx11::basic_string<char>]'
../thirdparty/include/AudioFile.h:289:3:   required from 'AudioFile<T>::AudioFile(std::string) [with T = float; std::string = std::__cxx11::basic_string<char>]'
../include/Interface/buddy/dap/AudioContainer.h:37:75:   required from 'dap::Audio<T, N>::Audio(std::string) [with T = float; long unsigned int N = 1; std::string = std::__cxx11::basic_string<char>]'
../tests/Interface/core/AudioContainerTest.cpp:29:76:   required from here
../thirdparty/include/AudioFile.h:545:50: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
  545 |           sample = (T) reinterpret_cast<float &>(sampleAsInt);
      |                                                  ^~~~~~~~~~~
../thirdparty/include/AudioFile.h: In instantiation of 'bool AudioFile<T>::decodeAiffFile(std::vector<unsigned char>&) [with T = float]':
../thirdparty/include/AudioFile.h:421:12:   required from 'bool AudioFile<T>::loadFromMemory(std::vector<unsigned char>&) [with T = float]'
../thirdparty/include/AudioFile.h:409:10:   required from 'bool AudioFile<T>::load(std::string) [with T = float; std::string = std::__cxx11::basic_string<char>]'
../thirdparty/include/AudioFile.h:289:3:   required from 'AudioFile<T>::AudioFile(std::string) [with T = float; std::string = std::__cxx11::basic_string<char>]'
../include/Interface/buddy/dap/AudioContainer.h:37:75:   required from 'dap::Audio<T, N>::Audio(std::string) [with T = float; long unsigned int N = 1; std::string = std::__cxx11::basic_string<char>]'
../tests/Interface/core/AudioContainerTest.cpp:29:76:   required from here
../thirdparty/include/AudioFile.h:697:50: warning: dereferencing type-punned pointer will break strict-aliasing rules [-Wstrict-aliasing]
  697 |           sample = (T) reinterpret_cast<float &>(sampleAsInt);
      |                                                  ^~~~~~~~~~~
[60/98] Building CXX object lib/Dialect/Bud/CMakeFiles/obj.BuddyBud.dir/BudDialect.cpp.o
In file included from ../lib/Dialect/Bud/BudDialect.cpp:84:
include/Dialect/Bud/BudOpsAttributes.cpp.inc:31:30: warning: 'mlir::LogicalResult generatedAttributePrinter(mlir::Attribute, mlir::AsmPrinter&)' defined but not used [-Wunused-function]
   31 | static ::mlir::LogicalResult generatedAttributePrinter(::mlir::Attribute def, ::mlir::AsmPrinter &printer) {
      |                              ^~~~~~~~~~~~~~~~~~~~~~~~~
In file included from ../lib/Dialect/Bud/BudDialect.cpp:84:
include/Dialect/Bud/BudOpsAttributes.cpp.inc:19:36: warning: 'mlir::OptionalParseResult generatedAttributeParser(mlir::AsmParser&, llvm::StringRef*, mlir::Type, mlir::Attribute&)' defined but not used [-Wunused-function]
   19 | static ::mlir::OptionalParseResult generatedAttributeParser(::mlir::AsmParser &parser, ::llvm::StringRef *mnemonic, ::mlir::Type type, ::mlir::Attribute &value) {
      |                                    ^~~~~~~~~~~~~~~~~~~~~~~~
[87/98] Building CXX object lib/Conversion/LowerDAP/CMakeFiles/obj.LowerDAPPass.dir/LowerDAPPass.cpp.o
../lib/Conversion/LowerDAP/LowerDAPPass.cpp: In member function 'virtual mlir::LogicalResult {anonymous}::DAPBiquadLowering::matchAndRewrite(buddy::dap::BiquadOp, mlir::PatternRewriter&) const':
../lib/Conversion/LowerDAP/LowerDAPPass.cpp:94:11: warning: variable 'a0' set but not used [-Wunused-but-set-variable]
   94 |     Value a0 = rewriter.create<memref::LoadOp>(loc, kernel, ValueRange{c3});
      |           ^~
../lib/Conversion/LowerDAP/LowerDAPPass.cpp: In lambda function:
../lib/Conversion/LowerDAP/LowerDAPPass.cpp:220:17: warning: variable 'a0' set but not used [-Wunused-but-set-variable]
  220 |           Value a0 = builder.create<memref::LoadOp>(loc, kernel,
      |                 ^~

Describe the solution you'd like
It would be good if these warnings are eliminated without affecting compile time and run-time performance.

[DIP] rotate2d operation may change picture size.

Describe the bug
Image size change after rotation.

  • input shape [1026, 1026]
  • output shape [1028, 1028]

To Reproduce
After compiling the image-processing-rotate-benchmark benchmark. (rotate2d PR)
RUN:
./image-processing-rotate-benchmark ../../benchmarks/ImageProcessing/Images/YuTu.png RADIAN 1.57

Expected behavior
An output image with [1026, 1026].

Additional context
For comparing the similarity of Buddy & OpenCV rotation image, we need their result image with the same size.

[Testing] Add tests for assert statements

Is your feature request related to a problem? Please describe.
There are no tests to check assert statements in the project.

Describe the solution you'd like
We should develop support for checking assert statements. This should ensure that our implementation(s) behave in an expected manner all the time as well as increase completeness of the testing mechanism.

[DIP] Create methods for splitting and merging channels in Image Container

Is your feature request related to a problem? Please describe.
As of now, we cannot work with individual channels directly in a multi-channel image stored in the image container.

Describe the solution you'd like
We should have separate methods for splitting and merging channels from/in the image container. Similar to cv2.split() and cv2.merge() for OpenCV's Mat container.

Fix dip.corr_2d for stride=13 and other similar cases

Is your feature request related to a problem? Please describe.
Our implementation for lowering pass of dip.corr_2d is based on the consideration that we would never require tail processing for col_mid
(Ex.

// colMid & rowMid
) elements.
However, in certain cases of stride values, this region requires tail processing as well (Ex. stride = 13).

Describe the solution you'd like
I have already solved this issue in this branch of my fork, but we might be able to fix this in a more elegant way (perhaps without adding extra machine instructions in each iteration).

Additional context
This feature request is created for future reference (as I don't have much time to deal with this problem now), feel free to start a discussion and share your ideas.

Add type inference mechanism in dip.corr_2d's lowering pass

Is your feature request related to a problem? Please describe.
As of now, pixel elements are stored in a vector having floating point 32 bit elements while lowering dip.corr_2d (Ref :

VectorType vectorTy32 = VectorType::get({stride}, f32);
).

Describe the solution you'd like
We should infer the type from user provided memref/image and use it for further calculations.

[Compiler]Define a generic set of pipelines

Is your feature request related to a problem? Please describe.
The reason I don't think the current pipeline is elegant is because all the pipelines are separated. For example, for a model, we use mlir upstream and buddy passes in order to lower that model, and many times different pipelines are used for different models, and I don't think this is elegant enough. I envision that we can design our own pipeline, in which case we only need to use this one pipeline when lowering code and models.

Describe the solution you'd like
I don't have too many good ideas at the moment, I don't know what you think @zhanghb97 .

build-rvv-env.sh script issues

The build-rvv-env.sh script In buddy-mlir/thirdparty directory has some problems.
First of all, in riscv-gnu-toolchain repo, there is no more rvv-next branch, which should be deleted.
Secondly, when we pull from large git repos, it alaways happens pull failure and takes a lot of time in pull process. For user-friendly reasons, I think it's better to run git clone or git submodule update commands in "while true" block. Like below:

#!/bin/bash
while true; do
    git submodule update --init --recursive
    if [ $? -eq 0 ]; then
        echo "Submodule update completed successfully."
        break
    else
        echo "Submodule update failed. Retrying in 5 seconds..."
        sleep 5
    fi
done

Thirdly, in "Build cross MLIR" and "Build cross Buddy-MLIR" process, leaving out the cd .. command.

Inquiry about good first issues

Greetings!

I have lots of interest in this project, especially the following two parts:

  • Convert DSL to MLIR.
  • Improve performance of a specified operation or a model.

I'm trying to learn and participate in it but I have found there is few good first issues. Could you please give me some guides or some new good first issues to start up? Many thanks!

Best regards,
Rinne

how to contact you in slack

I have downloaded slack, but i found that it's hard to join your group, especially i need a workspace url ๏ผŒIs there any guide for us to join the slack workspace to contact you for some more details?
Or can we contact in wechat group talk?

Design spec of buddy-mlir

Is your feature request related to a problem? Please describe.
A clear and concise description of what the problem is. Ex. I'm always frustrated when [...]

Thank you for the wonderful project. Can you provide a specification of the design specification of Buddy-MLIR? I really want to learn how to add feature into MLIR.

Describe the solution you'd like
A clear and concise description of what you want to happen.

Describe alternatives you've considered
A clear and concise description of any alternative solutions or features you've considered.

Additional context
Add any other context about the feature request here.

[Bug]MLIRSparseTensor example can't run

Describe the bug
~/buddy-mlir/examples/MLIRSparseTensor/sparse-tensor-binary.mlir can not run.The same problem will occur when running other files in this folder.

To Reproduce
Just run this command: ~/buddy-mlir/examples/MLIRSparseTensor$ make sparse-tensor-binary-run

Bug Info

./sparse-tensor-binary.mlir:1:22: error: unexpected key: dimLevelType
#SV = #sparse_tensor.encoding<{
                     ^
Error: entry point not found
makefile:96: recipe for target 'sparse-tensor-binary-run' failed
make: *** [sparse-tensor-binary-run] Error 1

Desktop (please complete the following information):

  • OS: Ubuntu
  • Version: 18.04
  • buddy-mlir version : faec06b
  • llvm version : 6c59f0e

Solution

#SV = #sparse_tensor.encoding<{
  dimLevelType = [ "compressed" ]
}>

should be replaced with

#SV = #sparse_tensor.encoding<{
  map = (i) -> (i : compressed)
}>

then run ~/buddy-mlir/examples/MLIRSparseTensor$ make sparse-tensor-binary-run, and you will get correct output like this:
image

[Build Failed] Edge Detection case doesn't work on macos.

Command:

/Applications/CLion.app/Contents/bin/cmake/mac/bin/cmake --build /Users/wanglei/buddy-mlir/build --target edge-detection

Log:

[0/1] Re-running CMake...
-- Building with -fPIC
-- 	SSE support - yes
-- 	AVX2 support - yes
-- 	AVX512 support - no
-- Spliting size: 256
/Users/wanglei/buddy-mlir/examples
-- boost include path is : /usr/local/include
-- Configuring done
-- Generating done
-- Build files have been written to: /Users/wanglei/buddy-mlir/build
[2/3] Linking C static library examples/ConvOpt/libConv2D.a
warning: /Library/Developer/CommandLineTools/usr/bin/ranlib: archive library: examples/ConvOpt/libConv2D.a the table of contents is empty (no object file members in the library define global symbols)
[3/3] Linking CXX executable bin/edge-detection
FAILED: bin/edge-detection 
: && /Library/Developer/CommandLineTools/usr/bin/c++ -fPIC -fvisibility-inlines-hidden -Werror=date-time -Werror=unguarded-availability-new -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wmissing-field-initializers -Wimplicit-fallthrough -Wcovered-switch-default -Wno-noexcept-type -Wnon-virtual-dtor -Wdelete-non-virtual-dtor -Wstring-conversion -Wmisleading-indentation -fdiagnostics-color -g -isysroot /Library/Developer/CommandLineTools/SDKs/MacOSX11.1.sdk -mmacosx-version-min=10.15 -Wl,-search_paths_first -Wl,-headerpad_max_install_names  examples/ConvOpt/CMakeFiles/edge-detection.dir/edge-detection.cpp.o -o bin/edge-detection  -Wl,-rpath,/usr/local/lib  /usr/local/lib/libopencv_gapi.4.5.4.dylib  /usr/local/lib/libopencv_stitching.4.5.4.dylib  /usr/local/lib/libopencv_alphamat.4.5.4.dylib  /usr/local/lib/libopencv_aruco.4.5.4.dylib  /usr/local/lib/libopencv_barcode.4.5.4.dylib  /usr/local/lib/libopencv_bgsegm.4.5.4.dylib  /usr/local/lib/libopencv_bioinspired.4.5.4.dylib  /usr/local/lib/libopencv_ccalib.4.5.4.dylib  /usr/local/lib/libopencv_dnn_objdetect.4.5.4.dylib  /usr/local/lib/libopencv_dnn_superres.4.5.4.dylib  /usr/local/lib/libopencv_dpm.4.5.4.dylib  /usr/local/lib/libopencv_face.4.5.4.dylib  /usr/local/lib/libopencv_freetype.4.5.4.dylib  /usr/local/lib/libopencv_fuzzy.4.5.4.dylib  /usr/local/lib/libopencv_hfs.4.5.4.dylib  /usr/local/lib/libopencv_img_hash.4.5.4.dylib  /usr/local/lib/libopencv_intensity_transform.4.5.4.dylib  /usr/local/lib/libopencv_line_descriptor.4.5.4.dylib  /usr/local/lib/libopencv_mcc.4.5.4.dylib  /usr/local/lib/libopencv_quality.4.5.4.dylib  /usr/local/lib/libopencv_rapid.4.5.4.dylib  /usr/local/lib/libopencv_reg.4.5.4.dylib  /usr/local/lib/libopencv_rgbd.4.5.4.dylib  /usr/local/lib/libopencv_saliency.4.5.4.dylib  /usr/local/lib/libopencv_sfm.4.5.4.dylib  /usr/local/lib/libopencv_stereo.4.5.4.dylib  /usr/local/lib/libopencv_structured_light.4.5.4.dylib  /usr/local/lib/libopencv_superres.4.5.4.dylib  /usr/local/lib/libopencv_surface_matching.4.5.4.dylib  /usr/local/lib/libopencv_tracking.4.5.4.dylib  /usr/local/lib/libopencv_videostab.4.5.4.dylib  /usr/local/lib/libopencv_viz.4.5.4.dylib  /usr/local/lib/libopencv_wechat_qrcode.4.5.4.dylib  /usr/local/lib/libopencv_xfeatures2d.4.5.4.dylib  /usr/local/lib/libopencv_xobjdetect.4.5.4.dylib  /usr/local/lib/libopencv_xphoto.4.5.4.dylib  examples/ConvOpt/libConv2D.a  /usr/local/lib/libopencv_shape.4.5.4.dylib  /usr/local/lib/libopencv_highgui.4.5.4.dylib  /usr/local/lib/libopencv_datasets.4.5.4.dylib  /usr/local/lib/libopencv_plot.4.5.4.dylib  /usr/local/lib/libopencv_text.4.5.4.dylib  /usr/local/lib/libopencv_ml.4.5.4.dylib  /usr/local/lib/libopencv_phase_unwrapping.4.5.4.dylib  /usr/local/lib/libopencv_optflow.4.5.4.dylib  /usr/local/lib/libopencv_ximgproc.4.5.4.dylib  /usr/local/lib/libopencv_video.4.5.4.dylib  /usr/local/lib/libopencv_videoio.4.5.4.dylib  /usr/local/lib/libopencv_imgcodecs.4.5.4.dylib  /usr/local/lib/libopencv_objdetect.4.5.4.dylib  /usr/local/lib/libopencv_calib3d.4.5.4.dylib  /usr/local/lib/libopencv_dnn.4.5.4.dylib  /usr/local/lib/libopencv_features2d.4.5.4.dylib  /usr/local/lib/libopencv_flann.4.5.4.dylib  /usr/local/lib/libopencv_photo.4.5.4.dylib  /usr/local/lib/libopencv_imgproc.4.5.4.dylib  /usr/local/lib/libopencv_core.4.5.4.dylib && :
ld: warning: ignoring file examples/ConvOpt/libConv2D.a, building for macOS-x86_64 but attempting to link with file built for unknown-unsupported file format ( 0x21 0x3C 0x61 0x72 0x63 0x68 0x3E 0x0A 0x23 0x31 0x2F 0x32 0x30 0x20 0x20 0x20 )
Undefined symbols for architecture x86_64:
  "__mlir_ciface_conv_2d", referenced from:
      _main in edge-detection.cpp.o
ld: symbol(s) not found for architecture x86_64
clang: error: linker command failed with exit code 1 (use -v to see invocation)
ninja: build stopped: subcommand failed.

Segmentation fault while running examples/BuddyLamma/buddy-llama-run

Describe the bug
Just trial run the example buddy-llama-run for the first time, while generating the first output token by _mlir_ciface_forward(), segmentation fault happens. From the backstrace, it crashed at __memmove_avx_unaligned_erms()

To Reproduce

  1. Build the code referred from the README
  2. get and run the app at build/bin/buddy-llama-run
  3. Segmentation fault happens

Expected behavior
Run outputs normally

Screenshots
Reading symbols from ./buddy-llama-run...
(gdb) r
Starting program: /home/code/buddy-mlir/build/bin/buddy-llama-run
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
LLaMA 2 Inference Powered by Buddy Compiler

Please send a message:

How old are you?

[Log] Vocab file: "/home/code/buddy-mlir/examples/BuddyLlama/vocab.txt"
[Log] Tokenize time: 39.2032ms
[Log] Loading params...
[Log] Params file: "/home/code/buddy-mlir/examples/BuddyLlama/arg0.data"
[Log] Params load time: 37.844s

[New Thread 0x7ff385662b80 (LWP 216262)]
[New Thread 0x7ff384e60c00 (LWP 216263)]
[New Thread 0x7ff37fffec80 (LWP 216264)]
[New Thread 0x7ff37f7fcd00 (LWP 216265)]
[New Thread 0x7ff37effad80 (LWP 216266)]
[New Thread 0x7ff37e7f8e00 (LWP 216267)]
[New Thread 0x7ff37dff6e80 (LWP 216268)]
[New Thread 0x7ff37d7f4f00 (LWP 216269)]
[New Thread 0x7ff37cff2f80 (LWP 216270)]
[New Thread 0x7ff363fff000 (LWP 216271)]
[New Thread 0x7ff3637fd080 (LWP 216272)]
[New Thread 0x7ff362ffb100 (LWP 216273)]
[New Thread 0x7ff3627f9180 (LWP 216274)]
[New Thread 0x7ff361ff7200 (LWP 216275)]
[New Thread 0x7ff3617f5280 (LWP 216276)]
[New Thread 0x7ff360ff3300 (LWP 216277)]
[New Thread 0x7ff337fff380 (LWP 216278)]
[New Thread 0x7ff3377fd400 (LWP 216279)]
[New Thread 0x7ff336ffb480 (LWP 216280)]

Thread 1 "buddy-llama-run" received signal SIGSEGV, Segmentation fault.
__memmove_avx_unaligned_erms () at ../sysdeps/x86_64/multiarch/memmove-vec-unaligned-erms.S:262
262 ../sysdeps/x86_64/multiarch/memmove-vec-unaligned-erms.S: No such file or directory.
(gdb) bt
#0 __memmove_avx_unaligned_erms () at ../sysdeps/x86_64/multiarch/memmove-vec-unaligned-erms.S:262
#1 0x0000555555592d31 in forward ()
#2 0x00005555557c9103 in _mlir_ciface_forward ()
#3 0x0000555555559055 in main () at /home/oppo/code/buddy-mlir/examples/BuddyLlama/llama-main.cpp:155

Desktop (please complete the following information):

  • OS: Ubuntu
  • Version: 20.04

Additional context
Add any other context about the problem here.

[RVV] Current solution to support i32 operands in RVVOp for generating RV32 intrinsics

I intend to make RVV dialect operations support i32 operands for generating RV32 intrinsics. The choice of solutions are related to user experience so I want to share and discuss about it.

Why it has to be i32 operands

Certain operands in current RVV dialect operations are Index type, such as vl (a crucial parameter for RVV to represent the current vector length to process). Once they are lowered, index will be converted to i64 in x86-64 machine and there is no way to customize it according to this discussion (correct me if there are any updates). If these operands are not presented in i32 type, it is RV64 intrinsics rather than RV32 intrinsics will be obtained after MLIR translation by default.

Applying integer conversion in lowering process seems not to work

To attain i32 operands, unconditional integer conversion from index to i32 can be applied by rewriter.create method like that:

Value vlCast = rewriter
                       .create<UnrealizedConversionCastOp>(
                           loadOp.getLoc(), rewriter.getIntegerType(32), vl)
                       .getResult(0);

However, the introduced builtin.unrealized_conversion_cast Op may be hard to remove under "i32->index->i64" patterns. Consider the following source code:

// test-original.mlir

func.func @main() -> index {
  %sew = arith.constant 2 : index
  %lmul = arith.constant 1 : index
  %avl = arith.constant 16 : index
  %vl = rvv.setvl %avl, %sew, %lmul : index
  %new_avl = arith.subi %avl, %vl : index
  return %new_avl : index
}

With integer conversion in lowering(see this experimental branch for more details), the generated code will be:

// test-unrealized.mlir

module attributes {llvm.data_layout = ""} {
  llvm.func @main() -> i64 {
    %0 = llvm.mlir.constant(16 : i32) : i32
    %1 = llvm.mlir.constant(2 : i32) : i32
    %2 = llvm.mlir.constant(1 : i32) : i32
    %3 = llvm.mlir.constant(16 : index) : i64
    %4 = "rvv.intr.vsetvli"(%0, %1, %2) : (i32, i32, i32) -> i32
    %5 = builtin.unrealized_conversion_cast %4 : i32 to index
    %6 = builtin.unrealized_conversion_cast %5 : index to i64
    %7 = llvm.sub %3, %6  : i64
    llvm.return %7 : i64
  }
}

After that, try to run ${BUDDY_OPT} test-unrealized.mlir -reconcile-unrealized-casts will lead to "failed to legalize operation" error, since -reconcile-unrealized-casts pass can only deal with "i64->index->i64"(ABA-like), "i32->i32"(AA-like) and similar kinds of patterns.

Noting that replacing builtin.unrealized_conversion_cast with arith.index_cast would not prevent this error, since it has nothing to do with the lowering process. It seems that it is just not a good idea to directly use integer conversion to get i32 operands.

Personal suggestion

Not long ago I had a discussion with @zhanghb97 and we originally thought that it would be user-friendly to add an option in -lower-rvv pass to enable rv32 intrinsics generations(e.g. -lower-rvv="rv32", see this experimental branch for more details). Nevertheless, as shown above, such a design requires the application of integer conversion, which is not an easy task.

If that doesn't work, users have to somewhat modify their.mlir scripts based on original RV64 scripts in order to generate RV32 intrinsics, and it would be more suitable to just extend operands' data type in RVV dialect instead of seeking for index->i32 conversion.

This is what this PR is about. I think this solution is good enough for now to support generation of RV32 intrinsics. Any suggestions and advice are welcome.

cmake project falid

โ€‹when rum this commard:

cmake -G Ninja ..
-DMLIR_DIR=$PWD/../llvm/build/lib/cmake/mlir
-DLLVM_DIR=$PWD/../llvm/build/lib/cmake/llvm
-DLLVM_ENABLE_ASSERTIONS=ON
-DCMAKE_BUILD_TYPE=RELEASE

Error:

CMake Error at llvm/build/lib/cmake/llvm/AddLLVM.cmake:612 (add_library):
Cannot find source file:

buddy-mlir-main/build/backend/llvm/lib/Target/RISCV/RISCVMacroFusion.cpp
Tried extensions .c .C .c++ .cc .cpp .cxx .cu .mpp .m .M .mm .ixx .cppm .h
.hh .h++ .hm .hpp .hxx .in .txx .f .F .for .f77 .f90 .f95 .f03 .hip .ispc
Call Stack (most recent call first):
llvm/build/lib/cmake/llvm/AddLLVM.cmake:874 (llvm_add_library)
llvm/build/lib/cmake/llvm/AddLLVM.cmake:849 (add_llvm_library)
backend/llvm/lib/Target/RISCV/CMakeLists.txt:20 (add_llvm_component_library)
backend/llvm/lib/Target/RISCV/CMakeLists.txt:87 (buddy_add_llvm_target)

CMake Error at llvm/build/lib/cmake/llvm/AddLLVM.cmake:612 (add_library):
Cannot find source file:

buddy-mlir-main/llvm/llvm/lib/CodeGen/CodeGenPassBuilder.cpp

Tried extensions .c .C .c++ .cc .cpp .cxx .cu .mpp .m .M .mm .ixx .cppm .h
.hh .h++ .hm .hpp .hxx .in .txx .f .F .for .f77 .f90 .f95 .f03 .hip .ispc
Call Stack (most recent call first):
llvm/build/lib/cmake/llvm/AddLLVM.cmake:874 (llvm_add_library)
llvm/build/lib/cmake/llvm/AddLLVM.cmake:849 (add_llvm_library)
backend/llvm/lib/CodeGen/CMakeLists.txt:3 (add_llvm_component_library)

CMake Error at llvm/build/lib/cmake/llvm/AddLLVM.cmake:612 (add_library):
Cannot find source file:

buddy-mlir-main/llvm/llvm/lib/Analysis/VFABIDemangling.cpp

โ€‹

[build system] buddy-translate fail to translate MLIR into LLVM IR

Describe the bug
When building buddy-mlir in one-step build mode (build as an external LLVM project), the provided buddy-translate tool fails to translate the RVV dialect to LLVM IR with the following error:

Expected an argument of Vector Type
UNREACHABLE executed at /build/buddy-mlir/llvm/llvm/lib/IR/Function.cpp:1402!
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0.      Program arguments: buddy-translate --buddy-to-llvmir
...

This issue cannot be reproduced using the two-step build strategy, so something may be missing when building buddy-mlir using the one-step build strategy.

To Reproduce

Build buddy-mlir following the one-step-build-strategy guide.

Then run the pipeline and compile into LLVM IR:

buddy-opt stripmining.mlir \
  --lower-affine  --convert-scf-to-cf --convert-math-to-llvm \
  --lower-vector-exp --lower-rvv=rv32 --convert-vector-to-llvm \
  --finalize-memref-to-llvm --convert-arith-to-llvm --convert-func-to-llvm --reconcile-unrealized-casts \
  -o stripmining-lowered.mlir

buddy-translate stripmining-lowered.mlir --buddy-to-llvmir stripminig.llvmir
  • The stripmining.mlir file
memref.global "private" @gv_i32 : memref<32768xi32>

func.func @test() -> i32 {
  %mem_i32 = memref.get_global @gv_i32 : memref<32768xi32>
  %c0 = arith.constant 0 : index
  %c0_i32 = arith.constant 0 : i32
  %i2 = arith.constant 2 : i32

  // Configure the register.
  // SEW = 32
  %sew = arith.constant 2 : i32
  // LMUL = 2
  %lmul = arith.constant 1 : i32

  %init_avl = memref.dim %mem_i32, %c0 : memref<32768xi32>
  %init_avl_i32 = arith.index_cast %init_avl : index to i32
  %init_idx = arith.constant 0 : index
  %res = memref.get_global @gv_i32 : memref<32768xi32>

  // While loop for strip-mining.
  %a1, %a2 = scf.while (%avl = %init_avl_i32, %idx = %init_idx) : (i32, index) -> (i32, index) {
    // If avl greater than zero.
    %cond = arith.cmpi sgt, %avl, %c0_i32 : i32
    // Pass avl, idx to the after region.
    scf.condition(%cond) %avl, %idx : i32, index
  } do {
  ^bb0(%avl : i32, %idx : index):
    // Perform the calculation according to the vl.
    %vl = rvv.setvl %avl, %sew, %lmul : i32
    %input_vector = rvv.load %mem_i32[%idx], %vl : memref<32768xi32>, vector<[8]xi32>, i32
    %result_vector = rvv.add %input_vector, %i2, %vl : vector<[8]xi32>, i32, i32
    rvv.store %result_vector, %res[%idx], %vl : vector<[8]xi32>, memref<32768xi32>, i32
    // Update idx and avl.
    %vl_ind = arith.index_cast %vl : i32 to index
    %new_idx = arith.addi %idx, %vl_ind : index
    %new_avl = arith.subi %avl, %vl : i32
    scf.yield %new_avl, %new_idx : i32, index
  }
  %result = vector.load %res[%c0] : memref<32768xi32>, vector<8xi32>

  %mask = arith.constant dense<1> : vector<8xi1>
  %c1_i32 = arith.constant 1 : i32
  %evl = arith.constant 8: i32
  %res_reduce_add_mask_driven = "llvm.intr.vp.reduce.add" (%c1_i32, %result, %mask, %evl) :
         (i32, vector<8xi32>, vector<8xi1>, i32) -> i32

  return %res_reduce_add_mask_driven : i32
}

[Build Failed] error: no viable conversion from 'ArrayRef<int64_t>' to 'ArrayRef<long>'

Describe the bug
Build failed on macOS 13.2
../midend/lib/Dialect/Gemmini/Transforms/LegalizeForLLVMExport.cpp

../midend/lib/Dialect/Gemmini/Transforms/LegalizeForLLVMExport.cpp:597:26: error: no viable conversion from 'ArrayRef<int64_t>' to 'ArrayRef<long>'
    llvm::ArrayRef<long> aArrayShape = aArrayType.getShape();
                         ^             ~~~~~~~~~~~~~~~~~~~~~

Desktop (please complete the following information):

  • OS: macOS
  • Version 13.2

Additional context

../midend/lib/Dialect/Gemmini/Transforms/LegalizeForLLVMExport.cpp:597:26: error: no viable conversion from 'ArrayRef<int64_t>' to 'ArrayRef<long>'
    llvm::ArrayRef<long> aArrayShape = aArrayType.getShape();
                         ^             ~~~~~~~~~~~~~~~~~~~~~
../llvm/llvm/include/llvm/ADT/ArrayRef.h:41:40: note: candidate constructor (the implicit copy constructor) not viable: no known conversion from '::llvm::ArrayRef<int64_t>' (aka 'ArrayRef<long long>') to 'const llvm::ArrayRef<long> &' for 1st argument
  class LLVM_GSL_POINTER [[nodiscard]] ArrayRef {
                                       ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:41:40: note: candidate constructor (the implicit move constructor) not viable: no known conversion from '::llvm::ArrayRef<int64_t>' (aka 'ArrayRef<long long>') to 'llvm::ArrayRef<long> &&' for 1st argument
../llvm/llvm/include/llvm/ADT/ArrayRef.h:70:18: note: candidate constructor not viable: no known conversion from '::llvm::ArrayRef<int64_t>' (aka 'ArrayRef<long long>') to 'std::nullopt_t' for 1st argument
    /*implicit*/ ArrayRef(std::nullopt_t) {}
                 ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:73:18: note: candidate constructor not viable: no known conversion from '::llvm::ArrayRef<int64_t>' (aka 'ArrayRef<long long>') to 'const long &' for 1st argument
    /*implicit*/ ArrayRef(const T &OneElt)
                 ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:114:28: note: candidate constructor not viable: no known conversion from '::llvm::ArrayRef<int64_t>' (aka 'ArrayRef<long long>') to 'const std::initializer_list<long> &' for 1st argument
    constexpr /*implicit*/ ArrayRef(const std::initializer_list<T> &Vec)
                           ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:88:18: note: candidate template ignored: could not match 'SmallVectorTemplateCommon' against 'ArrayRef'
    /*implicit*/ ArrayRef(const SmallVectorTemplateCommon<T, U> &Vec)
                 ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:94:18: note: candidate template ignored: could not match 'vector' against 'ArrayRef'
    /*implicit*/ ArrayRef(const std::vector<T, A> &Vec)
                 ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:99:28: note: candidate template ignored: could not match 'array' against 'ArrayRef'
    /*implicit*/ constexpr ArrayRef(const std::array<T, N> &Arr)
                           ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:104:28: note: candidate template ignored: could not match 'const long[N]' against '::llvm::ArrayRef<int64_t>' (aka 'ArrayRef<long long>')
    /*implicit*/ constexpr ArrayRef(const T (&Arr)[N]) : Data(Arr), Length(N) {}
                           ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:124:5: note: candidate template ignored: could not match 'type-parameter-0-0 *' against 'long long'
    ArrayRef(const ArrayRef<U *> &A,
    ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:133:18: note: candidate template ignored: could not match 'SmallVectorTemplateCommon' against 'ArrayRef'
    /*implicit*/ ArrayRef(
                 ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:142:5: note: candidate template ignored: could not match 'vector' against 'ArrayRef'
    ArrayRef(const std::vector<U *, A> &Vec,
    ^
../llvm/llvm/include/llvm/ADT/ArrayRef.h:285:5: note: candidate function
    operator std::vector<T>() const {
    ^

Novel warning introduced due to usage of `Img<float, 2>` instead of `MemRef_descriptor` in `correlation2D.cpp`

Describe the bug
The following warning is introduced when we use Img<float, 2> instead of MemRef_descriptor for the input variable in correlation2D.cpp.

root@3eb35749d59c:~/forks/buddy-mlir/build# ninja correlation2D
[1/2] Building CXX object examples/DIPDialect/CMakeFiles/correlation2D.dir/correlation2D.cpp.o
In file included from ../include/Interface/buddy/core/Container.h:88,
                 from ../include/Interface/buddy/core/ImageContainer.h:24,
                 from ../examples/DIPDialect/correlation2D.cpp:30:
../lib/Interface/core/Container.cpp: In function โ€˜bool testImplementation(int, char**, std::ptrdiff_t, std::ptrdiff_t, std::ptrdiff_t)โ€™:
../lib/Interface/core/Container.cpp:164:3: warning: โ€˜input.Img<float, 2>::<anonymous>.MemRef<float, 2>::allocatedโ€™ may be used uninitialized in this function [-Wmaybe-uninitialized]
  164 |   delete[] allocated;
      |   ^~~~~~
../examples/DIPDialect/correlation2D.cpp:106:17: note: โ€˜input.Img<float, 2>::<anonymous>.MemRef<float, 2>::allocatedโ€™ was declared here
  106 |   Img<float, 2> input(image);
      |                 ^~~~~
[2/2] Linking CXX executable bin/correlation2D

To Reproduce
Run correlation2d.cpp from this branch.

Expected behavior
Ideally, no new warning should be introduced due to porting of specialized container(s) in example(s).

Desktop (please complete the following information):

  • OS: Ubuntu
  • Version : 20.04
  • Head commit of local llvm build : 0083a02839c3365af4b406042d35b82ded6b505c

Additional context
Continuation of discussion from this comment.

[Testing] Improve Image Container tests

Is your feature request related to a problem? Please describe.
Current tests for the image container do not test the normalization functionality and the accuracy of stored pixel value when directly compared with the original image for all pixels.

Describe the solution you'd like
Add simple tests to check whether the stored values in wrapped memref (Img container) is in agreement with the original image pixel values (with and without norm flag).

MLIR GPU Makefile and gpu-mma-run

Describe the bug
The makefile gpu-mma-run comman doesn't actually run the gpu-mma.mlir file

To Reproduce
Follow the readme to setup buddy mlir.
cd examples/MLIRGPU
make gpu-mma-run

Expected behavior
Currently, it outputs the same thing as make gpu-launch-func-run. The examples/MLIRGPU/makefile has

gpu-mma-lower:
	@${MLIR_OPT} gpu-launch-func.mlir -gpu-kernel-outlining | \

but I think it's supposed to be

gpu-mma-lower:
	@${MLIR_OPT} gpu-mma.mlir -gpu-kernel-outlining | \

However, when I do change it to run the correct file i'm getting errors.

make gpu-mma-run       
LLVM ERROR: Cannot select: intrinsic %llvm.nvvm.wmma.m16n16k16.store.d.row.stride.f32                                                                                                                          
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.                                                                                                    Stack dump:                                                                                            
0.      Program arguments: ../../llvm/build/bin/mlir-opt -pass-pipeline=builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-nvvm,gpu-to-cubin))
1.      Running pass 'Function Pass Manager' on module 'LLVMDialectModule'.                                                                                                                                    
2.      Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@kernel1'                                                                                                                      #0 0x000055cd7a88aadf llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (../../llvm/build/bin/mlir-opt+0x214adf)
 #1 0x000055cd7a888594 SignalHandler(int) Signals.cpp:0:0                                                                                                                                                      
 #2 0x00007fa43ee66420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)                                                                                                                             #3 0x00007fa43d02a00b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1
 #4 0x00007fa43d009859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7                                                                                                                                 #5 0x000055cd7a80b3aa llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0
 #6 0x000055cd7c51533d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (../../llvm/build/bin/mlir-opt+0x1e9f33d)                                                                                         #7 0x000055cd7c517ca2 llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (../../llvm/build/bin/mlir-opt+0x1ea1ca2)                                                    #8 0x000055cd7c512a90 llvm::SelectionDAGISel::DoInstructionSelection() (../../llvm/build/bin/mlir-opt+0x1e9ca90)                                                                                               #9 0x000055cd7c520651 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (../../llvm/build/bin/mlir-opt+0x1eaa651)
#10 0x000055cd7c524a78 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (../../llvm/build/bin/mlir-opt+0x1eaea78)                                                                           #11 0x000055cd7c5269ff llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0         
#12 0x000055cd7c73344e llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0                                                                                         #13 0x000055cd7d350985 llvm::FPPassManager::runOnFunction(llvm::Function&) (../../llvm/build/bin/mlir-opt+0x2cda985)
#14 0x000055cd7d350bc9 llvm::FPPassManager::runOnModule(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdabc9)                              
#15 0x000055cd7d351082 llvm::legacy::PassManagerImpl::run(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdb082)
#16 0x000055cd7ac89d33 mlir::gpu::SerializeToBlobPass::translateToISA[abi:cxx11](llvm::Module&, llvm::TargetMachine&) (../../llvm/build/bin/mlir-opt+0x613d33)
#17 0x000055cd7ac89f7a mlir::gpu::SerializeToBlobPass::runOnOperation() (../../llvm/build/bin/mlir-opt+0x613f7a)
#18 0x000055cd7c00b6ba mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x19956ba)
#19 0x000055cd7c00bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::Pipe
lineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9)
#20 0x000055cd7c00a936 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) (../../llvm/build/bin/mlir-opt+0x1994936)
#21 0x000055cd7c00b59a mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x199559a)
#22 0x000055cd7c00bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::Pipe
lineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9)
#23 0x000055cd7c00c6d9 mlir::PassManager::run(mlir::Operation*) (../../llvm/build/bin/mlir-opt+0x19966d9)
#24 0x000055cd7bffd05b performActions(llvm::raw_ostream&, bool, bool, llvm::SourceMgr&, mlir::MLIRContext*, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, bool, bool) (.constprop.0) MlirOptMai
n.cpp:0:0                                                                                                                                                                                                      
#25 0x000055cd7bffd7e1 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, bool, bool, bool, bool, bool, bool, llvm::function_ref<mlir::LogicalResu
lt (mlir::PassManager&)>, mlir::DialectRegistry&, llvm::ThreadPool*) MlirOptMain.cpp:0:0
#26 0x000055cd7bffdbd3 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::callback_fn<mlir::MlirOp
tMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, bool, bool, bool,
 bool, bool, bool, bool)::'lambda'(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::Memor
yBuffer>>, llvm::raw_ostream&) MlirOptMain.cpp:0:0                                                     
#27 0x000055cd7c0b7ed4 mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, st
d::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>, llvm::raw_ostream&, bool, bool) (../../llvm/build/bin/mlir-opt+0x1a41ed4)
#28 0x000055cd7bffc2a1 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::PassPipelineCLParser const&, mlir::DialectRegistry&, bool, boo
l, bool, bool, bool, bool, bool, bool) (../../llvm/build/bin/mlir-opt+0x19862a1)
#29 0x000055cd7bffdfd7 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&, bool) (../../llvm/build/bin/mlir-opt+0x1987fd7)
#30 0x000055cd7a812386 main (../../llvm/build/bin/mlir-opt+0x19c386)
#31 0x00007fa43d00b083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3
#32 0x000055cd7a86642e _start (../../llvm/build/bin/mlir-opt+0x1f042e)
Aborted (core dumped)                                                                                                                                                                                                   Error: entry point not found                                                                                                                                                                                            make: *** [makefile:73: gpu-mma-run] Error 1                                                                                                                                                                      
abhi@abhi-MS-7C56:~/courses/ece527-SOC-Design/final_project/buddy-mlir/examples/MLIRGPU$ make gpu-mma-run                                                                                                               LLVM ERROR: Cannot select: intrinsic %llvm.nvvm.wmma.m16n16k16.store.d.row.stride.f32                                                                                                                                   PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.                                                                                                             Stack dump:                                                                                                                                                                                                         
0.      Program arguments: ../../llvm/build/bin/mlir-opt -pass-pipeline=builtin.module(gpu.module(strip-debuginfo,convert-gpu-to-nvvm,gpu-to-cubin))                                                                    1.      Running pass 'Function Pass Manager' on module 'LLVMDialectModule'.                                                                                                                                             2.      Running pass 'NVPTX DAG->DAG Pattern Instruction Selection' on function '@kernel1'                                                                                                                               #0 0x00005633cdd1aadf llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) (../../llvm/build/bin/mlir-opt+0x214adf)                                                                                                      #1 0x00005633cdd18594 SignalHandler(int) Signals.cpp:0:0                                              
 #2 0x00007f0246ca9420 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x14420)                                                                                                                                      #3 0x00007f0244e6d00b raise /build/glibc-SzIz7B/glibc-2.31/signal/../sysdeps/unix/sysv/linux/raise.c:51:1                                                                                                               #4 0x00007f0244e4c859 abort /build/glibc-SzIz7B/glibc-2.31/stdlib/abort.c:81:7                                                                                                                                          #5 0x00005633cdc9b3aa llvm::ConvertUTF8toUTF32(unsigned char const**, unsigned char const*, unsigned int**, unsigned int*, llvm::ConversionFlags) (.cold) ConvertUTF.cpp:0:0                                       
 #6 0x00005633cf9a533d llvm::SelectionDAGISel::CannotYetSelect(llvm::SDNode*) (../../llvm/build/bin/mlir-opt+0x1e9f33d)                                                                                                  #7 0x00005633cf9a7ca2 llvm::SelectionDAGISel::SelectCodeCommon(llvm::SDNode*, unsigned char const*, unsigned int) (../../llvm/build/bin/mlir-opt+0x1ea1ca2)                                                             #8 0x00005633cf9a2a90 llvm::SelectionDAGISel::DoInstructionSelection() (../../llvm/build/bin/mlir-opt+0x1e9ca90)                                                                                                        #9 0x00005633cf9b0651 llvm::SelectionDAGISel::CodeGenAndEmitDAG() (../../llvm/build/bin/mlir-opt+0x1eaa651)                                                                                                        
#10 0x00005633cf9b4a78 llvm::SelectionDAGISel::SelectAllBasicBlocks(llvm::Function const&) (../../llvm/build/bin/mlir-opt+0x1eaea78)                                                                                    #11 0x00005633cf9b69ff llvm::SelectionDAGISel::runOnMachineFunction(llvm::MachineFunction&) (.part.0) SelectionDAGISel.cpp:0:0                                                                                          #12 0x00005633cfbc344e llvm::MachineFunctionPass::runOnFunction(llvm::Function&) (.part.0) MachineFunctionPass.cpp:0:0                                                                                                  #13 0x00005633d07e0985 llvm::FPPassManager::runOnFunction(llvm::Function&) (../../llvm/build/bin/mlir-opt+0x2cda985)                                                                                                    #14 0x00005633d07e0bc9 llvm::FPPassManager::runOnModule(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdabc9)                                                                                                        #15 0x00005633d07e1082 llvm::legacy::PassManagerImpl::run(llvm::Module&) (../../llvm/build/bin/mlir-opt+0x2cdb082)                                                                                                      #16 0x00005633ce119d33 mlir::gpu::SerializeToBlobPass::translateToISA[abi:cxx11](llvm::Module&, llvm::TargetMachine&) (../../llvm/build/bin/mlir-opt+0x613d33)                                                   
#17 0x00005633ce119f7a mlir::gpu::SerializeToBlobPass::runOnOperation() (../../llvm/build/bin/mlir-opt+0x613f7a)                                                                                                        #18 0x00005633cf49b6ba mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x19956ba)                                         #19 0x00005633cf49bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9)                                                                                                                                                             
#20 0x00005633cf49a936 mlir::detail::OpToOpPassAdaptor::runOnOperationAsyncImpl(bool) (../../llvm/build/bin/mlir-opt+0x1994936)                                                                                         #21 0x00005633cf49b59a mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) (../../llvm/build/bin/mlir-opt+0x199559a)                                         #22 0x00005633cf49bcb9 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) (../../llvm/build/bin/mlir-opt+0x1995cb9)                                                                                                                                                                 #23 0x00005633cf49c6d9 mlir::PassManager::run(mlir::Operation*) (../../llvm/build/bin/mlir-opt+0x19966d9)                                                                                                          
#24 0x00005633cf48d05b performActions(llvm::raw_ostream&, bool, bool, llvm::SourceMgr&, mlir::MLIRContext*, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, bool, bool) (.constprop.0) MlirOptMain.cpp:0:0#25 0x00005633cf48d7e1 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, bool, bool, bool, bool, bool, bool, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, llvm::ThreadPool*) MlirOptMain.cpp:0:0                             
#26 0x00005633cf48dbd3 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::callback_fn<mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (mlir::PassManager&)>, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool)::'lambda'(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) MlirOptMain.cpp:0:0                                                                                                                                                                                                 #27 0x00005633cf547ed4 mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>, llvm::raw_ostream&, bool, bool) (../../llvm/build/bin/mlir-opt+0x1a41ed4)                                                                                          
#28 0x00005633cf48c2a1 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::PassPipelineCLParser const&, mlir::DialectRegistry&, bool, bool, bool, bool, bool, bool, bool, bool) (../../llvm/build/bin/mlir-opt+0x19862a1)                                                                                                                                                 #29 0x00005633cf48dfd7 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&, bool) (../../llvm/build/bin/mlir-opt+0x1987fd7)                                                                          #30 0x00005633cdca2386 main (../../llvm/build/bin/mlir-opt+0x19c386)                                    
#31 0x00007f0244e4e083 __libc_start_main /build/glibc-SzIz7B/glibc-2.31/csu/../csu/libc-start.c:342:3                                                                                                                   #32 0x00005633cdcf642e _start (../../llvm/build/bin/mlir-opt+0x1f042e)                                                                                                                                                  Aborted (core dumped)                                                                                                                                                                                                   Error: entry point not found                                                                            
make: *** [makefile:73: gpu-mma-run] Error 1

so it looks like this example is broken.

Desktop (please complete the following information):

  • OS: Ubuntu
  • Version 20.04

[build system] get rid of two-phase building

Is your feature request related to a problem? Please describe.

To build current version of buddy, two phases are required: first phase to build LLVM components, second phase to build buddy itself.

When trying to package buddy compiler with Nix, in my first attempt, I build LLVM as a separate package (or derivation in Nix jargon), then build buddy with LLVM_DIR and MLIR_DIR set to some path in the package installation directory. However, it failed to build buddy because what buddy requires is not only these installed files, but also some non-installed cpp files, even include/llvm/Config/config.h.

From the packager's perspective, the two-phase behavior and internal file dependency result in cumbersome and tricky packaging process (see https://github.com/sequencer/vector/blob/master/nix/buddy-mlir.nix#L34-L41). This is partially due to that Nix believes that one cmake will control everything. From a developer's perspective, the build manager (i.e. cmake) is not fully responsible for the state of the file tree. It will potentially cause cache inconsistency and portability issues.

Describe the solution you'd like

I propose two possible solutions:

  1. Treat LLVM as a cmake submodule or subdirectory. One CMakeList.txt takes care of both LLVM and buddy build. I.e. build LLVM things into main build directory instead of into llvm/build. This might also reduce the total building time, because unused parts in LLVM will not be built.
  2. Get rid of dependencies of LLVM internal files, use a installed LLVM tree. This makes things cleaner and perfect for packager. However, the developer requires more effort (execute cmake --install or something) unless first solution is applied.

I am not quite sure whether above two solutions will work, welcome for more suggestions.

The `torch` version used in `requirements.txt` seems to be removed

Describe the bug
I am trying to reproduce the example under examples/BuddyLlama, but I found that the torch version used in requirements.txt seems to be removed.

To Reproduce

  • Create a new conda env named llama and activate it
$ conda create -n llama
$ conda activate llama
  • Intall the packages needed
$ cd buddy-mlir
$ pip install -r requirements.txt

Then we can see the error.

Expected behavior
torch is successfully installed.

Screenshots
image

Desktop (please complete the following information):

  • OS: Ubuntu
  • Version 22.04

Additional context
I am trying another "dev" version of torch. However, for the sake of generality, maybe we should use a stable version of torch to make the reproduction of this example more stable?

[DAP] Bug in IIR benchmark.

Describe the bug
The buddy_iir benchmark cause the error "munmap_chunk(): invalid pointer".

To Reproduce
Run ./audio-processing-benchmark under /buddy-benchmark/build/bin.

Expected behavior

  1. Generate result picture for all benchmarks( include buddy_iir benchmark).
  2. The audio-processing-benchmark shouldnt crush.

Screenshots
c9565afe1aece62ca0322a4bce6d8f8d_

Server

  • OS: Ubuntu
  • Version 20.04

Additional context
This problem happens in the benchmark process, but it seems that something was wrong in the lowering pass.

Sync to latest llvm/mlir to resolve RVV environment configuration problems

Currently the llvm version@8f966ce that buddy-mlir relies on has problems when cross-compiling MLIR to RISC-V, which makes it unable to build "build-cross-mlir" to verify the RVV-related cases(see this discourse created by YangSen @linuxlonelyeagle for more details). After verification, using the most recent commit @55662b2 can fix this issue, so we may need to find a chance to sync to upstream in recent times. @zhanghb97

Allow destination image to have different type than input for Corr2d operation

Is your feature request related to a problem? Please describe.
During the review of #63 there was a suggestion to allow a destination image to be different from input and kernel.

Currently, the solution introduced by #63 expects that input, kernel, output and constant has the same type with a corresponding check in LowerDIPPass.cpp for the Corr2d operation.

Describe the solution you'd like
My suggestion would be to introduce a similar to openCV parameter ddepth to express a desired depth of output. See more
https://docs.opencv.org/4.x/d4/d86/group__imgproc__filter.html#filter_depths

CC @meshtag

Unify the coding style of using size_t or std::size_t

size_t and std::size_t are both aliases of unsigned long in normal cases but defined in different standards. Though differences between them are subtle, unifications of usage is good for code style and code readability. Should we need clarification about it?

Currently buddy-mlir is using size_t and std::size_t combinely (e.g. in Container.cpp), which is somethat misleading and one is better to be replaced with the other. @meshtag and I have discussed about it and summarized that std::size_t is more friendly and simple for the compiler while size_t is neater for developers to use. Considerations remain open.

examples mobilenet V3 result error

Describe the bug
A clear and concise description of what the bug is.
่ฏ†ๅˆซ็š„dog ไฝ†ๆ˜ฏ็ป“ๆžœ้”™่ฏฏ
To Reproduce
Steps to reproduce the behavior.
ๆ นๆฎๆ–‡ๆกฃๅฎŒๆˆ็ผ–่ฏ‘่ฟ่กŒ+
Expected behavior
A clear and concise description of what you expected to happen.

Screenshots
If applicable, add screenshots to help explain your problem.
image
Desktop (please complete the following information):

  • OS: Win11 pro workstation
  • Version 23H2

Additional context
Add any other context about the problem here.

typeinfo for mlir::Dialect

Describe the bug
when i use the mlir api and compile the code, i meet a problem undefined reference to `typeinfo for mlir::Dialect'.

To Reproduce
$:cmake -G Ninja .. \ -DMLIR_DIR=$PWD/../../llvm/build/lib/cmake/mlir \ -DLLVM_DIR=$PWD/../../llvm/build/lib/cmake/llvm
$:ninja
$:/usr/bin/ld: CMakeFiles/toyc.dir/mlir/Dialect.cpp.o:(.data.rel.ro._ZTIN4mlir3toy10ToyDialectE[_ZTIN4mlir3toy10ToyDialectE]+0x10): undefined reference to typeinfo for mlir::Dialect'

Desktop

  • OS: arch
  • Version [5.17.9]

Additional context
I create a storage,the storage has the code, you can put it in your buddy-mlir,then you can use cmake to build it, you will meet the problem, i hope you can help me solve it. thanks!

[DAP] Redesigning AudioFile interface.

Is your feature request related to a problem? Please describe.
Currently it is required to use a set of complicated functions to transfer the ownership between AudioFile and MemRef. It's not intuitive for end users.

Describe the solution you'd like

  1. Integrate AudioFile functions to AudioContainer. It requires heavy rework for all AudioFile functions, but the integration would be much better and no ownership transformation is needed. We may use a modular design, like custom loaders and savers for AudioContainer, thus the actually encoding and decoding functions could be added at any time.

Describe alternatives you've considered

  1. Let both of AudioFile and MemRef hold the pointer, then AudioContainer holds them both. Since they have similiar life cycle, this method is feasible and doesn't require much rework.
  2. Repair the function to handover ownership. Currently this function can be invoked only once. If encapsulated well, the user could use getMemRef and getAudioFile for unlimited times. This requires the least rework, but it would not be so intuitive for future developers.

Add support for Doxygen documentation

Is your feature request related to a problem? Please describe.
There is no developer documentation for the project at the moment. We should add this as the project becomes more mature.

Describe the solution you'd like
We can potentially use Doxygen to generate developer documentation for us. You can also use cool extensions like this one to get documentation boilerplate while writing the code. Perhaps we can make a rule of including developer documentation in all new PRs which will be accepted?
(/cc @zhanghb97 @FlagerLee @MbjYjbpivj @xlinsist @SForeKeeper @Avimitin et al)

Describe alternatives you've considered
Feel free to suggest any other alternatives that you might like, but we should stick to one tool for the entire codebase.

[CI] Add support for capturing code coverage

Is your feature request related to a problem? Please describe.
As of now, there is no mechanism to capture code coverage of the codebase embedded in the CI.

Describe the solution you'd like
We should have a GitHub Action job that takes care of this and submits coverage details as a comment in each PR.

Resolve unused variables.

There have been unused variables here that I hope to fix it.

buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:64:10: warning: unused variable 'typeConverter' [-Wunused-variable]
    auto typeConverter = getTypeConverter();
         ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:66:10: warning: unused variable 'awaitOp' [-Wunused-variable]
    auto awaitOp = rewriter.create<async::AwaitOp>(loc, operands[0]);
         ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:90:14: warning: unused variable 'host_register_op' [-Wunused-variable]
        auto host_register_op = rewriter.create<gpu::HostRegisterOp>(loc, memref_cast_op.getResult());
             ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:98:14: warning: unused variable 'host_register_op' [-Wunused-variable]
        auto host_register_op = rewriter.create<gpu::HostRegisterOp>(loc, memref_cast_op.getResult());
             ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:101:14: warning: unused variable 'vector_transfer_write_op' [-Wunused-variable]
        auto vector_transfer_write_op = rewriter.create<vector::TransferWriteOp>(loc, v, alloc_op.getResult(), indices);
             ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:105:14: warning: unused variable 'host_register_op' [-Wunused-variable]
        auto host_register_op = rewriter.create<gpu::HostRegisterOp>(loc, v);
             ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:110:14: warning: unused variable 'host_register_op' [-Wunused-variable]
        auto host_register_op = rewriter.create<gpu::HostRegisterOp>(loc, memref_cast_op.getResult());
             ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:162:12: warning: unused variable 'host_register_op' [-Wunused-variable]
      auto host_register_op = rewriter.create<gpu::HostRegisterOp>(loc, memref_cast_op.getResult());
           ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:172:14: warning: unused variable 'shape' [-Wunused-variable]
        auto shape = dyn_cast_or_null<TensorType>(t).getShape();
             ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:173:14: warning: unused variable 'ele_type' [-Wunused-variable]
        auto ele_type = dyn_cast_or_null<TensorType>(t).getElementType();
             ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:210:11: warning: unused variable 'bodyBlock' [-Wunused-variable]
    auto& bodyBlock = body.front();
          ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:227:11: warning: unused variable 'end' [-Wunused-variable]
    Value end = rewriter.create<arith::AddIOp>(loc, start, stepRangeInThread);
          ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:229:10: warning: unused variable 'sub_forOp' [-Wunused-variable]
    auto sub_forOp = rewriter.create<scf::ForOp>(loc, idx0, stepRangeInThread, idx1, forOp.getInitArgs(), 
         ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:328:20: warning: unused variable 'copy_op' [-Wunused-variable]
              auto copy_op = rewriter.create<memref::CopyOp>(loc, to_memref_op.getResult(), result_memrefs[i++]);
                   ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:333:20: warning: unused variable 'vector_transfer_write_op' [-Wunused-variable]
              auto vector_transfer_write_op = rewriter.create<vector::TransferWriteOp>(loc, mp.lookupOrNull<Value>(res), result_memrefs[i++], indices);
                   ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:336:20: warning: unused variable 'copy_op' [-Wunused-variable]
              auto copy_op = rewriter.create<memref::CopyOp>(loc, mp.lookupOrNull<Value>(res), result_memrefs[i++]);
                   ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:339:20: warning: unused variable 'store_op' [-Wunused-variable]
              auto store_op = rewriter.create<memref::StoreOp>(loc, mp.lookupOrNull<Value>(res), result_memrefs[i++]);
                   ^
buddy-mlir/midend/lib/Conversion/LowerSche/LowerSchePass.cpp:347:65: warning: format specifies type 'char *' but the argument has type 'llvm::StringRef' [-Wformat]
      printf("conversion from source %s has not implemented\n", sche_source);

[Midend]Use Declarative Pass Specification

Is your feature request related to a problem? Please describe.
Now, the midend Passes are written in c++, which I don't think is elegant.Such passes are too spread out.
Here are my reasons.

  • Using tablegen can make the code more stable
    The version of llvm is changing all the time, so maintaining the stability of the project is a problem, we should try our best to ensure stability and reduce the trouble when updating llvm.
  • Reduce the amount of code in c++
    As an example, the automatically generated pass class contains the following.
  static constexpr ::llvm::StringLiteral getArgumentName() {
    return ::llvm::StringLiteral("xxx");
  }
  ::llvm::StringRef getArgument() const override { return "xxxx"; }

  ::llvm::StringRef getDescription() const override { return "xxx"; }

  /// Returns the derived pass name.
  static constexpr ::llvm::StringLiteral getPassName() {
    return ::llvm::StringLiteral("xxxx");
  }
  ::llvm::StringRef getName() const override { return "xxx"; }

  static bool classof(const ::mlir::Pass *pass) {
    return pass->getTypeID() == ::mlir::TypeID::get<DerivedT>();
  }

  friend std::unique_ptr<::mlir::Pass> createxxx0Pass() {
    return std::make_unique<DerivedT>();
  }
  inline void registerPasses() {
  registerxxx0Pass();
  registerxxx1Pass();
  registerxxx2Pass();
 }
  • Make the code more intuitive
    This means I can use the pass like this
    buddy-opt xxx.mlir -my-pass="variable=false"
    I also know that this pass runs at the level of func.
def MyPass : Pass<"my-pass", "func::FuncOp"> {
  let summary = "xxxxx.";
  let options = [
   Option<"variable_", "variable", "bool",
           /*default=*/"true", "xxxxx">,
  ];
}
  • Use Declarative Pass Specification Many communities use it
    The above are just some of the benefits I can think of, but I don't think that's all of them.

Describe the solution you'd like
I think you just need to understand what's going on here,then define all passes in a tablegen file.If you have questions, feel free to leave a comment below.

Github Action update needed.

Describe the bug
In build action : HERE

Run echo "::set-output name=commit::$(git submodule status | awk '{print $1;}')"
  echo "::set-output name=commit::$(git submodule status | awk '{print $1;}')"
  shell: /usr/bin/bash -e {0}
Warning: The `set-output` command is deprecated and will be disabled soon. Please upgrade to using Environment Files. For more information see: https://github.blog/changelog/[2](https://github.com/buddy-compiler/buddy-mlir/actions/runs/3738006561/jobs/6343698885#step:4:2)022-10-11-github-actions-deprecating-save-state-and-set-output-commands/

This feature will be closed in 2023, so it's time to upgrade.

To Reproduce
Every Actions is effected now.

Bound check for operations with `memref`

Is your feature request related to a problem? Please describe.

Currently, elements accessing operations on memref do not provide any bound check, which I think there should be some check. For example, even this example below can successfully pass lowering/translate/run:

// Use global memref to ensure they are continuous in memory, for more 
// straight forward to show the result.
memref.global "private" @region_gv : memref<4xi32> = dense<[0, 1, 2, 3]>
memref.global "private" @canary_gv : memref<12xi32> = dense<202>
// 202 == 0xCA, means canary

func.func private @printMemrefI32(memref<*xi32>)

func.func @main() -> i32 {
  %region = memref.get_global @region_gv : memref<4xi32>
  %canary = memref.get_global @canary_gv : memref<12xi32>

  // ===================== test code ============================//
  %c1 = arith.constant 1 : i32
  %c5 = arith.constant 5 : index

  %region_print = memref.cast %region : memref<4xi32> to memref<*xi32>
  %canary_print = memref.cast %canary : memref<12xi32> to memref<*xi32>

  // NO ANY warning here! This type of access should be obvious for compiler.
  %val = memref.load %region[%c5] : memref<4xi32>
  vector.print %val : i32

  memref.store %c1, %region[%c5] : memref<4xi32>
  func.call @printMemrefI32(%region_print) : (memref<*xi32>) -> ()
  func.call @printMemrefI32(%canary_print) : (memref<*xi32>) -> ()

  %ret = arith.constant 0 : i32
  return %ret : i32
}

I have written more examples for this problem on my forked repo.

Describe the solution you'd like

A warning should be given when trying to lower/translate/run files contains out of bound access when an analysis is possible.

Describe alternatives you've considered

AFAIK, there are three ways we can do better about this:

  1. A pass may be added for bound checking. It seems easy to at least analyze the simplest cases (constant out-of-bound index for fixed length memref). And by keeping as long as the layout info inside or even between functions, we may also be possible to find out erased bound info like this examples.
  2. A pass to insert runtime bound check for every load/store or other operation on dynamic length memref. This is possible becasue striped layout dynamic length memref will be lowering with lengths in each dimension.
  3. Just leave it over, and let LLVM or the lower-level compiler check this. They should be able to detect gep or pointer out-of-bound operations.

Option 1 seems great, but it may "reinvent the wheel" by re-implementing the common check which is usually done at LLVM IR level. Option 2 is very costly in runtime, but having such a mode in debugging seems great. Option 3 will delay the time that such errors are detected, and make it harder to map the error back to the source language.

Additional context

Since most MLIR are created by compilers with higher level languages that may already ban unsafe memory operation, the bound check may be pushed up to be done before lowering to memref dialect or pushed down to be done after lowering to LLVM. The necessity of bound check in memref dialect level may need a deeper discussion.

In the current implement, the global memref will be continuous in memory, so it is the easiest way to see the effect of out-of-bound accesses. Stack allocated or normal memory allocated memref is not continuous in memory, which can be observed by reading printMemRefI32's output by mlir-cpu-runner.

Add caching mechanism for llvm build in GitHub Actions script

Is your feature request related to a problem? Please describe.
GitHub actions builds are terribly slow due to the initiation of fresh llvm build process on each commit.

Describe the solution you'd like
We should cache the llvm build data daily (since we update llvm dependency daily). This would enable faster CI builds when multiple commits are pushed during the same day.

Additional context
Ref :

# ToDo : Cache llvm build to reduce execution time on each run of virtual instance.

Since this issue doesn't require any significant knowledge of codebase (GitHub actions knowledge will be required), I am marking it as a Good first issue.

dip.corr_2d doesn't support correlation with irregular kernels.

Describe the bug

corr_2d op implemented in the DIP Dialect gives incorrect output for 1D kernels. For example kernels of the shape (n,1) or (1,n) where n is the number of elements.

To Reproduce

Steps to reproduce the behavior.
Test corr_2d with a 1D kernel an example could be [1, 2, 3] or any other kernel of shape (n, 1) or (1, n)

Expected behavior

OpenCV supports 1D kernels through filter2D. Currently, corr_2d is unable to support correlation with such 1D kernels.

Screenshots

image

Here a (3,1) kernel is used (it can be observed the results of corr_2d (above) and filter2D do not match.

image

Here a (1,3) kernel is used and the observations are similar to the case above.

(In the above example the input image is of shape (5, 5) with the anchor at (0,0) and stride size of 2.

Desktop (please complete the following information):

  • OS: Pop OS (Ubuntu)
  • Version: 22.04 LTS

Additional context

The corr_2d algorithm may require some changes to support 1D kernels; alternatively, a separate API for 1D convolutions can also be created.

[Examples] Bug in TOSA Examples

Describe the bug

  1. Currently pass-pipeline is being used with individual options in many examples of the TOSA dialect which is not allowed according to latest version of llvm.
  2. Current implementation of the resizeOp example doesn't work with latest llvm.

To Reproduce
Run Examples of TOSA dialect

Expected behavior
We should correct the pass-pipeline bug and make changes in the resizeOp example.

Screenshots
image

image

Desktop (please complete the following information):

  • OS: Pop Os (based on Ubuntu LTS)
  • Version 22.04 LTS

The problem of `affine.parallel`

Is your feature request related to a problem? Please describe.
Due to the difficulties in using affine.parallel, the outermost loop in the BatchMatMulOptimize of my PR does not utilize affine.parallel. However, this is crucial for the upcoming attempts to parallelize matmul using OpenMP. There's no open-source information available. I've tried writing code following the template, but it always results in errors

Describe the solution you'd like
Rewrite the outermost affine.for loop to affine.parallel

Incorrect result for small images

Obtained results are not in agreement with expected results in the case of 2D Correlation(conv2d in buddy MLIR) involving small images.

Input Image :

0 1 2 3 4 
5 6 7 8 9 
10 11 12 13 14 
15 16 17 18 19 
20 21 22 23 24 

Obtained output Image :

0 1 2 
8 59.3353 19.8358 
812552 9.5894e+11 12 

Used kernel :

float kernelAlign[9] = {1, 0, 0, 0, 0, 0, 0, 0, 0};

Commands used for compilation :

cmake -G Ninja .. -DBUDDY_CONV_OPT_EXAMPLES=ON -DBUDDY_CONV_OPT_STRIP_MINING=5 -DBUDDY_CONV_OPT_ATTR=avx512f

ninja edge-detection

Command used for execution :

./edge-detection ../../examples/conv-opt/images/test1.png result1.png

Add inbuilt support for multi-channel processing in dip.corr_2d lowering pass

Is your feature request related to a problem? Please describe.
Current implementation cannot process all channels of the input image at the same time. Thus, feeding each channel separately is a requirement/restriction of current API.

Describe the solution you'd like
The lowering pass should be able to handle multiple image channels on itself without the need of splitting and feeding individual channels of a multi-channelled image.

Additional context
Different alternatives such as creating dedicated operations for channel ordering, modifying inner loops of lowering pass, etc. can/should be considered before presenting a solution.

How to add new dialect to buddy-translate?

Hi, buddy-mlir! I want to translate the conv2d.mlir to llvm ir file to compare the difference of conv2d implemention between vectorization and non-vectorization(with linalg dialect). But when I try to run the command like this:

 buddy-translate --buddy-to-llvmir after-vec-conv2d.mlir -o after-vec-conv2d.ll 

It showed like:
image

Because the buddy-translate doesn't support the dialect linalg,affine and vector of mlir.
I've checked the CMakeLists.txt of buddy-translate:

get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
get_property(translation_libs GLOBAL PROPERTY MLIR_TRANSLATION_LIBS)

But I don't know how to add the dialect vector, affine and linalg to dialect_libs.

[DIP] [Examples] resize2D example is broken

Describe the bug
The resize2D build target appears to be broken. I am getting linking errors related to jpeg and png encoder and decoder methods.

To Reproduce
ninja resize2D

Expected behavior
The example should be built without any errors

Stack trace
Part containing the errors in the stacktrace is attached below

FAILED: bin/resize2D 
: && /usr/bin/c++ -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -fno-lifetime-dse -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wno-missing-field-initializers -Wimplicit-fallthrough -Wno-nonnull -Wno-class-memaccess -Wno-redundant-move -Wno-pessimizing-move -Wno-noexcept-type -Wdelete-non-virtual-dtor -Wsuggest-override -Wno-comment -Wno-misleading-indentation -Wctad-maybe-unsupported -fdiagnostics-color -ffunction-sections -fdata-sections -fPIC -fno-semantic-interposition -fvisibility-inlines-hidden -Werror=date-time -fno-lifetime-dse -Wall -Wextra -Wno-unused-parameter -Wwrite-strings -Wcast-qual -Wno-missing-field-initializers -Wimplicit-fallthrough -Wno-nonnull -Wno-class-memaccess -Wno-redundant-move -Wno-pessimizing-move -Wno-noexcept-type -Wdelete-non-virtual-dtor -Wsuggest-override -Wno-comment -Wno-misleading-indentation -Wctad-maybe-unsupported -fdiagnostics-color -ffunction-sections -fdata-sections -O3 -DNDEBUG  examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o -o bin/resize2D  lib/libBuddyLibDIP.a && :
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::JpegDecoder<float, 2ul>::~JpegDecoder()':
resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EED2Ev[_ZN3dip11JpegDecoderIfLm2EED5Ev]+0x2b): undefined reference to `jpeg_destroy_decompress'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::PngDecoder<float, 2ul>::~PngDecoder()':
resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EED2Ev[_ZN3dip10PngDecoderIfLm2EED5Ev]+0x7a): undefined reference to `png_destroy_read_struct'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::JpegDecoder<float, 2ul>::~JpegDecoder()':
resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EED0Ev[_ZN3dip11JpegDecoderIfLm2EED5Ev]+0x2c): undefined reference to `jpeg_destroy_decompress'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::PngDecoder<float, 2ul>::~PngDecoder()':
resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EED0Ev[_ZN3dip10PngDecoderIfLm2EED5Ev]+0x7a): undefined reference to `png_destroy_read_struct'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::JpegDecoder<float, 2ul>::readData(Img<float, 2ul>&)':
resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x91): undefined reference to `jpeg_destroy_decompress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x124): undefined reference to `jpeg_start_decompress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x1a1): undefined reference to `jpeg_read_scanlines'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip11JpegDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x28c): undefined reference to `jpeg_finish_decompress'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::JpegDecoder<float, 2ul>::readHeader()':
resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0x26): undefined reference to `jpeg_destroy_decompress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0x96): undefined reference to `jpeg_std_error'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0xf0): undefined reference to `jpeg_CreateDecompress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0x11d): undefined reference to `jpeg_stdio_src'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0x13e): undefined reference to `jpeg_save_markers'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0x14b): undefined reference to `jpeg_read_header'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0x16b): undefined reference to `jpeg_calc_output_dimensions'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegDecoderIfLm2EE10readHeaderEv[_ZN3dip11JpegDecoderIfLm2EE10readHeaderEv]+0x1b4): undefined reference to `jpeg_destroy_decompress'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::PngEncoder<float, 2ul>::write(Img<float, 2ul>&, std::vector<int, std::allocator<int> > const&)':
resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x39): undefined reference to `png_create_write_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0xbc): undefined reference to `png_create_info_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0xdc): undefined reference to `png_set_longjmp_fn'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0xfd): undefined reference to `png_destroy_write_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x16f): undefined reference to `png_init_io'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x18e): undefined reference to `png_set_filter'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x19d): undefined reference to `png_set_compression_level'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x1ac): undefined reference to `png_set_compression_strategy'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x1d6): undefined reference to `png_set_IHDR'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x1e9): undefined reference to `png_write_info'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x1f3): undefined reference to `png_set_bgr'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x1fd): undefined reference to `png_set_swap'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x60d): undefined reference to `png_write_image'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip10PngEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x61c): undefined reference to `png_write_end'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::PngDecoder<float, 2ul>::readData(Img<float, 2ul>&)':
resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0xe3): undefined reference to `png_set_longjmp_fn'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x115): undefined reference to `png_set_swap'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x11f): undefined reference to `png_set_strip_alpha'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x16b): undefined reference to `png_set_rgb_to_gray'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x173): undefined reference to `png_set_interlace_handling'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x180): undefined reference to `png_read_update_info'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x31f): undefined reference to `png_read_image'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x32c): undefined reference to `png_read_end'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x61d): undefined reference to `png_get_valid'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x63e): undefined reference to `png_get_eXIf_1'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x6ec): undefined reference to `png_destroy_read_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x745): undefined reference to `png_get_valid'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x766): undefined reference to `png_get_eXIf_1'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x770): undefined reference to `png_set_strip_16'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x77f): undefined reference to `png_set_expand_gray_1_2_4_to_8'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE[_ZN3dip10PngDecoderIfLm2EE8readDataER3ImgIfLm2EE]+0x78e): undefined reference to `png_set_palette_to_rgb'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::PngDecoder<float, 2ul>::readHeader()':
resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0x84): undefined reference to `png_destroy_read_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0xac): undefined reference to `png_create_read_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0xc1): undefined reference to `png_create_info_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0xd1): undefined reference to `png_create_info_struct'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0x11b): undefined reference to `png_set_longjmp_fn'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0x191): undefined reference to `png_init_io'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0x1b7): undefined reference to `png_read_info'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0x1e0): undefined reference to `png_get_IHDR'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0x251): undefined reference to `png_get_tRNS'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip10PngDecoderIfLm2EE10readHeaderEv[_ZN3dip10PngDecoderIfLm2EE10readHeaderEv]+0x2e3): undefined reference to `png_destroy_read_struct'
/usr/bin/ld: examples/DIPDialect/CMakeFiles/resize2D.dir/resize2D.cpp.o: in function `dip::JpegEncoder<float, 2ul>::write(Img<float, 2ul>&, std::vector<int, std::allocator<int> > const&)':
resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0xff): undefined reference to `jpeg_std_error'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x12d): undefined reference to `jpeg_CreateCompress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x257): undefined reference to `jpeg_destroy_compress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x2f4): undefined reference to `jpeg_set_defaults'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x316): undefined reference to `jpeg_set_quality'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x328): undefined reference to `jpeg_start_compress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x84e): undefined reference to `jpeg_write_scanlines'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x86f): undefined reference to `jpeg_finish_compress'
/usr/bin/ld: resize2D.cpp:(.text._ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE[_ZN3dip11JpegEncoderIfLm2EE5writeER3ImgIfLm2EERKSt6vectorIiSaIiEE]+0x8d8): undefined reference to `jpeg_stdio_dest'
collect2: error: ld returned 1 exit status
ninja: build stopped: subcommand failed.

Desktop (please complete the following information):
-- Target Triple: x86_64-unknown-linux-gnu

Remove upper and lower bound on BUDDY_DIP_OPT_STRIP_MINING

Is your feature request related to a problem? Please describe.
As of now, there is an upper and lower bound on stride size while using dip.corr_2d where upper bound is image width and lower bound is kernel width.

Describe the solution you'd like
I wish to eliminate the above mentioned bottleneck so that dip.corr_2d can be truly independent of any image/kernel.

Describe alternatives you've considered
I think we can develop support for a novel type of extrapolation (Refer this doc for context on type of extrapolation) which can fill corresponding values related to Upper/Mid/Lower left col, Upper/Mid/Lower mid col, Upper/Mid/Lower right col in a single if block.

Additional context
Above mentioned idea is just my personal opinion on how we can potentially solve this problem. Other solutions are highly encouraged and will be happily considered if they are successful in solving the above mentioned problem.

Implement separable convolution in DIP dialect for handling rank 1 kernels

Is your feature request related to a problem? Please describe.
A 2D rank 1 kernel can be divided in to two 1D kernels. Using the mathematical equation of convolution, we can conclude that two 1D convolutions with separated kernels is equivalent to one 2D convolution with original kernel for the same image. Theoretically, separable convolution might provide better performance boost for the rank 1 kernel case of 2D convolution.

Describe the solution you'd like
DIP dialect should have support for separable convolution. This can be provided in two ways :

  1. A separate API for separable convolution is created wherein the user provides both 1D kernels (Ref : sepFilter2D()).
  2. DIP dialect can detect whether a 2D kernel is rank 1 and then obtain separable kernels using Gaussian elimination. Separable convolution will then be applied with obtained separable kernels. Note : This method will have an overhead associated with Gaussian elimination.

Additional context
I have mentioned 2 possible ways of implementing the said feature, novel methods are encouraged and we are open for discussion.

Recommend Projects

  • React photo React

    A declarative, efficient, and flexible JavaScript library for building user interfaces.

  • Vue.js photo Vue.js

    ๐Ÿ–– Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.

  • Typescript photo Typescript

    TypeScript is a superset of JavaScript that compiles to clean JavaScript output.

  • TensorFlow photo TensorFlow

    An Open Source Machine Learning Framework for Everyone

  • Django photo Django

    The Web framework for perfectionists with deadlines.

  • D3 photo D3

    Bring data to life with SVG, Canvas and HTML. ๐Ÿ“Š๐Ÿ“ˆ๐ŸŽ‰

Recommend Topics

  • javascript

    JavaScript (JS) is a lightweight interpreted programming language with first-class functions.

  • web

    Some thing interesting about web. New door for the world.

  • server

    A server is a program made to process requests and deliver data to clients.

  • Machine learning

    Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.

  • Game

    Some thing interesting about game, make everyone happy.

Recommend Org

  • Facebook photo Facebook

    We are working to build community through open source technology. NB: members must have two-factor auth.

  • Microsoft photo Microsoft

    Open source projects and samples from Microsoft.

  • Google photo Google

    Google โค๏ธ Open Source for everyone.

  • D3 photo D3

    Data-Driven Documents codes.