diff --git a/Makefile b/Makefile index 8cb417b9dbebb0d76702e4bf51aab5485b90b8cb..4d8e3ae834fdf458b379e93751e1549383070205 100644 --- a/Makefile +++ b/Makefile @@ -1,6 +1,6 @@ BIN_DIR=bin -all: $(BIN_DIR)/matmul $(BIN_DIR)/pi +all: $(BIN_DIR)/matmul $(BIN_DIR)/pi $(BIN_DIR)/bilateral $(BIN_DIR): mkdir -p ./bin @@ -17,9 +17,16 @@ pi/out: $(BIN_DIR)/pi: $(BIN_DIR) pi/out cp ./pi/out $(BIN_DIR)/pi +bilateral/out: + cd ./bilateral/; make + +$(BIN_DIR)/bilateral: $(BIN_DIR) bilateral/out + cp ./bilateral/out $(BIN_DIR)/bilateral + clean: - rm -fv pi/out - rm -fv matmul/out + cd ./matmul/; make clean + cd ./pi; make clean + cd ./bilateral/; make clean rm -rfv $(BIN_DIR) diff --git a/README.md b/README.md index 175e53be595fd1b2e2b50f88cb46a0484cff85ef..6b12b8506d21763135ba70350969c7e3280437be 100644 --- a/README.md +++ b/README.md @@ -26,7 +26,7 @@ The source code of solution shown in [this file](./matmul/main.cu). # The PI number calculus -The problem is about how we can accelerate the approximation calculations of the Pi number using GPU using the Monte-Carlo method. +The problem is about how we can accelerate the approximation calculations of the Pi number on GPU using the Monte-Carlo method. The source code of solution shown in [this file](./pi/main.cu). @@ -39,4 +39,21 @@ The source code of solution shown in [this file](./pi/main.cu). | 300 | 0.045376 | 0.078464 | 0.578303 | | 1000 | 0.044672 | 0.097408 | 0.458607 | | 100000 | 0.807616 | 0.180928 | 4.46374 | -| 1000000 | 4.33581 | 1.43181 | 3.0282 | \ No newline at end of file +| 1000000 | 4.33581 | 1.43181 | 3.0282 | + + +# The bilateral image filtering + +The problem is about how we can accelerate the image filtering on GPU with bilateral filter. + +### Input parameters: +- img_path - path to source image_file +- sigma_r - smoothing parameter +- sigma_d - smoothing parameter + +### Calucatuion results +| Sigma r | Sigma d | CPU time | GPU time | Acceleration ratio | +|---------|---------|----------|----------|--------------------| +| 40 | 90 | 8.650950 | 0.022675 | 381.51066 | +| 75 | 75 | 9.288938 | 0.022062 | 421.02762 | +| 0.4 | 0.4 | 9.993022 | 0.022515 | 443.82321 | \ No newline at end of file diff --git a/bilateral/.gitignore b/bilateral/.gitignore new file mode 100644 index 0000000000000000000000000000000000000000..1dbd9e58e23f78bf5cf2e4f6f3a913d936c0838e --- /dev/null +++ b/bilateral/.gitignore @@ -0,0 +1,142 @@ +*.bmp +out_imgs/* +out +# Byte-compiled / optimized / DLL files +__pycache__/ +*.py[cod] +*$py.class + +# C extensions +*.so + +# Distribution / packaging +.Python +build/ +develop-eggs/ +dist/ +downloads/ +eggs/ +.eggs/ +lib/ +lib64/ +parts/ +sdist/ +var/ +wheels/ +share/python-wheels/ +*.egg-info/ +.installed.cfg +*.egg +MANIFEST + +# PyInstaller +# Usually these files are written by a python script from a template +# before PyInstaller builds the exe, so as to inject date/other infos into it. +*.manifest +*.spec + +# Installer logs +pip-log.txt +pip-delete-this-directory.txt + +# Unit test / coverage reports +htmlcov/ +.tox/ +.nox/ +.coverage +.coverage.* +.cache +nosetests.xml +coverage.xml +*.cover +*.py,cover +.hypothesis/ +.pytest_cache/ +cover/ + +# Translations +*.mo +*.pot + +# Django stuff: +*.log +local_settings.py +db.sqlite3 +db.sqlite3-journal + +# Flask stuff: +instance/ +.webassets-cache + +# Scrapy stuff: +.scrapy + +# Sphinx documentation +docs/_build/ + +# PyBuilder +.pybuilder/ +target/ + +# Jupyter Notebook +.ipynb_checkpoints + +# IPython +profile_default/ +ipython_config.py + +# pyenv +# For a library or package, you might want to ignore these files since the code is +# intended to run in multiple environments; otherwise, check them in: +# .python-version + +# pipenv +# According to pypa/pipenv#598, it is recommended to include Pipfile.lock in version control. +# However, in case of collaboration, if having platform-specific dependencies or dependencies +# having no cross-platform support, pipenv may install dependencies that don't work, or not +# install all needed dependencies. +#Pipfile.lock + +# PEP 582; used by e.g. github.com/David-OConnor/pyflow +__pypackages__/ + +# Celery stuff +celerybeat-schedule +celerybeat.pid + +# SageMath parsed files +*.sage.py + +# Environments +.env +.venv +env/ +venv/ +ENV/ +env.bak/ +venv.bak/ + +# Spyder project settings +.spyderproject +.spyproject + +# Rope project settings +.ropeproject + +# mkdocs documentation +/site + +# mypy +.mypy_cache/ +.dmypy.json +dmypy.json + +# Pyre type checker +.pyre/ + +# pytype static type analyzer +.pytype/ + +# Cython debug symbols +cython_debug/ + diff --git a/bilateral/Makefile b/bilateral/Makefile new file mode 100644 index 0000000000000000000000000000000000000000..d906516af98c96d4f229d226c0f9f601523cf407 --- /dev/null +++ b/bilateral/Makefile @@ -0,0 +1,26 @@ +EBMP_DIR=EasyBMP + +all: out post_install + +poetry: + pip install --user poetry + +out: + poetry run pyinstaller -F \ + --clean \ + --add-data 'src/kernel.cu:src' \ + --additional-hooks-dir='./hooks/' \ + --exclude-module='FixTk' \ + --exclude-module='tcl' \ + --exclude-module='tk' \ + --exclude-module='_tkinter' \ + --exclude-module='tkinter' \ + --exclude-module='Tkinter' \ + main.py + cp ./dist/main out + +post_install: + rm -rf ./__pycache__/ ./build/ ./dist/ main.spec + +clean: + rm -rf out out_imgs/ diff --git a/bilateral/hooks/hook-pycuda.py b/bilateral/hooks/hook-pycuda.py new file mode 100644 index 0000000000000000000000000000000000000000..bf41b49344c444e2bd3136d8ed4bb9322f33b034 --- /dev/null +++ b/bilateral/hooks/hook-pycuda.py @@ -0,0 +1,2 @@ +from PyInstaller.utils.hooks import copy_metadata +datas = copy_metadata('pycuda') diff --git a/bilateral/main.py b/bilateral/main.py new file mode 100644 index 0000000000000000000000000000000000000000..fb2c56272b8e7761ef7bc9302e6f6c894d11d7fe --- /dev/null +++ b/bilateral/main.py @@ -0,0 +1,62 @@ +import numpy as np +from PIL import Image +from src import gpu, cpu +import fire +import logging +from pathlib import Path +logger = logging.getLogger(__name__) + + +def read_img(file_path: str) -> np.ndarray: + logger.debug(f'Reading image {file_path}') + src_img = Image.open(file_path).convert('L') + return np.asarray(src_img) + + +def write_img(img_arr: np.ndarray, img_name: str): + Path('out_imgs').mkdir(exist_ok=True) + res_img = Image.fromarray(img_arr).convert('L') + filename = f'out_imgs/{img_name}.png' + logger.debug(f'Writing image to {filename}') + res_img.save(filename) + + +class CLI(object): + def run(self, + img_path: str, + sigma_r: float = 75, + sigma_d: float = 75): + gpu_time = self.run_gpu(img_path, sigma_r, sigma_d) + cpu_time = self.run_cpu(img_path, sigma_r, sigma_d) + logger.info(f'CPU execution time: {cpu_time}') + logger.info(f'GPU execution time: {gpu_time}') + logger.info(f'Acceleration: {cpu_time / gpu_time:.5f}') + + def run_gpu(self, + img_path: str, + sigma_r: float = 75, + sigma_d: float = 75): + src_img = read_img(img_path) + res_img, execution_time = gpu.filter_image(src_img, sigma_r, sigma_d) + write_img(res_img, 'gpu') + return execution_time + + def run_cpu(self, + img_path: str, + sigma_r: float = 75, + sigma_d: float = 75): + src_img = read_img(img_path) + res_img, execution_time = cpu.filter_image(src_img, sigma_r, sigma_d) + write_img(res_img, 'cpu') + return execution_time + + +def setup_logging(): + logging.basicConfig( + format="[%(levelname)8s] %(msg)s {%(filename)s:%(lineno)s}", + level=logging.DEBUG) + + +if __name__ == '__main__': + setup_logging() + fire.Fire(CLI) diff --git a/bilateral/poetry.lock b/bilateral/poetry.lock new file mode 100644 index 0000000000000000000000000000000000000000..2a642461844679b62603294a0b1fbadacde95e12 --- /dev/null +++ b/bilateral/poetry.lock @@ -0,0 +1,276 @@ +[[package]] +category = "main" +description = "Python graph (network) package" +name = "altgraph" +optional = false +python-versions = "*" +version = "0.17" + +[[package]] +category = "main" +description = "A small Python module for determining appropriate platform-specific dirs, e.g. a \"user data dir\"." +name = "appdirs" +optional = false +python-versions = "*" +version = "1.4.4" + +[[package]] +category = "main" +description = "Decorators for Humans" +name = "decorator" +optional = false +python-versions = ">=2.6, !=3.0.*, !=3.1.*" +version = "4.4.2" + +[[package]] +category = "main" +description = "Python 2.7 backport of the \"dis\" module from Python 3.5+" +name = "dis3" +optional = false +python-versions = "*" +version = "0.1.3" + +[[package]] +category = "main" +description = "A library for automatically generating command line interfaces." +name = "fire" +optional = false +python-versions = "*" +version = "0.3.1" + +[package.dependencies] +six = "*" +termcolor = "*" + +[[package]] +category = "main" +description = "A super-fast templating language that borrows the best ideas from the existing templating languages." +name = "mako" +optional = false +python-versions = ">=2.7, !=3.0.*, !=3.1.*, !=3.2.*, !=3.3.*" +version = "1.1.3" + +[package.dependencies] +MarkupSafe = ">=0.9.2" + +[package.extras] +babel = ["babel"] +lingua = ["lingua"] + +[[package]] +category = "main" +description = "Safely add untrusted strings to HTML/XML markup." +name = "markupsafe" +optional = false +python-versions = ">=2.7,!=3.0.*,!=3.1.*,!=3.2.*,!=3.3.*" +version = "1.1.1" + +[[package]] +category = "main" +description = "NumPy is the fundamental package for array computing with Python." +name = "numpy" +optional = false +python-versions = ">=3.6" +version = "1.19.0" + +[[package]] +category = "main" +description = "Python Imaging Library (Fork)" +name = "pillow" +optional = false +python-versions = ">=3.5" +version = "7.2.0" + +[[package]] +category = "main" +description = "Python wrapper for Nvidia CUDA" +name = "pycuda" +optional = false +python-versions = "*" +version = "2019.1.2" + +[package.dependencies] +appdirs = ">=1.4.0" +decorator = ">=3.2.0" +mako = "*" +pytools = ">=2011.2" + +[[package]] +category = "main" +description = "PyInstaller bundles a Python application and all its dependencies into a single package." +name = "pyinstaller" +optional = false +python-versions = ">=2.7, !=3.0.*, !=3.1.*, !=3.2.*, !=3.3.*, !=3.4.*" +version = "3.6" + +[package.dependencies] +altgraph = "*" +dis3 = "*" +setuptools = "*" + +[[package]] +category = "main" +description = "A collection of tools for Python" +name = "pytools" +optional = false +python-versions = "~=3.6" +version = "2020.3.1" + +[package.dependencies] +appdirs = ">=1.4.0" +decorator = ">=3.2.0" +numpy = ">=1.6.0" +six = ">=1.8.0" + +[[package]] +category = "main" +description = "Python 2 and 3 compatibility utilities" +name = "six" +optional = false +python-versions = ">=2.7, !=3.0.*, !=3.1.*, !=3.2.*" +version = "1.15.0" + +[[package]] +category = "main" +description = "ANSII Color formatting for output in terminal." +name = "termcolor" +optional = false +python-versions = "*" +version = "1.1.0" + +[metadata] +content-hash = "df15e4cc962947ea0257fd079fe5702ea74be601b1f1afdcce9598c491df4018" +python-versions = "^3.8" + +[metadata.files] +altgraph = [ + {file = "altgraph-0.17-py2.py3-none-any.whl", hash = "sha256:c623e5f3408ca61d4016f23a681b9adb100802ca3e3da5e718915a9e4052cebe"}, + {file = "altgraph-0.17.tar.gz", hash = "sha256:1f05a47122542f97028caf78775a095fbe6a2699b5089de8477eb583167d69aa"}, +] +appdirs = [ + {file = "appdirs-1.4.4-py2.py3-none-any.whl", hash = "sha256:a841dacd6b99318a741b166adb07e19ee71a274450e68237b4650ca1055ab128"}, + {file = "appdirs-1.4.4.tar.gz", hash = "sha256:7d5d0167b2b1ba821647616af46a749d1c653740dd0d2415100fe26e27afdf41"}, +] +decorator = [ + {file = "decorator-4.4.2-py2.py3-none-any.whl", hash = "sha256:41fa54c2a0cc4ba648be4fd43cff00aedf5b9465c9bf18d64325bc225f08f760"}, + {file = "decorator-4.4.2.tar.gz", hash = "sha256:e3a62f0520172440ca0dcc823749319382e377f37f140a0b99ef45fecb84bfe7"}, +] +dis3 = [ + {file = "dis3-0.1.3-py2-none-any.whl", hash = "sha256:61f7720dd0d8749d23fda3d7227ce74d73da11c2fade993a67ab2f9852451b14"}, + {file = "dis3-0.1.3-py3-none-any.whl", hash = "sha256:30b6412d33d738663e8ded781b138f4b01116437f0872aa56aa3adba6aeff218"}, + {file = "dis3-0.1.3.tar.gz", hash = "sha256:9259b881fc1df02ed12ac25f82d4a85b44241854330b1a651e40e0c675cb2d1e"}, +] +fire = [ + {file = "fire-0.3.1.tar.gz", hash = "sha256:9736a16227c3d469e5d2d296bce5b4d8fa8d7851e953bda327a455fc2994307f"}, +] +mako = [ + {file = "Mako-1.1.3-py2.py3-none-any.whl", hash = "sha256:93729a258e4ff0747c876bd9e20df1b9758028946e976324ccd2d68245c7b6a9"}, + {file = "Mako-1.1.3.tar.gz", hash = "sha256:8195c8c1400ceb53496064314c6736719c6f25e7479cd24c77be3d9361cddc27"}, +] +markupsafe = [ + {file = "MarkupSafe-1.1.1-cp27-cp27m-macosx_10_6_intel.whl", hash = "sha256:09027a7803a62ca78792ad89403b1b7a73a01c8cb65909cd876f7fcebd79b161"}, + {file = "MarkupSafe-1.1.1-cp27-cp27m-manylinux1_i686.whl", hash = "sha256:e249096428b3ae81b08327a63a485ad0878de3fb939049038579ac0ef61e17e7"}, + {file = "MarkupSafe-1.1.1-cp27-cp27m-manylinux1_x86_64.whl", hash = "sha256:500d4957e52ddc3351cabf489e79c91c17f6e0899158447047588650b5e69183"}, + {file = "MarkupSafe-1.1.1-cp27-cp27m-win32.whl", hash = "sha256:b2051432115498d3562c084a49bba65d97cf251f5a331c64a12ee7e04dacc51b"}, + {file = "MarkupSafe-1.1.1-cp27-cp27m-win_amd64.whl", hash = "sha256:98c7086708b163d425c67c7a91bad6e466bb99d797aa64f965e9d25c12111a5e"}, + {file = "MarkupSafe-1.1.1-cp27-cp27mu-manylinux1_i686.whl", hash = "sha256:cd5df75523866410809ca100dc9681e301e3c27567cf498077e8551b6d20e42f"}, + {file = "MarkupSafe-1.1.1-cp27-cp27mu-manylinux1_x86_64.whl", hash = "sha256:43a55c2930bbc139570ac2452adf3d70cdbb3cfe5912c71cdce1c2c6bbd9c5d1"}, + {file = "MarkupSafe-1.1.1-cp34-cp34m-macosx_10_6_intel.whl", hash = "sha256:1027c282dad077d0bae18be6794e6b6b8c91d58ed8a8d89a89d59693b9131db5"}, + {file = "MarkupSafe-1.1.1-cp34-cp34m-manylinux1_i686.whl", hash = "sha256:62fe6c95e3ec8a7fad637b7f3d372c15ec1caa01ab47926cfdf7a75b40e0eac1"}, + {file = "MarkupSafe-1.1.1-cp34-cp34m-manylinux1_x86_64.whl", hash = "sha256:88e5fcfb52ee7b911e8bb6d6aa2fd21fbecc674eadd44118a9cc3863f938e735"}, + {file = "MarkupSafe-1.1.1-cp34-cp34m-win32.whl", hash = "sha256:ade5e387d2ad0d7ebf59146cc00c8044acbd863725f887353a10df825fc8ae21"}, + {file = "MarkupSafe-1.1.1-cp34-cp34m-win_amd64.whl", hash = "sha256:09c4b7f37d6c648cb13f9230d847adf22f8171b1ccc4d5682398e77f40309235"}, + {file = "MarkupSafe-1.1.1-cp35-cp35m-macosx_10_6_intel.whl", hash = "sha256:79855e1c5b8da654cf486b830bd42c06e8780cea587384cf6545b7d9ac013a0b"}, + {file = "MarkupSafe-1.1.1-cp35-cp35m-manylinux1_i686.whl", hash = "sha256:c8716a48d94b06bb3b2524c2b77e055fb313aeb4ea620c8dd03a105574ba704f"}, + {file = "MarkupSafe-1.1.1-cp35-cp35m-manylinux1_x86_64.whl", hash = "sha256:7c1699dfe0cf8ff607dbdcc1e9b9af1755371f92a68f706051cc8c37d447c905"}, + {file = "MarkupSafe-1.1.1-cp35-cp35m-win32.whl", hash = "sha256:6dd73240d2af64df90aa7c4e7481e23825ea70af4b4922f8ede5b9e35f78a3b1"}, + {file = "MarkupSafe-1.1.1-cp35-cp35m-win_amd64.whl", hash = "sha256:9add70b36c5666a2ed02b43b335fe19002ee5235efd4b8a89bfcf9005bebac0d"}, + {file = "MarkupSafe-1.1.1-cp36-cp36m-macosx_10_6_intel.whl", hash = "sha256:24982cc2533820871eba85ba648cd53d8623687ff11cbb805be4ff7b4c971aff"}, + {file = "MarkupSafe-1.1.1-cp36-cp36m-manylinux1_i686.whl", hash = "sha256:00bc623926325b26bb9605ae9eae8a215691f33cae5df11ca5424f06f2d1f473"}, + {file = "MarkupSafe-1.1.1-cp36-cp36m-manylinux1_x86_64.whl", hash = "sha256:717ba8fe3ae9cc0006d7c451f0bb265ee07739daf76355d06366154ee68d221e"}, + {file = "MarkupSafe-1.1.1-cp36-cp36m-win32.whl", hash = "sha256:535f6fc4d397c1563d08b88e485c3496cf5784e927af890fb3c3aac7f933ec66"}, + {file = "MarkupSafe-1.1.1-cp36-cp36m-win_amd64.whl", hash = "sha256:b1282f8c00509d99fef04d8ba936b156d419be841854fe901d8ae224c59f0be5"}, + {file = "MarkupSafe-1.1.1-cp37-cp37m-macosx_10_6_intel.whl", hash = "sha256:8defac2f2ccd6805ebf65f5eeb132adcf2ab57aa11fdf4c0dd5169a004710e7d"}, + {file = "MarkupSafe-1.1.1-cp37-cp37m-manylinux1_i686.whl", hash = "sha256:46c99d2de99945ec5cb54f23c8cd5689f6d7177305ebff350a58ce5f8de1669e"}, + {file = "MarkupSafe-1.1.1-cp37-cp37m-manylinux1_x86_64.whl", hash = "sha256:ba59edeaa2fc6114428f1637ffff42da1e311e29382d81b339c1817d37ec93c6"}, + {file = "MarkupSafe-1.1.1-cp37-cp37m-win32.whl", hash = "sha256:b00c1de48212e4cc9603895652c5c410df699856a2853135b3967591e4beebc2"}, + {file = "MarkupSafe-1.1.1-cp37-cp37m-win_amd64.whl", hash = "sha256:9bf40443012702a1d2070043cb6291650a0841ece432556f784f004937f0f32c"}, + {file = "MarkupSafe-1.1.1-cp38-cp38-macosx_10_9_x86_64.whl", hash = "sha256:6788b695d50a51edb699cb55e35487e430fa21f1ed838122d722e0ff0ac5ba15"}, + {file = "MarkupSafe-1.1.1-cp38-cp38-manylinux1_i686.whl", hash = "sha256:cdb132fc825c38e1aeec2c8aa9338310d29d337bebbd7baa06889d09a60a1fa2"}, + {file = "MarkupSafe-1.1.1-cp38-cp38-manylinux1_x86_64.whl", hash = "sha256:13d3144e1e340870b25e7b10b98d779608c02016d5184cfb9927a9f10c689f42"}, + {file = "MarkupSafe-1.1.1-cp38-cp38-win32.whl", hash = "sha256:596510de112c685489095da617b5bcbbac7dd6384aeebeda4df6025d0256a81b"}, + {file = "MarkupSafe-1.1.1-cp38-cp38-win_amd64.whl", hash = "sha256:e8313f01ba26fbbe36c7be1966a7b7424942f670f38e666995b88d012765b9be"}, + {file = "MarkupSafe-1.1.1.tar.gz", hash = "sha256:29872e92839765e546828bb7754a68c418d927cd064fd4708fab9fe9c8bb116b"}, +] +numpy = [ + {file = "numpy-1.19.0-cp36-cp36m-macosx_10_9_x86_64.whl", hash = "sha256:63d971bb211ad3ca37b2adecdd5365f40f3b741a455beecba70fd0dde8b2a4cb"}, + {file = "numpy-1.19.0-cp36-cp36m-manylinux1_i686.whl", hash = "sha256:b6aaeadf1e4866ca0fdf7bb4eed25e521ae21a7947c59f78154b24fc7abbe1dd"}, + {file = "numpy-1.19.0-cp36-cp36m-manylinux1_x86_64.whl", hash = "sha256:13af0184177469192d80db9bd02619f6fa8b922f9f327e077d6f2a6acb1ce1c0"}, + {file = "numpy-1.19.0-cp36-cp36m-manylinux2010_i686.whl", hash = "sha256:356f96c9fbec59974a592452ab6a036cd6f180822a60b529a975c9467fcd5f23"}, + {file = "numpy-1.19.0-cp36-cp36m-manylinux2010_x86_64.whl", hash = "sha256:fa1fe75b4a9e18b66ae7f0b122543c42debcf800aaafa0212aaff3ad273c2596"}, + {file = "numpy-1.19.0-cp36-cp36m-manylinux2014_aarch64.whl", hash = "sha256:cbe326f6d364375a8e5a8ccb7e9cd73f4b2f6dc3b2ed205633a0db8243e2a96a"}, + {file = "numpy-1.19.0-cp36-cp36m-win32.whl", hash = "sha256:a2e3a39f43f0ce95204beb8fe0831199542ccab1e0c6e486a0b4947256215632"}, + {file = "numpy-1.19.0-cp36-cp36m-win_amd64.whl", hash = "sha256:7b852817800eb02e109ae4a9cef2beda8dd50d98b76b6cfb7b5c0099d27b52d4"}, + {file = "numpy-1.19.0-cp37-cp37m-macosx_10_9_x86_64.whl", hash = "sha256:d97a86937cf9970453c3b62abb55a6475f173347b4cde7f8dcdb48c8e1b9952d"}, + {file = "numpy-1.19.0-cp37-cp37m-manylinux1_i686.whl", hash = "sha256:a86c962e211f37edd61d6e11bb4df7eddc4a519a38a856e20a6498c319efa6b0"}, + {file = "numpy-1.19.0-cp37-cp37m-manylinux1_x86_64.whl", hash = "sha256:d34fbb98ad0d6b563b95de852a284074514331e6b9da0a9fc894fb1cdae7a79e"}, + {file = "numpy-1.19.0-cp37-cp37m-manylinux2010_i686.whl", hash = "sha256:658624a11f6e1c252b2cd170d94bf28c8f9410acab9f2fd4369e11e1cd4e1aaf"}, + {file = "numpy-1.19.0-cp37-cp37m-manylinux2010_x86_64.whl", hash = "sha256:4d054f013a1983551254e2379385e359884e5af105e3efe00418977d02f634a7"}, + {file = "numpy-1.19.0-cp37-cp37m-manylinux2014_aarch64.whl", hash = "sha256:26a45798ca2a4e168d00de75d4a524abf5907949231512f372b217ede3429e98"}, + {file = "numpy-1.19.0-cp37-cp37m-win32.whl", hash = "sha256:3c40c827d36c6d1c3cf413694d7dc843d50997ebffbc7c87d888a203ed6403a7"}, + {file = "numpy-1.19.0-cp37-cp37m-win_amd64.whl", hash = "sha256:be62aeff8f2f054eff7725f502f6228298891fd648dc2630e03e44bf63e8cee0"}, + {file = "numpy-1.19.0-cp38-cp38-macosx_10_9_x86_64.whl", hash = "sha256:dd53d7c4a69e766e4900f29db5872f5824a06827d594427cf1a4aa542818b796"}, + {file = "numpy-1.19.0-cp38-cp38-manylinux1_i686.whl", hash = "sha256:30a59fb41bb6b8c465ab50d60a1b298d1cd7b85274e71f38af5a75d6c475d2d2"}, + {file = "numpy-1.19.0-cp38-cp38-manylinux1_x86_64.whl", hash = "sha256:df1889701e2dfd8ba4dc9b1a010f0a60950077fb5242bb92c8b5c7f1a6f2668a"}, + {file = "numpy-1.19.0-cp38-cp38-manylinux2010_i686.whl", hash = "sha256:33c623ef9ca5e19e05991f127c1be5aeb1ab5cdf30cb1c5cf3960752e58b599b"}, + {file = "numpy-1.19.0-cp38-cp38-manylinux2010_x86_64.whl", hash = "sha256:26f509450db547e4dfa3ec739419b31edad646d21fb8d0ed0734188b35ff6b27"}, + {file = "numpy-1.19.0-cp38-cp38-manylinux2014_aarch64.whl", hash = "sha256:7b57f26e5e6ee2f14f960db46bd58ffdca25ca06dd997729b1b179fddd35f5a3"}, + {file = "numpy-1.19.0-cp38-cp38-win32.whl", hash = "sha256:a8705c5073fe3fcc297fb8e0b31aa794e05af6a329e81b7ca4ffecab7f2b95ef"}, + {file = "numpy-1.19.0-cp38-cp38-win_amd64.whl", hash = "sha256:c2edbb783c841e36ca0fa159f0ae97a88ce8137fb3a6cd82eae77349ba4b607b"}, + {file = "numpy-1.19.0-pp36-pypy36_pp73-manylinux2010_x86_64.whl", hash = "sha256:8cde829f14bd38f6da7b2954be0f2837043e8b8d7a9110ec5e318ae6bf706610"}, + {file = "numpy-1.19.0.zip", hash = "sha256:76766cc80d6128750075378d3bb7812cf146415bd29b588616f72c943c00d598"}, +] +pillow = [ + {file = "Pillow-7.2.0-cp35-cp35m-macosx_10_10_intel.whl", hash = "sha256:1ca594126d3c4def54babee699c055a913efb01e106c309fa6b04405d474d5ae"}, + {file = "Pillow-7.2.0-cp35-cp35m-manylinux1_i686.whl", hash = "sha256:c92302a33138409e8f1ad16731568c55c9053eee71bb05b6b744067e1b62380f"}, + {file = "Pillow-7.2.0-cp35-cp35m-manylinux1_x86_64.whl", hash = "sha256:8dad18b69f710bf3a001d2bf3afab7c432785d94fcf819c16b5207b1cfd17d38"}, + {file = "Pillow-7.2.0-cp35-cp35m-manylinux2014_aarch64.whl", hash = "sha256:431b15cffbf949e89df2f7b48528be18b78bfa5177cb3036284a5508159492b5"}, + {file = "Pillow-7.2.0-cp35-cp35m-win32.whl", hash = "sha256:09d7f9e64289cb40c2c8d7ad674b2ed6105f55dc3b09aa8e4918e20a0311e7ad"}, + {file = "Pillow-7.2.0-cp35-cp35m-win_amd64.whl", hash = "sha256:0295442429645fa16d05bd567ef5cff178482439c9aad0411d3f0ce9b88b3a6f"}, + {file = "Pillow-7.2.0-cp36-cp36m-macosx_10_10_x86_64.whl", hash = "sha256:ec29604081f10f16a7aea809ad42e27764188fc258b02259a03a8ff7ded3808d"}, + {file = "Pillow-7.2.0-cp36-cp36m-manylinux1_i686.whl", hash = "sha256:612cfda94e9c8346f239bf1a4b082fdd5c8143cf82d685ba2dba76e7adeeb233"}, + {file = "Pillow-7.2.0-cp36-cp36m-manylinux1_x86_64.whl", hash = "sha256:0a80dd307a5d8440b0a08bd7b81617e04d870e40a3e46a32d9c246e54705e86f"}, + {file = "Pillow-7.2.0-cp36-cp36m-manylinux2014_aarch64.whl", hash = "sha256:06aba4169e78c439d528fdeb34762c3b61a70813527a2c57f0540541e9f433a8"}, + {file = "Pillow-7.2.0-cp36-cp36m-win32.whl", hash = "sha256:f7e30c27477dffc3e85c2463b3e649f751789e0f6c8456099eea7ddd53be4a8a"}, + {file = "Pillow-7.2.0-cp36-cp36m-win_amd64.whl", hash = "sha256:ffe538682dc19cc542ae7c3e504fdf54ca7f86fb8a135e59dd6bc8627eae6cce"}, + {file = "Pillow-7.2.0-cp37-cp37m-macosx_10_10_x86_64.whl", hash = "sha256:94cf49723928eb6070a892cb39d6c156f7b5a2db4e8971cb958f7b6b104fb4c4"}, + {file = "Pillow-7.2.0-cp37-cp37m-manylinux1_i686.whl", hash = "sha256:6edb5446f44d901e8683ffb25ebdfc26988ee813da3bf91e12252b57ac163727"}, + {file = "Pillow-7.2.0-cp37-cp37m-manylinux1_x86_64.whl", hash = "sha256:52125833b070791fcb5710fabc640fc1df07d087fc0c0f02d3661f76c23c5b8b"}, + {file = "Pillow-7.2.0-cp37-cp37m-manylinux2014_aarch64.whl", hash = "sha256:9ad7f865eebde135d526bb3163d0b23ffff365cf87e767c649550964ad72785d"}, + {file = "Pillow-7.2.0-cp37-cp37m-win32.whl", hash = "sha256:c79f9c5fb846285f943aafeafda3358992d64f0ef58566e23484132ecd8d7d63"}, + {file = "Pillow-7.2.0-cp37-cp37m-win_amd64.whl", hash = "sha256:d350f0f2c2421e65fbc62690f26b59b0bcda1b614beb318c81e38647e0f673a1"}, + {file = "Pillow-7.2.0-cp38-cp38-macosx_10_10_x86_64.whl", hash = "sha256:6d7741e65835716ceea0fd13a7d0192961212fd59e741a46bbed7a473c634ed6"}, + {file = "Pillow-7.2.0-cp38-cp38-manylinux1_i686.whl", hash = "sha256:edf31f1150778abd4322444c393ab9c7bd2af271dd4dafb4208fb613b1f3cdc9"}, + {file = "Pillow-7.2.0-cp38-cp38-manylinux1_x86_64.whl", hash = "sha256:d08b23fdb388c0715990cbc06866db554e1822c4bdcf6d4166cf30ac82df8c41"}, + {file = "Pillow-7.2.0-cp38-cp38-manylinux2014_aarch64.whl", hash = "sha256:5e51ee2b8114def244384eda1c82b10e307ad9778dac5c83fb0943775a653cd8"}, + {file = "Pillow-7.2.0-cp38-cp38-win32.whl", hash = "sha256:725aa6cfc66ce2857d585f06e9519a1cc0ef6d13f186ff3447ab6dff0a09bc7f"}, + {file = "Pillow-7.2.0-cp38-cp38-win_amd64.whl", hash = "sha256:a060cf8aa332052df2158e5a119303965be92c3da6f2d93b6878f0ebca80b2f6"}, + {file = "Pillow-7.2.0-pp36-pypy36_pp73-win32.whl", hash = "sha256:25930fadde8019f374400f7986e8404c8b781ce519da27792cbe46eabec00c4d"}, + {file = "Pillow-7.2.0.tar.gz", hash = "sha256:97f9e7953a77d5a70f49b9a48da7776dc51e9b738151b22dacf101641594a626"}, +] +pycuda = [ + {file = "pycuda-2019.1.2.tar.gz", hash = "sha256:ada56ce98a41f9f95fe18809f38afbae473a5c62d346cfa126a2d5477f24cc8a"}, +] +pyinstaller = [ + {file = "PyInstaller-3.6.tar.gz", hash = "sha256:3730fa80d088f8bb7084d32480eb87cbb4ddb64123363763cf8f2a1378c1c4b7"}, +] +pytools = [ + {file = "pytools-2020.3.1.tar.gz", hash = "sha256:86ebb27e8d946b30bc4479f97862066eb26e305d5ad4327230b2b2f8cbf110f9"}, +] +six = [ + {file = "six-1.15.0-py2.py3-none-any.whl", hash = "sha256:8b74bedcbbbaca38ff6d7491d76f2b06b3592611af620f8426e82dddb04a5ced"}, + {file = "six-1.15.0.tar.gz", hash = "sha256:30639c035cdb23534cd4aa2dd52c3bf48f06e5f4a941509c8bafd8ce11080259"}, +] +termcolor = [ + {file = "termcolor-1.1.0.tar.gz", hash = "sha256:1d6d69ce66211143803fbc56652b41d73b4a400a2891d7bf7a1cdf4c02de613b"}, +] diff --git a/bilateral/pyproject.toml b/bilateral/pyproject.toml new file mode 100644 index 0000000000000000000000000000000000000000..9a374369db4d29cc007a86b65d2f423d41d173e4 --- /dev/null +++ b/bilateral/pyproject.toml @@ -0,0 +1,19 @@ +[tool.poetry] +name = "bilateral" +version = "0.1.0" +description = "" +authors = ["Pavel Kirilin <win10@list.ru>"] + +[tool.poetry.dependencies] +python = "^3.8" +numpy = "^1.19.0" +pyinstaller = "^3.6" +pycuda = "^2019.1.2" +fire = "^0.3.1" +pillow = "^7.2.0" + +[tool.poetry.dev-dependencies] + +[build-system] +requires = ["poetry>=0.12"] +build-backend = "poetry.masonry.api" diff --git a/bilateral/src/__init__.py b/bilateral/src/__init__.py new file mode 100644 index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391 diff --git a/bilateral/src/cpu.py b/bilateral/src/cpu.py new file mode 100644 index 0000000000000000000000000000000000000000..8b2065588bdce7b0ef87efdfae25bc8ff8decd68 --- /dev/null +++ b/bilateral/src/cpu.py @@ -0,0 +1,25 @@ +import numpy as np +import time +import logging + +logger = logging.getLogger(__name__) + + +def filter_image(image: np.ndarray, sigma_r: float, sigma_d: float): + result = np.zeros(image.shape) + logger.debug("Started CPU filtering") + start_time = time.time() + for i in range(1, image.shape[0] - 1): + for j in range(1, image.shape[1] - 1): + c = 0 + s = 0 + for k in range(i - 1, i + 2): + for l in range(j - 1, j + 2): + g = np.exp(-((k - i)**2 + (l - j)**2) / sigma_d**2) + r = np.exp(-(image[k, l] - image[i, j])**2 / sigma_r**2) + c += g * r + s += g * r * image[k, l] + result[i, j] = s / c + result_time = time.time() - start_time + logger.debug(f"CPU execution time: {result_time} ms") + return result, result_time diff --git a/bilateral/src/gpu.py b/bilateral/src/gpu.py new file mode 100644 index 0000000000000000000000000000000000000000..b183a2cd7370b17efee06ec6a743633f92016648 --- /dev/null +++ b/bilateral/src/gpu.py @@ -0,0 +1,48 @@ +#!/usr/bin/env python +# -*- coding: utf-8 -*- +import time +import numpy as np +import pkgutil +import pycuda.autoinit +import pycuda.driver as drv +from pycuda.compiler import SourceModule + +import logging + +logger = logging.getLogger(__name__) + + +def filter_image(image: np.ndarray, sigma_r: float, sigma_d: float): + kernel_code = pkgutil.get_data('src', + 'kernel.cu').decode('UTF-8', 'ignore') + logger.debug('Read kernel code.') + logger.debug(kernel_code) + kernel = SourceModule(kernel_code) + N, M = image.shape + logger.debug("Image shape: {N}, {M}") + result = np.zeros((N, M), dtype=np.uint32) + block_size = (8, 8, 1) + logger.debug(f'Using block size: {block_size}') + grid_size = (int(np.ceil(N / block_size[0])), + int(np.ceil(M / block_size[1]))) + logger.debug(f'Using block size: {block_size}') + filt_gpu = kernel.get_function("filter") + tex = kernel.get_texref("tex") + tex.set_filter_mode(drv.filter_mode.LINEAR) + tex.set_address_mode(0, drv.address_mode.MIRROR) + tex.set_address_mode(1, drv.address_mode.MIRROR) + drv.matrix_to_texref(image.astype(np.uint32), tex, order="C") + start_time = time.time() + filt_gpu(drv.Out(result), + np.int32(N), + np.int32(M), + np.float32(sigma_d), + np.float32(sigma_r), + block=block_size, + grid=grid_size, + texrefs=[tex]) + drv.Context.synchronize() + end_time = time.time() + execution_time = end_time - start_time + logging.info(f"Execution time: {execution_time}") + return result, execution_time diff --git a/bilateral/src/kernel.cu b/bilateral/src/kernel.cu new file mode 100644 index 0000000000000000000000000000000000000000..19217a5108a88a55d3eb81b204244c526d9654aa --- /dev/null +++ b/bilateral/src/kernel.cu @@ -0,0 +1,21 @@ +texture<unsigned int, 2, cudaReadModeElementType> tex; +__global__ void filter(unsigned int* result, const int M, const int N, const float sigma_d, const float sigma_r) +{ + const int i = threadIdx.x + blockDim.x * blockIdx.x; + const int j = threadIdx.y + blockDim.y * blockIdx.y; + if ((i < M) && (j < N)) { + float s = 0; + float c = 0; + for (int l = i - 1; l <= i + 1; l++){ + for (int k = j - 1; k <= j + 1; k++){ + float img1 = tex2D(tex, k, l) / 255; + float img2 = tex2D(tex, i, j) / 255; + float g = exp(-(pow(k - i, 2) + pow(l - j, 2)) / pow(sigma_d, 2)); + float r = exp(-pow((img1 - img2) * 255, 2) / pow(sigma_r, 2)); + c += g * r; + s += g * r * tex2D(tex, k, l); + } + } + result[i * N + j] = s / c; + } +} diff --git a/bilateral/test_image.jpg b/bilateral/test_image.jpg new file mode 100644 index 0000000000000000000000000000000000000000..0f93a62aeba189bcfe93c41f461526c5f613f65a Binary files /dev/null and b/bilateral/test_image.jpg differ diff --git a/matmul/Makefile b/matmul/Makefile index 8485037b8b587c2b3a8df779dc39225f6b877d2f..2bb2ef87379058a187a2ae339514941416e2b943 100644 --- a/matmul/Makefile +++ b/matmul/Makefile @@ -2,11 +2,11 @@ OUT_FILE_NAME=out all: $(OUT_FILE_NAME) -out: main.cu - nvcc -O3 main.cu -o $(OUT_FILE_NAME) +$(OUT_FILE_NAME): main.cu + nvcc -O3 main.cu -o $@ clean: - rm $(OUT_FILE_NAME) + rm -f $(OUT_FILE_NAME) run: $(OUT_FILE_NAME) ./$(OUT_FILE_NAME) diff --git a/pi/Makefile b/pi/Makefile index bd58546b1adbb29117a4c3fac1e2c2c3e3aea677..92aa6c5248e8db2a3ccd04bd4a86b442847b336f 100644 --- a/pi/Makefile +++ b/pi/Makefile @@ -2,8 +2,8 @@ OUT_FILE_NAME=out all: $(OUT_FILE_NAME) -out: main.cu - nvcc -O2 -o $(OUT_FILE_NAME) main.cu -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA +$(OUT_FILE_NAME): main.cu + nvcc -O2 -o $@ main.cu -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CUDA clean: rm $(OUT_FILE_NAME)