diff --git a/.gitignore b/.gitignore index 259148f..77ee40e 100644 --- a/.gitignore +++ b/.gitignore @@ -30,3 +30,7 @@ *.exe *.out *.app + +# Directories +build/ +.vscode/ \ No newline at end of file diff --git a/rolling-mean-filter/CMakeLists.txt b/rolling-mean-filter/CMakeLists.txt new file mode 100644 index 0000000..9991263 --- /dev/null +++ b/rolling-mean-filter/CMakeLists.txt @@ -0,0 +1,19 @@ +set(CMAKE_CXX_COMPILER "dpcpp") +if(WIN32) + set(CMAKE_C_COMPILER "clang-cl") +endif() + +cmake_minimum_required (VERSION 2.8) + +project(RollingMeanFilter) + +set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) +set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) +set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) + +# copy image files +file(COPY ${CMAKE_CURRENT_SOURCE_DIR}/input/ DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/input/) +file(COPY ${CMAKE_CURRENT_SOURCE_DIR}/output/ DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/output/) + +add_subdirectory (src) +add_subdirectory (src/Utils) diff --git a/rolling-mean-filter/License.txt b/rolling-mean-filter/License.txt new file mode 100644 index 0000000..7c8b8a3 --- /dev/null +++ b/rolling-mean-filter/License.txt @@ -0,0 +1,23 @@ +Copyright Intel Corporation + +SPDX-License-Identifier: MIT +https://opensource.org/licenses/MIT + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. + diff --git a/rolling-mean-filter/README.md b/rolling-mean-filter/README.md new file mode 100644 index 0000000..22cb262 --- /dev/null +++ b/rolling-mean-filter/README.md @@ -0,0 +1,111 @@ +# Rolling Mean Filter +This application details an implementation of a Rolling Mean Filter, and was created by combining the Buffered Host Stream and Zero Copy Data Transfer tutorials from Intels FPGA Code sample github. + +***Documentation***: The [DPC++ FPGA Code Samples Guide](https://software.intel.com/content/www/us/en/develop/articles/explore-dpcpp-through-intel-fpga-code-samples.html) helps you to navigate the samples and build your knowledge of DPC++ for FPGA.
+The [oneAPI DPC++ FPGA Optimization Guide](https://software.intel.com/content/www/us/en/develop/documentation/oneapi-fpga-optimization-guide) is the reference manual for targeting FPGAs through DPC++.
+The [oneAPI Programming Guide](https://software.intel.com/en-us/oneapi-programming-guide) is a general resource for target-independent DPC++ programming. + +| Optimized for | Description +--- |--- +| OS | Linux* Ubuntu* 18.04; Windows* 10 +| Hardware | Intel® FPGA Programmable Acceleration Card (PAC) D5005 (with Intel Stratix® 10 SX) +| Software | Intel® oneAPI DPC++ Compiler +| What you will learn | How to optimally stream data between the host and device to maximize throughput +| Time to complete | 3 hours (On target hardware) + +_Notice: SYCL USM host allocations (and therefore this tutorial) are only supported for the Intel® FPGA PAC D5005 (with Intel Stratix 10 SX)_ +_Notice: This tutorial demonstrates an implementation of host streaming that will be supplanted by better techniques in a future release. See the [Drawbacks and Future Work](#drawbacks-and-future-work)_ + +## Purpose +This application details an implementation of a Weighted Rolling Mean Filter. The application streams data from the host, to the device where it calculates the mean of an image, and applies the weightes, and back to the host. The techniques employed in this application are not specific to a CPU-FPGA system (like the one used in this tutorial); they apply to GPUs, multi-core CPUs, and other processing units. + +### Key Implementation Details +In this application, we will create a design where a *Producer* (running on the CPU) produces data into USM host allocations, a *Kernel* (running on the FPGA) processes this data and produces output into host allocations, and a *Consumer* (running on the CPU) consumes the data. Data is shared between the host and FPGA device via host pointers (pointers to USM host allocations). + +![](block_diagram.png) + + +## License +Code samples are licensed under the MIT license. + +Third party program Licenses can be found here: [third-party-programs.txt](https://github.com/oneapi-src/oneAPI-samples/blob/master/third-party-programs.txt) + +## Building the `rolling-mean` Tutorial + +### Running Samples in DevCloud +If running a sample in the Intel DevCloud, remember that you must specify the compute node (fpga_compile, fpga_runtime:arria10, or fpga_runtime:stratix10) as well as whether to run in batch or interactive mode. For more information see the Intel® oneAPI Base Toolkit Get Started Guide ([https://devcloud.intel.com/oneapi/documentation/base-toolkit/](https://devcloud.intel.com/oneapi/documentation/base-toolkit/)). + +When compiling for FPGA hardware, it is recommended to increase the job timeout to 12h. + +### On a Linux* System + +1. Generate the `Makefile` by running `cmake`. + ``` + mkdir build + cd build + ``` + To compile for the Intel® FPGA PAC D5005 (with Intel Stratix® 10 SX), run `cmake` using the command: + ``` + cmake .. + ``` + +2. Compile the design through the generated `Makefile`. The following build targets are provided, matching the recommended development flow: + + * Compile for emulation (fast compile time, targets emulated FPGA device): + ``` + make fpga_emu + ``` + * Generate the optimization report: + ``` + make report + ``` + * Compile for FPGA hardware (longer compile time, targets FPGA device): + ``` + make fpga + ``` + +## Examining the Reports +Locate `report.html` in the `rolling-mean.prj/reports/` directory. Open the report in any of Chrome*, Firefox*, Edge*, or Internet Explorer*. + +## Running the Sample + + 1. Run the sample on the FPGA emulator (the kernel executes on the CPU): + ``` + ./rolling-mean.fpga_emu (Linux) + ``` +2. Run the sample on the FPGA device: + ``` + ./rolling-mean.fpga (Linux) + ``` + +### Example of Output + +You should see the following output in the console: + +1. When running on the FPGA emulator + ``` + Iterations: 0 + + Starting kernel processing. + Kernel processing done. + PASSED + ``` + NOTE: The FPGA emulator does not accurately represent the performance (throughput or latency) of the kernels. + +2. When running on the FPGA device + ``` + Iterations: 4 + + Starting kernel processing. + Kernel processing done. + Starting kernel processing. + Kernel processing done. + Starting kernel processing. + Kernel processing done. + Starting kernel processing. + Kernel processing done. + + Average latency for the restricted USM kernel: 2830ms + + PASSED + ``` diff --git a/rolling-mean-filter/input/cat.bmp b/rolling-mean-filter/input/cat.bmp new file mode 100644 index 0000000..64b4512 Binary files /dev/null and b/rolling-mean-filter/input/cat.bmp differ diff --git a/rolling-mean-filter/input/cat_0.bmp b/rolling-mean-filter/input/cat_0.bmp new file mode 100644 index 0000000..2761b27 Binary files /dev/null and b/rolling-mean-filter/input/cat_0.bmp differ diff --git a/rolling-mean-filter/input/cat_1.bmp b/rolling-mean-filter/input/cat_1.bmp new file mode 100644 index 0000000..ab5d422 Binary files /dev/null and b/rolling-mean-filter/input/cat_1.bmp differ diff --git a/rolling-mean-filter/input/cat_10.bmp b/rolling-mean-filter/input/cat_10.bmp new file mode 100644 index 0000000..02a7b9b Binary files /dev/null and b/rolling-mean-filter/input/cat_10.bmp differ diff --git a/rolling-mean-filter/input/cat_100.bmp b/rolling-mean-filter/input/cat_100.bmp new file mode 100644 index 0000000..4c17d13 Binary files /dev/null and b/rolling-mean-filter/input/cat_100.bmp differ diff --git a/rolling-mean-filter/input/cat_101.bmp b/rolling-mean-filter/input/cat_101.bmp new file mode 100644 index 0000000..f442370 Binary files /dev/null and b/rolling-mean-filter/input/cat_101.bmp differ diff --git a/rolling-mean-filter/input/cat_102.bmp b/rolling-mean-filter/input/cat_102.bmp new file mode 100644 index 0000000..013591a Binary files /dev/null and b/rolling-mean-filter/input/cat_102.bmp differ diff --git a/rolling-mean-filter/input/cat_103.bmp b/rolling-mean-filter/input/cat_103.bmp new file mode 100644 index 0000000..cc41feb Binary files /dev/null and b/rolling-mean-filter/input/cat_103.bmp differ diff --git a/rolling-mean-filter/input/cat_104.bmp b/rolling-mean-filter/input/cat_104.bmp new file mode 100644 index 0000000..5b30e85 Binary files /dev/null and b/rolling-mean-filter/input/cat_104.bmp differ diff --git a/rolling-mean-filter/input/cat_105.bmp b/rolling-mean-filter/input/cat_105.bmp new file mode 100644 index 0000000..db819a7 Binary files /dev/null and b/rolling-mean-filter/input/cat_105.bmp differ diff --git a/rolling-mean-filter/input/cat_106.bmp b/rolling-mean-filter/input/cat_106.bmp new file mode 100644 index 0000000..d3e652e Binary files /dev/null and b/rolling-mean-filter/input/cat_106.bmp differ diff --git a/rolling-mean-filter/input/cat_107.bmp b/rolling-mean-filter/input/cat_107.bmp new file mode 100644 index 0000000..199f52e Binary files /dev/null and b/rolling-mean-filter/input/cat_107.bmp differ diff --git a/rolling-mean-filter/input/cat_11.bmp b/rolling-mean-filter/input/cat_11.bmp new file mode 100644 index 0000000..12e609f Binary files /dev/null and b/rolling-mean-filter/input/cat_11.bmp differ diff --git a/rolling-mean-filter/input/cat_12.bmp b/rolling-mean-filter/input/cat_12.bmp new file mode 100644 index 0000000..b2d4e02 Binary files /dev/null and b/rolling-mean-filter/input/cat_12.bmp differ diff --git a/rolling-mean-filter/input/cat_13.bmp b/rolling-mean-filter/input/cat_13.bmp new file mode 100644 index 0000000..f28a728 Binary files /dev/null and b/rolling-mean-filter/input/cat_13.bmp differ diff --git a/rolling-mean-filter/input/cat_14.bmp b/rolling-mean-filter/input/cat_14.bmp new file mode 100644 index 0000000..fff835d Binary files /dev/null and b/rolling-mean-filter/input/cat_14.bmp differ diff --git a/rolling-mean-filter/input/cat_15.bmp b/rolling-mean-filter/input/cat_15.bmp new file mode 100644 index 0000000..a495204 Binary files /dev/null and b/rolling-mean-filter/input/cat_15.bmp differ diff --git a/rolling-mean-filter/input/cat_16.bmp b/rolling-mean-filter/input/cat_16.bmp new file mode 100644 index 0000000..58bd439 Binary files /dev/null and b/rolling-mean-filter/input/cat_16.bmp differ diff --git a/rolling-mean-filter/input/cat_17.bmp b/rolling-mean-filter/input/cat_17.bmp new file mode 100644 index 0000000..c386924 Binary files /dev/null and b/rolling-mean-filter/input/cat_17.bmp differ diff --git a/rolling-mean-filter/input/cat_18.bmp b/rolling-mean-filter/input/cat_18.bmp new file mode 100644 index 0000000..e34a63a Binary files /dev/null and b/rolling-mean-filter/input/cat_18.bmp differ diff --git a/rolling-mean-filter/input/cat_19.bmp b/rolling-mean-filter/input/cat_19.bmp new file mode 100644 index 0000000..3efa039 Binary files /dev/null and b/rolling-mean-filter/input/cat_19.bmp differ diff --git a/rolling-mean-filter/input/cat_2.bmp b/rolling-mean-filter/input/cat_2.bmp new file mode 100644 index 0000000..db95371 Binary files /dev/null and b/rolling-mean-filter/input/cat_2.bmp differ diff --git a/rolling-mean-filter/input/cat_20.bmp b/rolling-mean-filter/input/cat_20.bmp new file mode 100644 index 0000000..dadaf61 Binary files /dev/null and b/rolling-mean-filter/input/cat_20.bmp differ diff --git a/rolling-mean-filter/input/cat_21.bmp b/rolling-mean-filter/input/cat_21.bmp new file mode 100644 index 0000000..ef69c8c Binary files /dev/null and b/rolling-mean-filter/input/cat_21.bmp differ diff --git a/rolling-mean-filter/input/cat_22.bmp b/rolling-mean-filter/input/cat_22.bmp new file mode 100644 index 0000000..60f97c9 Binary files /dev/null and b/rolling-mean-filter/input/cat_22.bmp differ diff --git a/rolling-mean-filter/input/cat_23.bmp b/rolling-mean-filter/input/cat_23.bmp new file mode 100644 index 0000000..80a6782 Binary files /dev/null and b/rolling-mean-filter/input/cat_23.bmp differ diff --git a/rolling-mean-filter/input/cat_24.bmp b/rolling-mean-filter/input/cat_24.bmp new file mode 100644 index 0000000..372dfa9 Binary files /dev/null and b/rolling-mean-filter/input/cat_24.bmp differ diff --git a/rolling-mean-filter/input/cat_25.bmp b/rolling-mean-filter/input/cat_25.bmp new file mode 100644 index 0000000..44efb98 Binary files /dev/null and b/rolling-mean-filter/input/cat_25.bmp differ diff --git a/rolling-mean-filter/input/cat_26.bmp b/rolling-mean-filter/input/cat_26.bmp new file mode 100644 index 0000000..4dcc293 Binary files /dev/null and b/rolling-mean-filter/input/cat_26.bmp differ diff --git a/rolling-mean-filter/input/cat_27.bmp b/rolling-mean-filter/input/cat_27.bmp new file mode 100644 index 0000000..360560a Binary files /dev/null and b/rolling-mean-filter/input/cat_27.bmp differ diff --git a/rolling-mean-filter/input/cat_28.bmp b/rolling-mean-filter/input/cat_28.bmp new file mode 100644 index 0000000..033f429 Binary files /dev/null and b/rolling-mean-filter/input/cat_28.bmp differ diff --git a/rolling-mean-filter/input/cat_29.bmp b/rolling-mean-filter/input/cat_29.bmp new file mode 100644 index 0000000..2e32491 Binary files /dev/null and b/rolling-mean-filter/input/cat_29.bmp differ diff --git a/rolling-mean-filter/input/cat_3.bmp b/rolling-mean-filter/input/cat_3.bmp new file mode 100644 index 0000000..982e007 Binary files /dev/null and b/rolling-mean-filter/input/cat_3.bmp differ diff --git a/rolling-mean-filter/input/cat_30.bmp b/rolling-mean-filter/input/cat_30.bmp new file mode 100644 index 0000000..dea459d Binary files /dev/null and b/rolling-mean-filter/input/cat_30.bmp differ diff --git a/rolling-mean-filter/input/cat_31.bmp b/rolling-mean-filter/input/cat_31.bmp new file mode 100644 index 0000000..51ce7f3 Binary files /dev/null and b/rolling-mean-filter/input/cat_31.bmp differ diff --git a/rolling-mean-filter/input/cat_32.bmp b/rolling-mean-filter/input/cat_32.bmp new file mode 100644 index 0000000..8114468 Binary files /dev/null and b/rolling-mean-filter/input/cat_32.bmp differ diff --git a/rolling-mean-filter/input/cat_33.bmp b/rolling-mean-filter/input/cat_33.bmp new file mode 100644 index 0000000..d6257b0 Binary files /dev/null and b/rolling-mean-filter/input/cat_33.bmp differ diff --git a/rolling-mean-filter/input/cat_34.bmp b/rolling-mean-filter/input/cat_34.bmp new file mode 100644 index 0000000..ace4d3f Binary files /dev/null and b/rolling-mean-filter/input/cat_34.bmp differ diff --git a/rolling-mean-filter/input/cat_35.bmp b/rolling-mean-filter/input/cat_35.bmp new file mode 100644 index 0000000..3293ec0 Binary files /dev/null and b/rolling-mean-filter/input/cat_35.bmp differ diff --git a/rolling-mean-filter/input/cat_36.bmp b/rolling-mean-filter/input/cat_36.bmp new file mode 100644 index 0000000..bb9b0b9 Binary files /dev/null and b/rolling-mean-filter/input/cat_36.bmp differ diff --git a/rolling-mean-filter/input/cat_37.bmp b/rolling-mean-filter/input/cat_37.bmp new file mode 100644 index 0000000..8842c22 Binary files /dev/null and b/rolling-mean-filter/input/cat_37.bmp differ diff --git a/rolling-mean-filter/input/cat_38.bmp b/rolling-mean-filter/input/cat_38.bmp new file mode 100644 index 0000000..6e48547 Binary files /dev/null and b/rolling-mean-filter/input/cat_38.bmp differ diff --git a/rolling-mean-filter/input/cat_39.bmp b/rolling-mean-filter/input/cat_39.bmp new file mode 100644 index 0000000..00de937 Binary files /dev/null and b/rolling-mean-filter/input/cat_39.bmp differ diff --git a/rolling-mean-filter/input/cat_4.bmp b/rolling-mean-filter/input/cat_4.bmp new file mode 100644 index 0000000..4f6e06e Binary files /dev/null and b/rolling-mean-filter/input/cat_4.bmp differ diff --git a/rolling-mean-filter/input/cat_40.bmp b/rolling-mean-filter/input/cat_40.bmp new file mode 100644 index 0000000..a0a5d8e Binary files /dev/null and b/rolling-mean-filter/input/cat_40.bmp differ diff --git a/rolling-mean-filter/input/cat_41.bmp b/rolling-mean-filter/input/cat_41.bmp new file mode 100644 index 0000000..563dff8 Binary files /dev/null and b/rolling-mean-filter/input/cat_41.bmp differ diff --git a/rolling-mean-filter/input/cat_42.bmp b/rolling-mean-filter/input/cat_42.bmp new file mode 100644 index 0000000..06513a6 Binary files /dev/null and b/rolling-mean-filter/input/cat_42.bmp differ diff --git a/rolling-mean-filter/input/cat_43.bmp b/rolling-mean-filter/input/cat_43.bmp new file mode 100644 index 0000000..23d34a7 Binary files /dev/null and b/rolling-mean-filter/input/cat_43.bmp differ diff --git a/rolling-mean-filter/input/cat_44.bmp b/rolling-mean-filter/input/cat_44.bmp new file mode 100644 index 0000000..0bbc1ee Binary files /dev/null and b/rolling-mean-filter/input/cat_44.bmp differ diff --git a/rolling-mean-filter/input/cat_45.bmp b/rolling-mean-filter/input/cat_45.bmp new file mode 100644 index 0000000..bcd5099 Binary files /dev/null and b/rolling-mean-filter/input/cat_45.bmp differ diff --git a/rolling-mean-filter/input/cat_46.bmp b/rolling-mean-filter/input/cat_46.bmp new file mode 100644 index 0000000..a9d1232 Binary files /dev/null and b/rolling-mean-filter/input/cat_46.bmp differ diff --git a/rolling-mean-filter/input/cat_47.bmp b/rolling-mean-filter/input/cat_47.bmp new file mode 100644 index 0000000..3f3d499 Binary files /dev/null and b/rolling-mean-filter/input/cat_47.bmp differ diff --git a/rolling-mean-filter/input/cat_48.bmp b/rolling-mean-filter/input/cat_48.bmp new file mode 100644 index 0000000..1351519 Binary files /dev/null and b/rolling-mean-filter/input/cat_48.bmp differ diff --git a/rolling-mean-filter/input/cat_49.bmp b/rolling-mean-filter/input/cat_49.bmp new file mode 100644 index 0000000..0901eef Binary files /dev/null and b/rolling-mean-filter/input/cat_49.bmp differ diff --git a/rolling-mean-filter/input/cat_5.bmp b/rolling-mean-filter/input/cat_5.bmp new file mode 100644 index 0000000..b111696 Binary files /dev/null and b/rolling-mean-filter/input/cat_5.bmp differ diff --git a/rolling-mean-filter/input/cat_50.bmp b/rolling-mean-filter/input/cat_50.bmp new file mode 100644 index 0000000..dff0b9f Binary files /dev/null and b/rolling-mean-filter/input/cat_50.bmp differ diff --git a/rolling-mean-filter/input/cat_51.bmp b/rolling-mean-filter/input/cat_51.bmp new file mode 100644 index 0000000..44d1c27 Binary files /dev/null and b/rolling-mean-filter/input/cat_51.bmp differ diff --git a/rolling-mean-filter/input/cat_52.bmp b/rolling-mean-filter/input/cat_52.bmp new file mode 100644 index 0000000..df9f897 Binary files /dev/null and b/rolling-mean-filter/input/cat_52.bmp differ diff --git a/rolling-mean-filter/input/cat_53.bmp b/rolling-mean-filter/input/cat_53.bmp new file mode 100644 index 0000000..bfe41d1 Binary files /dev/null and b/rolling-mean-filter/input/cat_53.bmp differ diff --git a/rolling-mean-filter/input/cat_54.bmp b/rolling-mean-filter/input/cat_54.bmp new file mode 100644 index 0000000..25d2799 Binary files /dev/null and b/rolling-mean-filter/input/cat_54.bmp differ diff --git a/rolling-mean-filter/input/cat_55.bmp b/rolling-mean-filter/input/cat_55.bmp new file mode 100644 index 0000000..38ff51b Binary files /dev/null and b/rolling-mean-filter/input/cat_55.bmp differ diff --git a/rolling-mean-filter/input/cat_56.bmp b/rolling-mean-filter/input/cat_56.bmp new file mode 100644 index 0000000..2c4799b Binary files /dev/null and b/rolling-mean-filter/input/cat_56.bmp differ diff --git a/rolling-mean-filter/input/cat_57.bmp b/rolling-mean-filter/input/cat_57.bmp new file mode 100644 index 0000000..07c9d70 Binary files /dev/null and b/rolling-mean-filter/input/cat_57.bmp differ diff --git a/rolling-mean-filter/input/cat_58.bmp b/rolling-mean-filter/input/cat_58.bmp new file mode 100644 index 0000000..8d9d792 Binary files /dev/null and b/rolling-mean-filter/input/cat_58.bmp differ diff --git a/rolling-mean-filter/input/cat_59.bmp b/rolling-mean-filter/input/cat_59.bmp new file mode 100644 index 0000000..c36e789 Binary files /dev/null and b/rolling-mean-filter/input/cat_59.bmp differ diff --git a/rolling-mean-filter/input/cat_6.bmp b/rolling-mean-filter/input/cat_6.bmp new file mode 100644 index 0000000..eecd4e6 Binary files /dev/null and b/rolling-mean-filter/input/cat_6.bmp differ diff --git a/rolling-mean-filter/input/cat_60.bmp b/rolling-mean-filter/input/cat_60.bmp new file mode 100644 index 0000000..b2709c2 Binary files /dev/null and b/rolling-mean-filter/input/cat_60.bmp differ diff --git a/rolling-mean-filter/input/cat_61.bmp b/rolling-mean-filter/input/cat_61.bmp new file mode 100644 index 0000000..5807898 Binary files /dev/null and b/rolling-mean-filter/input/cat_61.bmp differ diff --git a/rolling-mean-filter/input/cat_62.bmp b/rolling-mean-filter/input/cat_62.bmp new file mode 100644 index 0000000..f8e28fc Binary files /dev/null and b/rolling-mean-filter/input/cat_62.bmp differ diff --git a/rolling-mean-filter/input/cat_63.bmp b/rolling-mean-filter/input/cat_63.bmp new file mode 100644 index 0000000..30af06d Binary files /dev/null and b/rolling-mean-filter/input/cat_63.bmp differ diff --git a/rolling-mean-filter/input/cat_64.bmp b/rolling-mean-filter/input/cat_64.bmp new file mode 100644 index 0000000..eb00ab1 Binary files /dev/null and b/rolling-mean-filter/input/cat_64.bmp differ diff --git a/rolling-mean-filter/input/cat_65.bmp b/rolling-mean-filter/input/cat_65.bmp new file mode 100644 index 0000000..dab7d0d Binary files /dev/null and b/rolling-mean-filter/input/cat_65.bmp differ diff --git a/rolling-mean-filter/input/cat_66.bmp b/rolling-mean-filter/input/cat_66.bmp new file mode 100644 index 0000000..0fc8072 Binary files /dev/null and b/rolling-mean-filter/input/cat_66.bmp differ diff --git a/rolling-mean-filter/input/cat_67.bmp b/rolling-mean-filter/input/cat_67.bmp new file mode 100644 index 0000000..c05ec6b Binary files /dev/null and b/rolling-mean-filter/input/cat_67.bmp differ diff --git a/rolling-mean-filter/input/cat_68.bmp b/rolling-mean-filter/input/cat_68.bmp new file mode 100644 index 0000000..2713aac Binary files /dev/null and b/rolling-mean-filter/input/cat_68.bmp differ diff --git a/rolling-mean-filter/input/cat_69.bmp b/rolling-mean-filter/input/cat_69.bmp new file mode 100644 index 0000000..5f8a6c3 Binary files /dev/null and b/rolling-mean-filter/input/cat_69.bmp differ diff --git a/rolling-mean-filter/input/cat_7.bmp b/rolling-mean-filter/input/cat_7.bmp new file mode 100644 index 0000000..6ac67a2 Binary files /dev/null and b/rolling-mean-filter/input/cat_7.bmp differ diff --git a/rolling-mean-filter/input/cat_70.bmp b/rolling-mean-filter/input/cat_70.bmp new file mode 100644 index 0000000..d6fbde5 Binary files /dev/null and b/rolling-mean-filter/input/cat_70.bmp differ diff --git a/rolling-mean-filter/input/cat_71.bmp b/rolling-mean-filter/input/cat_71.bmp new file mode 100644 index 0000000..cc09ed0 Binary files /dev/null and b/rolling-mean-filter/input/cat_71.bmp differ diff --git a/rolling-mean-filter/input/cat_72.bmp b/rolling-mean-filter/input/cat_72.bmp new file mode 100644 index 0000000..41d5655 Binary files /dev/null and b/rolling-mean-filter/input/cat_72.bmp differ diff --git a/rolling-mean-filter/input/cat_73.bmp b/rolling-mean-filter/input/cat_73.bmp new file mode 100644 index 0000000..b1599a4 Binary files /dev/null and b/rolling-mean-filter/input/cat_73.bmp differ diff --git a/rolling-mean-filter/input/cat_74.bmp b/rolling-mean-filter/input/cat_74.bmp new file mode 100644 index 0000000..b613e52 Binary files /dev/null and b/rolling-mean-filter/input/cat_74.bmp differ diff --git a/rolling-mean-filter/input/cat_75.bmp b/rolling-mean-filter/input/cat_75.bmp new file mode 100644 index 0000000..b439f6c Binary files /dev/null and b/rolling-mean-filter/input/cat_75.bmp differ diff --git a/rolling-mean-filter/input/cat_76.bmp b/rolling-mean-filter/input/cat_76.bmp new file mode 100644 index 0000000..67e7b66 Binary files /dev/null and b/rolling-mean-filter/input/cat_76.bmp differ diff --git a/rolling-mean-filter/input/cat_77.bmp b/rolling-mean-filter/input/cat_77.bmp new file mode 100644 index 0000000..7282c04 Binary files /dev/null and b/rolling-mean-filter/input/cat_77.bmp differ diff --git a/rolling-mean-filter/input/cat_78.bmp b/rolling-mean-filter/input/cat_78.bmp new file mode 100644 index 0000000..4beaac6 Binary files /dev/null and b/rolling-mean-filter/input/cat_78.bmp differ diff --git a/rolling-mean-filter/input/cat_79.bmp b/rolling-mean-filter/input/cat_79.bmp new file mode 100644 index 0000000..0102d39 Binary files /dev/null and b/rolling-mean-filter/input/cat_79.bmp differ diff --git a/rolling-mean-filter/input/cat_8.bmp b/rolling-mean-filter/input/cat_8.bmp new file mode 100644 index 0000000..cd1acc5 Binary files /dev/null and b/rolling-mean-filter/input/cat_8.bmp differ diff --git a/rolling-mean-filter/input/cat_80.bmp b/rolling-mean-filter/input/cat_80.bmp new file mode 100644 index 0000000..1984bac Binary files /dev/null and b/rolling-mean-filter/input/cat_80.bmp differ diff --git a/rolling-mean-filter/input/cat_81.bmp b/rolling-mean-filter/input/cat_81.bmp new file mode 100644 index 0000000..2284816 Binary files /dev/null and b/rolling-mean-filter/input/cat_81.bmp differ diff --git a/rolling-mean-filter/input/cat_82.bmp b/rolling-mean-filter/input/cat_82.bmp new file mode 100644 index 0000000..af23f88 Binary files /dev/null and b/rolling-mean-filter/input/cat_82.bmp differ diff --git a/rolling-mean-filter/input/cat_83.bmp b/rolling-mean-filter/input/cat_83.bmp new file mode 100644 index 0000000..2361a4f Binary files /dev/null and b/rolling-mean-filter/input/cat_83.bmp differ diff --git a/rolling-mean-filter/input/cat_84.bmp b/rolling-mean-filter/input/cat_84.bmp new file mode 100644 index 0000000..7670906 Binary files /dev/null and b/rolling-mean-filter/input/cat_84.bmp differ diff --git a/rolling-mean-filter/input/cat_85.bmp b/rolling-mean-filter/input/cat_85.bmp new file mode 100644 index 0000000..16ea73f Binary files /dev/null and b/rolling-mean-filter/input/cat_85.bmp differ diff --git a/rolling-mean-filter/input/cat_86.bmp b/rolling-mean-filter/input/cat_86.bmp new file mode 100644 index 0000000..6f85ae5 Binary files /dev/null and b/rolling-mean-filter/input/cat_86.bmp differ diff --git a/rolling-mean-filter/input/cat_87.bmp b/rolling-mean-filter/input/cat_87.bmp new file mode 100644 index 0000000..b719985 Binary files /dev/null and b/rolling-mean-filter/input/cat_87.bmp differ diff --git a/rolling-mean-filter/input/cat_88.bmp b/rolling-mean-filter/input/cat_88.bmp new file mode 100644 index 0000000..defb89c Binary files /dev/null and b/rolling-mean-filter/input/cat_88.bmp differ diff --git a/rolling-mean-filter/input/cat_89.bmp b/rolling-mean-filter/input/cat_89.bmp new file mode 100644 index 0000000..ecb65eb Binary files /dev/null and b/rolling-mean-filter/input/cat_89.bmp differ diff --git a/rolling-mean-filter/input/cat_9.bmp b/rolling-mean-filter/input/cat_9.bmp new file mode 100644 index 0000000..169b075 Binary files /dev/null and b/rolling-mean-filter/input/cat_9.bmp differ diff --git a/rolling-mean-filter/input/cat_90.bmp b/rolling-mean-filter/input/cat_90.bmp new file mode 100644 index 0000000..58b0b92 Binary files /dev/null and b/rolling-mean-filter/input/cat_90.bmp differ diff --git a/rolling-mean-filter/input/cat_91.bmp b/rolling-mean-filter/input/cat_91.bmp new file mode 100644 index 0000000..25667b8 Binary files /dev/null and b/rolling-mean-filter/input/cat_91.bmp differ diff --git a/rolling-mean-filter/input/cat_92.bmp b/rolling-mean-filter/input/cat_92.bmp new file mode 100644 index 0000000..b85a11f Binary files /dev/null and b/rolling-mean-filter/input/cat_92.bmp differ diff --git a/rolling-mean-filter/input/cat_93.bmp b/rolling-mean-filter/input/cat_93.bmp new file mode 100644 index 0000000..11cfff9 Binary files /dev/null and b/rolling-mean-filter/input/cat_93.bmp differ diff --git a/rolling-mean-filter/input/cat_94.bmp b/rolling-mean-filter/input/cat_94.bmp new file mode 100644 index 0000000..258e72e Binary files /dev/null and b/rolling-mean-filter/input/cat_94.bmp differ diff --git a/rolling-mean-filter/input/cat_95.bmp b/rolling-mean-filter/input/cat_95.bmp new file mode 100644 index 0000000..87de2b9 Binary files /dev/null and b/rolling-mean-filter/input/cat_95.bmp differ diff --git a/rolling-mean-filter/input/cat_96.bmp b/rolling-mean-filter/input/cat_96.bmp new file mode 100644 index 0000000..9744c97 Binary files /dev/null and b/rolling-mean-filter/input/cat_96.bmp differ diff --git a/rolling-mean-filter/input/cat_97.bmp b/rolling-mean-filter/input/cat_97.bmp new file mode 100644 index 0000000..c5efa14 Binary files /dev/null and b/rolling-mean-filter/input/cat_97.bmp differ diff --git a/rolling-mean-filter/input/cat_98.bmp b/rolling-mean-filter/input/cat_98.bmp new file mode 100644 index 0000000..5538c49 Binary files /dev/null and b/rolling-mean-filter/input/cat_98.bmp differ diff --git a/rolling-mean-filter/input/cat_99.bmp b/rolling-mean-filter/input/cat_99.bmp new file mode 100644 index 0000000..9e3aa36 Binary files /dev/null and b/rolling-mean-filter/input/cat_99.bmp differ diff --git a/rolling-mean-filter/output/tmp.txt b/rolling-mean-filter/output/tmp.txt new file mode 100644 index 0000000..ee3ef64 --- /dev/null +++ b/rolling-mean-filter/output/tmp.txt @@ -0,0 +1 @@ +Empty file to add directory to git repo. diff --git a/rolling-mean-filter/src/CMakeLists.txt b/rolling-mean-filter/src/CMakeLists.txt new file mode 100644 index 0000000..b5d48e3 --- /dev/null +++ b/rolling-mean-filter/src/CMakeLists.txt @@ -0,0 +1,73 @@ +set(SOURCE_FILE rolling-mean-filter.cpp) +set(TARGET_NAME rolling-mean-filter) +set(EMULATOR_TARGET ${TARGET_NAME}.fpga_emu) +set(FPGA_TARGET ${TARGET_NAME}.fpga) +set(REPORTS_TARGET ${TARGET_NAME}_report) + +# USM is only supported on the PAC S10 board +set(S10_PAC_USM_BOARD_NAME "intel_s10sx_pac:pac_s10_usm") + +# Flags +if(WIN32) + set(THREAD_FLAG "") +else() + set(THREAD_FLAG "-lpthread") +endif() + +set(EMULATOR_COMPILE_FLAGS "-fintelfpga -Wall -DFPGA_EMULATOR") +set(EMULATOR_LINK_FLAGS "-fintelfpga ${THREAD_FLAG}") +set(HARDWARE_COMPILE_FLAGS "-fintelfpga -c") +set(HARDWARE_LINK_FLAGS "-fintelfpga ${THREAD_FLAG} -Wall -Xshardware -Xsboard=${S10_PAC_USM_BOARD_NAME} -reuse-exe=${CMAKE_BINARY_DIR}/${FPGA_TARGET} ${USER_HARDWARE_FLAGS}") +# use cmake -D USER_HARDWARE_FLAGS= to set extra flags for FPGA backend compilation + + +# FPGA emulator +if(WIN32) + set(WIN_EMULATOR_TARGET ${EMULATOR_TARGET}.exe) + add_custom_target(fpga_emu DEPENDS ${WIN_EMULATOR_TARGET}) + separate_arguments(WIN_EMULATOR_COMPILE_FLAGS WINDOWS_COMMAND "${EMULATOR_COMPILE_FLAGS}") + add_custom_command(OUTPUT ${WIN_EMULATOR_TARGET} + COMMAND ${CMAKE_CXX_COMPILER} /EHsc ${WIN_EMULATOR_COMPILE_FLAGS} ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_FILE} -o ${CMAKE_BINARY_DIR}/${WIN_EMULATOR_TARGET} + DEPENDS ${SOURCE_FILE}) +else() + add_executable(${EMULATOR_TARGET} ${SOURCE_FILE}) + add_custom_target(fpga_emu DEPENDS ${EMULATOR_TARGET}) + set_target_properties(${EMULATOR_TARGET} PROPERTIES COMPILE_FLAGS ${EMULATOR_COMPILE_FLAGS}) + set_target_properties(${EMULATOR_TARGET} PROPERTIES LINK_FLAGS ${EMULATOR_LINK_FLAGS}) + # to link with functions in Utils library (definded in Utils/ folder) + target_link_libraries (${EMULATOR_TARGET} LINK_PUBLIC Utils) +endif() + +# FPGA hardware +if(WIN32) + add_custom_target(fpga COMMAND echo "An FPGA hardware target is not provided on Windows. See README for details.") +else() + add_executable(${FPGA_TARGET} EXCLUDE_FROM_ALL ${SOURCE_FILE}) + add_custom_target(fpga DEPENDS ${FPGA_TARGET}) + set_target_properties(${FPGA_TARGET} PROPERTIES COMPILE_FLAGS ${HARDWARE_COMPILE_FLAGS}) + set_target_properties(${FPGA_TARGET} PROPERTIES LINK_FLAGS ${HARDWARE_LINK_FLAGS}) + # to link with functions in Utils library (definded in Utils/ folder) + target_link_libraries (${FPGA_TARGET} LINK_PUBLIC Utils) + +endif() + +# FPGA hardware report +if(WIN32) + set(DEVICE_OBJ_FILE ${TARGET_NAME}_report.a) + add_custom_target(report DEPENDS ${DEVICE_OBJ_FILE}) + separate_arguments(HARDWARE_LINK_FLAGS_LIST WINDOWS_COMMAND "${HARDWARE_LINK_FLAGS}") + add_custom_command(OUTPUT ${DEVICE_OBJ_FILE} + COMMAND ${CMAKE_CXX_COMPILER} /EHsc ${HARDWARE_LINK_FLAGS_LIST} -fsycl-link ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_FILE} -o ${CMAKE_BINARY_DIR}/${DEVICE_OBJ_FILE} + DEPENDS ${SOURCE_FILE}) +else() + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_FILE} ${SOURCE_FILE} COPYONLY) + + separate_arguments(HARDWARE_LINK_FLAGS_LIST UNIX_COMMAND "${HARDWARE_LINK_FLAGS}") + separate_arguments(CMAKE_CXX_FLAGS_LIST UNIX_COMMAND "${CMAKE_CXX_FLAGS}") + list(APPEND CMAKE_CXX_FLAGS_LIST -I${CMAKE_CURRENT_SOURCE_DIR}/Utils) + add_custom_command(OUTPUT ${REPORTS_TARGET} + COMMAND ${CMAKE_CXX_COMPILER} ${CMAKE_CXX_FLAGS_LIST} ${HARDWARE_LINK_FLAGS_LIST} -fsycl-link ${SOURCE_FILE} -o ${CMAKE_BINARY_DIR}/${REPORTS_TARGET} + DEPENDS ${SOURCE_FILE}) + add_custom_target(report DEPENDS ${REPORTS_TARGET}) +endif() + diff --git a/rolling-mean-filter/src/Utils/CMakeLists.txt b/rolling-mean-filter/src/Utils/CMakeLists.txt new file mode 100644 index 0000000..0a03b28 --- /dev/null +++ b/rolling-mean-filter/src/Utils/CMakeLists.txt @@ -0,0 +1,9 @@ +# Create a library called "Utils" which includes the source file *.c. +# The extension is already found. Any number of sources could be listed here. +file( GLOB UTILS_LIB_SOURCE_FILES *.c) +#message(CHECK_START "Found source files in Utils/ - ${UTILS_LIB_SOURCE_FILES}") +add_library (Utils STATIC ${UTILS_LIB_SOURCE_FILES}) + +# Make sure the compiler can find include files for our Utils library +# when other libraries or executables link to Hello +target_include_directories (Utils PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}) \ No newline at end of file diff --git a/rolling-mean-filter/src/Utils/bmp-utils.c b/rolling-mean-filter/src/Utils/bmp-utils.c new file mode 100644 index 0000000..6d528b9 --- /dev/null +++ b/rolling-mean-filter/src/Utils/bmp-utils.c @@ -0,0 +1,431 @@ +#include +#include + +#include "bmp-utils.h" + +void writeBmp(int *imageOut, const char *filename, int rows, int cols, + const char* refFilename) { + + FILE *ifp, *ofp; + unsigned char tmp; + int offset; + unsigned char *buffer; + int i, j; + int bytes; + int height, width; + size_t itemsRead; + + ifp = fopen(refFilename, "rb"); + if(ifp == NULL) { + perror(filename); + exit(-1); + } + + fseek(ifp, 10, SEEK_SET); + itemsRead = fread(&offset, 4, 1, ifp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + + + fseek(ifp, 18, SEEK_SET); + itemsRead = fread(&width, 4, 1, ifp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + itemsRead = fread(&height, 4, 1, ifp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + + fseek(ifp, 0, SEEK_SET); + + buffer = (unsigned char *)malloc(offset); + if(buffer == NULL) { + perror("malloc"); + exit(-1); + } + + itemsRead = fread(buffer, 1, offset, ifp); + if (itemsRead != offset) { + perror(filename); + exit(-1); + } + + //printf("Writing output image to %s\n", filename); + ofp = fopen(filename, "wb"); + if(ofp == NULL) { + perror("opening output file"); + exit(-1); + } + bytes = fwrite(buffer, 1, offset, ofp); + if(bytes != offset) { + perror("error writing header"); + exit(-1); + } + + // NOTE bmp formats store data in reverse raster order (see comment in + // readImage function), so we need to flip it upside down here. + int mod = width % 4; + if(mod != 0) { + mod = 4 - mod; + } + // printf("mod = %d\n", mod); + for(i = height-1; i >= 0; i--) { + for(j = 0; j < width; j++) { + tmp = (unsigned char)imageOut[i*cols+j]; + fwrite(&tmp, sizeof(char), 1, ofp); + } + // In bmp format, rows must be a multiple of 4-bytes. + // So if we're not at a multiple of 4, add junk padding. + for(j = 0; j < mod; j++) { + fwrite(&tmp, sizeof(char), 1, ofp); + } + } + + fclose(ofp); + fclose(ifp); + + free(buffer); +} + +void writeBmpFloat(float *imageOut, const char *filename, int rows, int cols, + const char* refFilename) { + + FILE *ifp, *ofp; + unsigned char tmp; + int offset; + unsigned char *buffer; + int i, j; + int bytes; + int height, width; + size_t itemsRead; + + ifp = fopen(refFilename, "rb"); + if(ifp == NULL) { + perror(filename); + exit(-1); + } + + fseek(ifp, 10, SEEK_SET); + itemsRead = fread(&offset, 4, 1, ifp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + + fseek(ifp, 18, SEEK_SET); + itemsRead = fread(&width, 4, 1, ifp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + itemsRead = fread(&height, 4, 1, ifp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + + fseek(ifp, 0, SEEK_SET); + + buffer = (unsigned char *)malloc(offset); + if(buffer == NULL) { + perror("malloc"); + exit(-1); + } + + itemsRead = fread(buffer, 1, offset, ifp); + if (itemsRead != offset) { + perror(filename); + exit(-1); + } + + //printf("Writing output image to %s\n", filename); + ofp = fopen(filename, "wb"); + if(ofp == NULL) { + perror("opening output file"); + exit(-1); + } + bytes = fwrite(buffer, 1, offset, ofp); + if(bytes != offset) { + perror("error writing header"); + exit(-1); + } + + // NOTE bmp formats store data in reverse raster order (see comment in + // readImage function), so we need to flip it upside down here. + int mod = width % 4; + if(mod != 0) { + mod = 4 - mod; + } + // printf("mod = %d\n", mod); + for(i = height-1; i >= 0; i--) { + for(j = 0; j < width; j++) { + tmp = (unsigned char)imageOut[i*cols+j]; + fwrite(&tmp, sizeof(char), 1, ofp); + } + // In bmp format, rows must be a multiple of 4-bytes. + // So if we're not at a multiple of 4, add junk padding. + for(j = 0; j < mod; j++) { + fwrite(&tmp, sizeof(char), 1, ofp); + } + } + + fclose(ofp); + fclose(ifp); + + free(buffer); +} + +/* + * Read bmp image and convert to byte array. Also output the width and height + */ +int* readBmp(const char *filename, int* rows, int* cols) { + + uchar* imageData; + int height, width; + uchar tmp; + int offset; + int i, j; + size_t itemsRead; + + //printf("Reading input image from %s\n", filename); + FILE *fp = fopen(filename, "rb"); + if(fp == NULL) { + perror(filename); + exit(-1); + } + + fseek(fp, 10, SEEK_SET); + itemsRead = fread(&offset, 4, 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + + fseek(fp, 18, SEEK_SET); + itemsRead = fread(&width, 4, 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + itemsRead = fread(&height, 4, 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + + //printf("width = %d\n", width); + //printf("height = %d\n", height); + + *cols = width; + *rows = height; + + imageData = (uchar*)malloc(width*height); + if(imageData == NULL) { + perror("malloc"); + exit(-1); + } + + fseek(fp, offset, SEEK_SET); + fflush(NULL); + + int mod = width % 4; + if(mod != 0) { + mod = 4 - mod; + } + + // NOTE bitmaps are stored in upside-down raster order. So we begin + // reading from the bottom left pixel, then going from left-to-right, + // read from the bottom to the top of the image. For image analysis, + // we want the image to be right-side up, so we'll modify it here. + + // First we read the image in upside-down + + // Read in the actual image + for(i = 0; i < height; i++) { + + // add actual data to the image + for(j = 0; j < width; j++) { + itemsRead = fread(&tmp, sizeof(char), 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + imageData[i*width + j] = tmp; + } + // For the bmp format, each row has to be a multiple of 4, + // so I need to read in the junk data and throw it away + for(j = 0; j < mod; j++) { + itemsRead = fread(&tmp, sizeof(char), 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + } + } + + // Then we flip it over + int flipRow; + for(i = 0; i < height/2; i++) { + flipRow = height - (i+1); + for(j = 0; j < width; j++) { + tmp = imageData[i*width+j]; + imageData[i*width+j] = imageData[flipRow*width+j]; + imageData[flipRow*width+j] = tmp; + } + } + + fclose(fp); + + // Input image on the host + int* intImage = NULL; + intImage = (int*)malloc(sizeof(int)*width*height); + if(intImage == NULL) { + perror("malloc"); + exit(-1); + } + + // Convert the BMP image to int (not required) + for(i = 0; i < height; i++) { + for(j = 0; j < width; j++) { + intImage[i*width+j] = (int)imageData[i*width+j]; + } + } + + free(imageData); + return intImage; +} + +/* + * Read bmp image and convert to byte array. Also output the width and height + */ +float* readBmpFloat(const char *filename, int* rows, int* cols) { + + uchar* imageData; + int height, width; + uchar tmp; + int offset; + int i, j; + short bits_per_pixel; + size_t itemsRead; + + //printf("Reading input image from %s\n", filename); + FILE *fp = fopen(filename, "rb"); + if(fp == NULL) { + perror(filename); + exit(-1); + } + + fseek(fp, 10, SEEK_SET); + itemsRead = fread(&offset, 4, 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + //printf("offset = %d\n", offset); + + fseek(fp, 18, SEEK_SET); + itemsRead = fread(&width, 4, 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + itemsRead = fread(&height, 4, 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + + //printf("width = %d\n", width); + //printf("height = %d\n", height); + + *cols = width; + *rows = height; + + fseek(fp, 28, SEEK_SET); + itemsRead = fread(&bits_per_pixel, 2, 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + //printf("bits per pixel = %d\n", bits_per_pixel); + + imageData = (uchar*)malloc(width*height); + if(imageData == NULL) { + perror("malloc"); + exit(-1); + } + + fseek(fp, offset, SEEK_SET); + fflush(NULL); + + int mod = width % 4; + if(mod != 0) { + mod = 4 - mod; + } + + // NOTE bitmaps are stored in upside-down raster order. So we begin + // reading from the bottom left pixel, then going from left-to-right, + // read from the bottom to the top of the image. For image analysis, + // we want the image to be right-side up, so we'll modify it here. + + // First we read the image in upside-down + + // Read in the actual image + for(i = 0; i < height; i++) { + + // add actual data to the image + for(j = 0; j < width; j++) { + itemsRead = fread(&tmp, sizeof(char), 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + imageData[i*width + j] = tmp; + } + // For the bmp format, each row has to be a multiple of 4, + // so I need to read in the junk data and throw it away + for(j = 0; j < mod; j++) { + itemsRead = fread(&tmp, sizeof(char), 1, fp); + if (itemsRead != 1) { + perror(filename); + exit(-1); + } + } + } + + // Then we flip it over + int flipRow; + for(i = 0; i < height/2; i++) { + flipRow = height - (i+1); + for(j = 0; j < width; j++) { + tmp = imageData[i*width+j]; + imageData[i*width+j] = imageData[flipRow*width+j]; + imageData[flipRow*width+j] = tmp; + } + } + + fclose(fp); + + // Input image on the host + float* floatImage = NULL; + floatImage = (float*)malloc(sizeof(float)*width*height); + if(floatImage == NULL) { + perror("malloc"); + exit(-1); + } + + // Convert the BMP image to float (not required) + for(i = 0; i < height; i++) { + for(j = 0; j < width; j++) { + floatImage[i*width+j] = (float)imageData[i*width+j]; + } + } + + free(imageData); + return floatImage; +} diff --git a/rolling-mean-filter/src/Utils/bmp-utils.h b/rolling-mean-filter/src/Utils/bmp-utils.h new file mode 100644 index 0000000..e1f8762 --- /dev/null +++ b/rolling-mean-filter/src/Utils/bmp-utils.h @@ -0,0 +1,21 @@ +#ifndef __BMPFUNCS_H__ +#define __BMPFUNCS_H__ + +#ifdef __cplusplus +extern "C" { +#endif + +typedef unsigned char uchar; + +int* readBmp(const char *filename, int* rows, int* cols); +void writeBmp(int *imageOut, const char *filename, int rows, int cols, + const char* refFilename); +float* readBmpFloat(const char *filename, int* rows, int* cols); +void writeBmpFloat(float *imageOut, const char *filename, int rows, int cols, + const char* refFilename); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/rolling-mean-filter/src/Utils/gold.c b/rolling-mean-filter/src/Utils/gold.c new file mode 100644 index 0000000..54e8d60 --- /dev/null +++ b/rolling-mean-filter/src/Utils/gold.c @@ -0,0 +1,160 @@ +#include +#include + +#include "gold.h" + +int* convolutionGold(int *image, int rows, int cols, float *filter, int filterWidth) +{ + /* Allocate space for the filtered image */ + int *outputImage; + outputImage = (int*)malloc(rows*cols*sizeof(int)); + if (!outputImage) { exit(-1); } + + /* Compute the filter width (intentionally truncate) */ + int halfFilterWidth = (int)filterWidth/2; + + /* Iterate over the rows of the source image */ + int i; + int j; + int k; + int l; + for (i = 0; i < rows; i++) + { + /* Iterate over the columns of the source image */ + for (j = 0; j < cols; j++) + { + /* Reset sum for new source pixel */ + int sum = 0; + + /* Apply the filter to the neighborhood */ + for (k = -halfFilterWidth; k <= halfFilterWidth; k++) + { + for (l = -halfFilterWidth; l <= halfFilterWidth; l++) + { + /* Indices used to access the image */ + int r = i+k; + int c = j+l; + + /* Handle out-of-bounds locations by clamping to + * the border pixel */ + r = (r < 0) ? 0 : r; + c = (c < 0) ? 0 : c; + r = (r >= rows) ? rows-1 : r; + c = (c >= cols) ? cols-1 : c; + + sum += image[r*cols+c] * + filter[(k+halfFilterWidth)*filterWidth + + (l+halfFilterWidth)]; + } + } + + /* Write the new pixel value */ + outputImage[i*cols+j] = sum; + } + } + + return outputImage; +} + +float* convolutionGoldFloat(float *image, int rows, int cols, float *filter, int filterWidth) +{ + /* Allocate space for the filtered image */ + float *outputImage; + outputImage = (float*)malloc(rows*cols*sizeof(float)); + if (!outputImage) { exit(-1); } + + /* Compute the filter width (intentionally truncate) */ + int halfFilterWidth = (int)filterWidth/2; + + /* Iterate over the rows of the source image */ + int i; + int j; + int k; + int l; + for (i = 0; i < rows; i++) + { + /* Iterate over the columns of the source image */ + for (j = 0; j < cols; j++) + { + /* Reset sum for new source pixel */ + float sum = 0; + + /* Apply the filter to the neighborhood */ + for (k = -halfFilterWidth; k <= halfFilterWidth; k++) + { + for (l = -halfFilterWidth; l <= halfFilterWidth; l++) + { + /* Indices used to access the image */ + int r = i+k; + int c = j+l; + + /* Handle out-of-bounds locations by clamping to + * the border pixel */ + r = (r < 0) ? 0 : r; + c = (c < 0) ? 0 : c; + r = (r >= rows) ? rows-1 : r; + c = (c >= cols) ? cols-1 : c; + + sum += image[r*cols+c] * + filter[(k+halfFilterWidth)*filterWidth + + (l+halfFilterWidth)]; + } + } + + /* Write the new pixel value */ + outputImage[i*cols+j] = sum; + } + } + + return outputImage; +} + +int* histogramGold(int *data, int items, int bins) +{ + int *refHistogram; + + /* Allocate space for the histogram */ + refHistogram = (int*)malloc(bins*sizeof(int)); + if (!refHistogram) { exit(-1); } + + /* Initialize the histogram to zero */ + int i; + for (i = 0; i < bins; i++) { + refHistogram[i] = 0; + } + + /* Compute the histogram */ + for (i = 0; i < items; i++) { + if (data[i] >= bins) { + printf("Warning: Histogram data value out-of-bounds\n"); + } + refHistogram[data[i]]++; + } + + return refHistogram; +} + +int* histogramGoldFloat(float *data, int items, int bins) +{ + int *refHistogram; + + /* Allocate space for the histogram */ + refHistogram = (int*)malloc(bins*sizeof(int)); + if (!refHistogram) { exit(-1); } + + /* Initialize the histogram to zero */ + int i; + for (i = 0; i < bins; i++) { + refHistogram[i] = 0; + } + + /* Compute the histogram */ + for (i = 0; i < items; i++) { + if ((int)data[i] >= bins) { + printf("Warning: Histogram data value out-of-bounds\n"); + } + refHistogram[(int)data[i]]++; + } + + return refHistogram; +} diff --git a/rolling-mean-filter/src/Utils/gold.h b/rolling-mean-filter/src/Utils/gold.h new file mode 100644 index 0000000..a83f99a --- /dev/null +++ b/rolling-mean-filter/src/Utils/gold.h @@ -0,0 +1,20 @@ +#ifndef __GOLD_H__ +#define __GOLD_H__ + +#ifdef __cplusplus +extern "C" { +#endif + +int* convolutionGold(int *image, int rows, int cols, + float *filter, int filterWidth); +float* convolutionGoldFloat(float *image, int rows, int cols, + float *filter, int filterWidth); + +int* histogramGold(int *data, int items, int bins); +int* histogramGoldFloat(float *data, int items, int bins); + +#ifdef __cplusplus +} +#endif + +#endif diff --git a/rolling-mean-filter/src/Utils/utils.c b/rolling-mean-filter/src/Utils/utils.c new file mode 100644 index 0000000..dd9ed40 --- /dev/null +++ b/rolling-mean-filter/src/Utils/utils.c @@ -0,0 +1,53 @@ +/* System includes */ +#include +#include + +char* readFile(const char *filename) { + + FILE *fp; + char *fileData; + long fileSize; + + /* Open the file */ + fp = fopen(filename, "r"); + if (!fp) { + printf("Could not open file: %s\n", filename); + exit(-1); + } + + /* Determine the file size */ + if (fseek(fp, 0, SEEK_END)) { + printf("Error reading the file\n"); + exit(-1); + } + fileSize = ftell(fp); + if (fileSize < 0) { + printf("Error reading the file\n"); + exit(-1); + } + if (fseek(fp, 0, SEEK_SET)) { + printf("Error reading the file\n"); + exit(-1); + } + + /* Read the contents */ + fileData = (char*)malloc(fileSize + 1); + if (!fileData) { + exit(-1); + } + if (fread(fileData, fileSize, 1, fp) != 1) { + printf("Error reading the file\n"); + exit(-1); + } + + /* Terminate the string */ + fileData[fileSize] = '\0'; + + /* Close the file */ + if (fclose(fp)) { + printf("Error closing the file\n"); + exit(-1); + } + + return fileData; +} diff --git a/rolling-mean-filter/src/Utils/utils.h b/rolling-mean-filter/src/Utils/utils.h new file mode 100644 index 0000000..02836e5 --- /dev/null +++ b/rolling-mean-filter/src/Utils/utils.h @@ -0,0 +1,17 @@ +#ifndef __UTILS_H__ +#define __UTILS_H__ + +/* OpenCL includes */ +#ifdef __APPLE__ +#include +#else +#include +#endif + +void check(cl_int); + +void printCompilerError(cl_program program, cl_device_id device); + +char* readFile(const char *filename); + +#endif diff --git a/rolling-mean-filter/src/kernel_defs.hpp b/rolling-mean-filter/src/kernel_defs.hpp new file mode 100644 index 0000000..145544c --- /dev/null +++ b/rolling-mean-filter/src/kernel_defs.hpp @@ -0,0 +1,301 @@ +/////////////////////////////////////////////////////////////////////////////// +// +// The structure of this file is modeled after restricted_usm_kernel.hpp +// in the DPC++-FPGA Tutorial for Zero Copy Data Transfer +// +// +// +// The structure of the kernels in this design is shown in the diagram below. +// The Producer kernel reads the data from CPU memory (via PCIe), producing it +// for the RestrictedUSM via a pipe. The Worker does the computation on the +// input data and writes it to the ConsumePipe. The consumer reads the data +// from this pipe and writes the output back to the CPU memory (via PCIe). +// +// |---------------| |-------------| |----------------------------------| +// | CPU | | | | FPGA | +// | | | | | | +// | |-----------| | | | | |----------| |---------------| | +// | | Producer |--->| |--->| Producer |==>| | | +// | |-----------| | | | | |----------| | | | +// | | | Host Memory | | | RestrictedUSM | | +// | |-----------| | | | | |----------| | | | +// | | Consumer |<---| |<---| Consumer |<==| | | +// | |-----------| | | | | |----------| |---------------| | +// | | | | | | +// |---------------| |-------------| |----------------------------------| +// +// +// As shown in the image above and the code below, we have split this design +// into three kernels: +// 1) Producer +// 2) RestrictedUSM +// 3) Consumer +// We do this to decouple the reads/writes from/to the Host Memory over PCIe. +// Decoupling the memory accesses and using SYCL pipes with a substantial +// depth ('kPipeDepth' below) allows the kernel to be more resilient against +// stalls while reading/writing from/to the Host Memory over PCIe. +// +/////////////////////////////////////////////////////////////////////////////// + +#ifndef __KERNEL_DEFS_HPP__ +#define __KERNEL_DEFS_HPP__ +#pragma once + +#include + +#include "Utils/bmp-utils.h" + +#include +#include + +using namespace sycl; +using namespace std::chrono; + +// Forward declare the kernel names in the global scope. +// This FPGA best practice reduces name mangling in the optimization reports. +class RestrictedUSM; +class Producer; +class Consumer; + +// pipes +constexpr size_t kPipeDepth = 64; +template +using ProducePipe = pipe; +template +using ConsumePipe = pipe; + +static const char* inputImagePath = "./input/cat_0.bmp"; + +static int g_imageRows = 0; +static int g_imageCols = 0; + +static float meanFilter[9] = { + 0.11f, 0.11f, 0.11f, + 0.11f, 0.11f, 0.11f, + 0.11f, 0.11f, 0.11f}; + +/* Compute the filter width (intentionally truncate) */ +static const int FilterWidth = 3; +static const int halfFilterWidth = (int)FilterWidth/2; + +const bool debug = false; + +// +// The producer thread function. All production of data happens in this function +// which is run in a separate CPU thread from the launching of kernels and +// consumption of data. +// +template +void ProducerThread(T* in_buffer, size_t size) +{ + // In our case, the Producer's job is simple. It has an input stream of data + // ('in_stream') whose size is buffer_count * reps elements + // (i.e. ALL of the data). + // When signalled to, it produces buffer_count elements to 'out_ptr', which is + // a shared variable between the Producer thread and the thread launching the + // kernels. + int imageRows; + int imageCols; + + size_t rep = 0; + + while (rep < size) + { + std::string input_filename = "./input/cat_" + std::to_string(rep) + ".bmp"; + + if (debug) + std::cout << "Reading image " << input_filename.c_str() << std::endl; + + T in_image = readBmpFloat(input_filename.c_str(), &imageRows, &imageCols); + + if (imageRows == g_imageRows && imageCols == g_imageCols) + { + in_buffer[rep] = in_image; + } + else + { + std::cerr << "Image row/col size does not match the first image." << std::endl; + std::terminate(); + } + rep++; + } +} + + +// +// The consumer thread function. All consumption of data happens in this function +// which is run in a separate CPU thread from the launching of kernels and +// production of data. +// +template +void ConsumerThread(T* out_buffer, size_t size) +{ + // In our case, the Producer's job is simple. It has an input stream of data + // ('in_stream') whose size is buffer_count * reps elements + // (i.e. ALL of the data). + // When signalled to, it produces buffer_count elements to 'out_ptr', which is + // a shared variable between the Producer thread and the thread launching the + // kernels. + size_t rep = 0; + + while (rep < size) + { + std::string output_filename = "./output/cat_" + std::to_string(rep) + ".bmp"; + if (debug) { + std::cout << "Output image saved as " << output_filename.c_str() << std::endl; + } + writeBmpFloat(out_buffer[rep], output_filename.c_str(), + g_imageRows, g_imageCols, inputImagePath); + rep++; + } +} + +// +// reads the input data from the hosts memory +// and writes it to the ProducePipe +// +template +event SubmitProducer(queue& q, T* in_data, size_t size) { + auto e = q.submit([&](handler& h) { + h.single_task([=]() [[intel::kernel_args_restrict]] { + // using a host_ptr tells the compiler that this pointer lives in the + // hosts address space + host_ptr h_in_data(in_data); + + for (size_t i = 0; i < size; i++) { + T data_from_host_memory = *(h_in_data + i); + ProducePipe::write(data_from_host_memory); + } + }); + }); + + return e; +} + +// +// The worker kernel in the device: +// 1) read input data from the ProducePipe +// 2) perform computation +// 3) write the output data to the ConsumePipe +// +template +event SubmitWorker(queue& q, size_t size, int imageRows, int imageCols) +{ + auto e = q.submit([&](handler& h) + { + h.single_task([=]() [[intel::kernel_args_restrict]] + { + for (size_t i = 0; i < size; i++) + { + T data = ProducePipe::read(); + + T value = data; + + for (int row = 0; row < imageRows; row++) + { + for (int col = 0; col < imageCols; col++) + { + float sum = 0.0f; + + /* Apply the filter to the neighborhood */ + for (int k = -halfFilterWidth; k <= halfFilterWidth; k++) + { + for (int l = -halfFilterWidth; l <= halfFilterWidth; l++) + { + /* Indices used to access the image */ + int r = row+k; + int c = col+l; + + /* Handle out-of-bounds locations by clamping to + * the border pixel */ + r = (r < 0) ? 0 : r; + c = (c < 0) ? 0 : c; + r = (r >= imageRows) ? imageRows-1 : r; + c = (c >= imageCols) ? imageCols-1 : c; + + sum += data[r*imageCols+c] * + meanFilter[(k+halfFilterWidth)*FilterWidth + (l+halfFilterWidth)]; + } + } + + /* Write the new pixel value */ + value[row*imageCols+col] = sum; + } + } + + ConsumePipe::write(value); + } + }); + }); + + return e; +} + +// +// reads output data from the device via ConsumePipe +// and writes it to the hosts memory +// +template +event SubmitConsumer(queue& q, T* out_data, size_t size) +{ + auto e = q.submit([&](handler& h) + { + // using a host_ptr tells the compiler that this pointer lives in the + // hosts address space + host_ptr h_out_data(out_data); + + h.single_task([=]() [[intel::kernel_args_restrict]] + { + for (size_t i = 0; i < size; i++) + { + T data_to_host_memory = ConsumePipe::read(); + *(h_out_data + i) = data_to_host_memory; + } + }); + }); + + return e; +} + +template +double Run_Iteration(queue& q, T* in, T* out, size_t size) +{ + // Read the first image sepatately. + // This allows us to set the global image rows/cols so that + // we can add error checking for reading subsequent images. + T test_image = readBmpFloat(inputImagePath, &g_imageRows, &g_imageCols); + + // start the timer + auto start = high_resolution_clock::now(); + + // start the Producer in a new thread (starts in the ProducerThread function) + std::thread producer_thread(ProducerThread, in, size); + + // wait for producer thread to finish + producer_thread.join(); + + std::cout << "Starting kernel processing." << std::endl; + + // start the kernels + auto worker_event = SubmitWorker(q, size, g_imageRows, g_imageCols); + auto producer_event = SubmitProducer(q, in, size); + auto consumer_event = SubmitConsumer(q, out, size); + + // wait for all the kernels to finish + q.wait(); + std::cout << "Kernel processing done." << std::endl; + + // start the Consumer in a new thread (starts in the ConsumerThread function) + std::thread consumer_thread(ConsumerThread, out, size); + + // wait for consumer thread to finish + consumer_thread.join(); + + // stop the timer + auto end = high_resolution_clock::now(); + duration diff = end - start; + + return diff.count(); +} + +#endif /* __KERNEL_DEFS_HPP__ */ \ No newline at end of file diff --git a/rolling-mean-filter/src/rolling-mean-filter.cpp b/rolling-mean-filter/src/rolling-mean-filter.cpp new file mode 100644 index 0000000..d561f5e --- /dev/null +++ b/rolling-mean-filter/src/rolling-mean-filter.cpp @@ -0,0 +1,178 @@ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +// dpc_common.hpp can be found in the dev-utilities include folder. +// e.g., $ONEAPI_ROOT/dev-utilities//include/dpc_common.hpp +#include "dpc_common.hpp" + +#include "kernel_defs.hpp" + +using namespace sycl; + +// the type used +// NOTE: the tutorial assumes the use of a sycl::vec datatype (like long8). +// Therefore, 'Type' must be a sycl::vec datatype (e.g. int8, char64, etc). +using Type = float*; +#define IMAGE_SIZE (720*1080) + +// the main function +int main(int argc, char* argv[]) +{ + // parse command line arguments + #if defined(FPGA_EMULATOR) + size_t size = 108; + size_t iterations = 1; + #else + size_t size = 108; + size_t iterations = 5; + #endif + + bool need_help = false; + + // parse the command line arguments + for (int i = 1; i < argc; i++) + { + std::string arg(argv[i]); + + if (arg == "--help" || arg == "-h") + { + need_help = true; + } + else + { + std::string str_after_equals = arg.substr(arg.find("=") + 1); + + if (arg.find("--iterations=") == 0) + { + iterations = std::max(2, atoi(str_after_equals.c_str()) + 1); + } + else + { + std::cout << "WARNING: ignoring unknown argument '" << arg << "'\n"; + } + } + } + + // print help is asked + if (need_help) + { + std::cout << "USAGE: " + << "./rolling-mean-filter " + << "[--iterations=]\n"; + return 0; + } + + // check the number of iterations + if (iterations <= 0) + { + std::cerr << "ERROR: 'iterations' must be positive\n"; + std::terminate(); + } + + // print info + std::cout << "Iterations: " << iterations-1 << "\n"; + std::cout << "\n"; + + try + { + // device selector + #if defined(FPGA_EMULATOR) + INTEL::fpga_emulator_selector selector; + #else + INTEL::fpga_selector selector; + #endif + + // queue properties to enable profiling + property_list prop_list { property::queue::enable_profiling() }; + + // create the device queue + queue q(selector, dpc_common::exception_handler, prop_list); + + // make sure the device supports USM host allocations + device d = q.get_device(); + if (!d.get_info()) + { + std::cerr << "ERROR: The selected device does not support USM host allocations\n"; + std::terminate(); + } + + // + // Setup the USM input/outputs + // malloc_host allocates memory specifically in the host's address space + auto in_restricted_usm = malloc_host(IMAGE_SIZE * size, q.get_context()); + auto out_restricted_usm = malloc_host(IMAGE_SIZE * size, q.get_context()); + + // Make sure we were able to allocate space for the input and output + if (in_restricted_usm == NULL) + { + std::cerr << "ERROR:\tFailed to allocate space for 'in_restricted_usm'." << std::endl; + return 1; + } + if (out_restricted_usm == NULL) + { + std::cerr << "ERROR:\tFailed to allocate space for 'out_restricted_usm'." << std::endl; + return 1; + } + + // Buffer to hold latency data + std::vector usm_latency(iterations); + + for (size_t i = 0; i < iterations; i++) + { + usm_latency[i] = Run_Iteration(q, in_restricted_usm, out_restricted_usm, size); + } + + // The FPGA emulator does not accurately represent the hardware performance + // so we don't print performance results when running with the emulator + #ifndef FPGA_EMULATOR + // Compute the average latency across all iterations. + // We use the first iteration as a 'warmup' for the FPGA, + // so we ignore its results. + double usm_avg_lat = + std::accumulate(usm_latency.begin() + 1, + usm_latency.end(), 0.0) / + (iterations - 1); + + std::cout << "Average latency for the restricted USM kernel: " + << usm_avg_lat << " ms\n"; + #endif + + // free the allocated host usm memory + // note that these are calls to sycl::free() + free(in_restricted_usm, q); + free(out_restricted_usm, q); + + } + catch (exception const& e) + { + // Catches exceptions in the host code + std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n"; + + // Most likely the runtime couldn't find FPGA hardware! + if (e.get_cl_code() == CL_DEVICE_NOT_FOUND) + { + std::cerr << "If you are targeting an FPGA, please ensure that your " + "system has a correctly configured FPGA board.\n"; + std::cerr << "Run sys_check in the oneAPI root directory to verify.\n"; + std::cerr << "If you are targeting the FPGA emulator, compile with " + "-DFPGA_EMULATOR.\n"; + } + std::terminate(); + } + + std::cout << "PASSED\n"; + return 0; + +} + diff --git a/rolling-mean-filter/third-party-programs.txt b/rolling-mean-filter/third-party-programs.txt new file mode 100644 index 0000000..90daff4 --- /dev/null +++ b/rolling-mean-filter/third-party-programs.txt @@ -0,0 +1,253 @@ +oneAPI Code Samples - Third Party Programs File + +This file contains the list of third party software ("third party programs") +contained in the Intel software and their required notices and/or license +terms. This third party software, even if included with the distribution of the +Intel software, may be governed by separate license terms, including without +limitation, third party license terms, other Intel software license terms, and +open source software license terms. These separate license terms govern your use +of the third party programs as set forth in the “third-party-programs.txt” or +other similarly named text file. + +Third party programs and their corresponding required notices and/or license +terms are listed below. + +-------------------------------------------------------------------------------- + +1. Nothings STB Libraries + +stb/LICENSE + + This software is available under 2 licenses -- choose whichever you prefer. + ------------------------------------------------------------------------------ + ALTERNATIVE A - MIT License + Copyright (c) 2017 Sean Barrett + Permission is hereby granted, free of charge, to any person obtaining a copy of + this software and associated documentation files (the "Software"), to deal in + the Software without restriction, including without limitation the rights to + use, copy, modify, merge, publish, distribute, sublicense, and/or sell copies + of the Software, and to permit persons to whom the Software is furnished to do + so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in all + copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE + SOFTWARE. + ------------------------------------------------------------------------------ + ALTERNATIVE B - Public Domain (www.unlicense.org) + This is free and unencumbered software released into the public domain. + Anyone is free to copy, modify, publish, use, compile, sell, or distribute this + software, either in source code form or as a compiled binary, for any purpose, + commercial or non-commercial, and by any means. + In jurisdictions that recognize copyright laws, the author or authors of this + software dedicate any and all copyright interest in the software to the public + domain. We make this dedication for the benefit of the public at large and to + the detriment of our heirs and successors. We intend this dedication to be an + overt act of relinquishment in perpetuity of all present and future rights to + this software under copyright law. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN + ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION + WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. + +-------------------------------------------------------------------------------- + +2. FGPA example designs-gzip + + SDL2.0 + +zlib License + + + This software is provided 'as-is', without any express or implied + warranty. In no event will the authors be held liable for any damages + arising from the use of this software. + + Permission is granted to anyone to use this software for any purpose, + including commercial applications, and to alter it and redistribute it + freely, subject to the following restrictions: + + 1. The origin of this software must not be misrepresented; you must not + claim that you wrote the original software. If you use this software + in a product, an acknowledgment in the product documentation would be + appreciated but is not required. + 2. Altered source versions must be plainly marked as such, and must not be + misrepresented as being the original software. + 3. This notice may not be removed or altered from any source distribution. + + +-------------------------------------------------------------------------------- + +3. Nbody + (c) 2019 Fabio Baruffa + + Plotly.js + Copyright (c) 2020 Plotly, Inc + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. +© 2020 GitHub, Inc. + +-------------------------------------------------------------------------------- + +4. GNU-EFI + Copyright (c) 1998-2000 Intel Corporation + +The files in the "lib" and "inc" subdirectories are using the EFI Application +Toolkit distributed by Intel at http://developer.intel.com/technology/efi + +This code is covered by the following agreement: + +Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: + +Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. + +Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. + +THIS SOFTWARE IS PROVIDED ``AS IS'' AND ANY EXPRESS OR IMPLIED WARRANTIES, +INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND +FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL INTEL BE +LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE +POSSIBILITY OF SUCH DAMAGE. THE EFI SPECIFICATION AND ALL OTHER INFORMATION +ON THIS WEB SITE ARE PROVIDED "AS IS" WITH NO WARRANTIES, AND ARE SUBJECT +TO CHANGE WITHOUT NOTICE. + +-------------------------------------------------------------------------------- + +5. Edk2 + Copyright (c) 2019, Intel Corporation. All rights reserved. + + Edk2 Basetools + Copyright (c) 2019, Intel Corporation. All rights reserved. + +SPDX-License-Identifier: BSD-2-Clause-Patent + +-------------------------------------------------------------------------------- + +6. Heat Transmission + +GNU LESSER GENERAL PUBLIC LICENSE +Version 3, 29 June 2007 + +Copyright © 2007 Free Software Foundation, Inc. + +Everyone is permitted to copy and distribute verbatim copies of this license document, but changing it is not allowed. + +This version of the GNU Lesser General Public License incorporates the terms and conditions of version 3 of the GNU General Public License, supplemented by the additional permissions listed below. + +0. Additional Definitions. +As used herein, “this License” refers to version 3 of the GNU Lesser General Public License, and the “GNU GPL” refers to version 3 of the GNU General Public License. + +“The Library” refers to a covered work governed by this License, other than an Application or a Combined Work as defined below. + +An “Application” is any work that makes use of an interface provided by the Library, but which is not otherwise based on the Library. Defining a subclass of a class defined by the Library is deemed a mode of using an interface provided by the Library. + +A “Combined Work” is a work produced by combining or linking an Application with the Library. The particular version of the Library with which the Combined Work was made is also called the “Linked Version”. + +The “Minimal Corresponding Source” for a Combined Work means the Corresponding Source for the Combined Work, excluding any source code for portions of the Combined Work that, considered in isolation, are based on the Application, and not on the Linked Version. + +The “Corresponding Application Code” for a Combined Work means the object code and/or source code for the Application, including any data and utility programs needed for reproducing the Combined Work from the Application, but excluding the System Libraries of the Combined Work. + +1. Exception to Section 3 of the GNU GPL. +You may convey a covered work under sections 3 and 4 of this License without being bound by section 3 of the GNU GPL. + +2. Conveying Modified Versions. +If you modify a copy of the Library, and, in your modifications, a facility refers to a function or data to be supplied by an Application that uses the facility (other than as an argument passed when the facility is invoked), then you may convey a copy of the modified version: + +a) under this License, provided that you make a good faith effort to ensure that, in the event an Application does not supply the function or data, the facility still operates, and performs whatever part of its purpose remains meaningful, or +b) under the GNU GPL, with none of the additional permissions of this License applicable to that copy. +3. Object Code Incorporating Material from Library Header Files. +The object code form of an Application may incorporate material from a header file that is part of the Library. You may convey such object code under terms of your choice, provided that, if the incorporated material is not limited to numerical parameters, data structure layouts and accessors, or small macros, inline functions and templates (ten or fewer lines in length), you do both of the following: + +a) Give prominent notice with each copy of the object code that the Library is used in it and that the Library and its use are covered by this License. +b) Accompany the object code with a copy of the GNU GPL and this license document. +4. Combined Works. +You may convey a Combined Work under terms of your choice that, taken together, effectively do not restrict modification of the portions of the Library contained in the Combined Work and reverse engineering for debugging such modifications, if you also do each of the following: + +a) Give prominent notice with each copy of the Combined Work that the Library is used in it and that the Library and its use are covered by this License. +b) Accompany the Combined Work with a copy of the GNU GPL and this license document. +c) For a Combined Work that displays copyright notices during execution, include the copyright notice for the Library among these notices, as well as a reference directing the user to the copies of the GNU GPL and this license document. +d) Do one of the following: +0) Convey the Minimal Corresponding Source under the terms of this License, and the Corresponding Application Code in a form suitable for, and under terms that permit, the user to recombine or relink the Application with a modified version of the Linked Version to produce a modified Combined Work, in the manner specified by section 6 of the GNU GPL for conveying Corresponding Source. +1) Use a suitable shared library mechanism for linking with the Library. A suitable mechanism is one that (a) uses at run time a copy of the Library already present on the user's computer system, and (b) will operate properly with a modified version of the Library that is interface-compatible with the Linked Version. +e) Provide Installation Information, but only if you would otherwise be required to provide such information under section 6 of the GNU GPL, and only to the extent that such information is necessary to install and execute a modified version of the Combined Work produced by recombining or relinking the Application with a modified version of the Linked Version. (If you use option 4d0, the Installation Information must accompany the Minimal Corresponding Source and Corresponding Application Code. If you use option 4d1, you must provide the Installation Information in the manner specified by section 6 of the GNU GPL for conveying Corresponding Source.) +5. Combined Libraries. +You may place library facilities that are a work based on the Library side by side in a single library together with other library facilities that are not Applications and are not covered by this License, and convey such a combined library under terms of your choice, if you do both of the following: + +a) Accompany the combined library with a copy of the same work based on the Library, uncombined with any other library facilities, conveyed under the terms of this License. +b) Give prominent notice with the combined library that part of it is a work based on the Library, and explaining where to find the accompanying uncombined form of the same work. +6. Revised Versions of the GNU Lesser General Public License. +The Free Software Foundation may publish revised and/or new versions of the GNU Lesser General Public License from time to time. Such new versions will be similar in spirit to the present version, but may differ in detail to address new problems or concerns. + +Each version is given a distinguishing version number. If the Library as you received it specifies that a certain numbered version of the GNU Lesser General Public License “or any later version” applies to it, you have the option of following the terms and conditions either of that published version or of any later version published by the Free Software Foundation. If the Library as you received it does not specify a version number of the GNU Lesser General Public License, you may choose any version of the GNU Lesser General Public License ever published by the Free Software Foundation. + +If the Library as you received it specifies that a proxy can decide whether future versions of the GNU Lesser General Public License shall apply, that proxy's public statement of acceptance of any version is permanent authorization for you to choose that version for the Library. + +-------------------------------------------------------------------------------- +7. Rodinia + Copyright (c)2008-2011 University of Virginia +All rights reserved. + +Redistribution and use in source and binary forms, with or without modification, are permitted without royalty fees or other restrictions, provided that the following conditions are met: + + * Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. + * Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. + * Neither the name of the University of Virginia, the Dept. of Computer Science, nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE UNIVERSITY OF VIRGINIA OR THE SOFTWARE AUTHORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + +If you use this software or a modified version of it, please cite the most relevant among the following papers: + + - M. A. Goodrum, M. J. Trotter, A. Aksel, S. T. Acton, and K. Skadron. Parallelization of Particle Filter Algorithms. In Proceedings of the 3rd Workshop on Emerging Applications and Many-core Architecture (EAMA), in conjunction with the IEEE/ACM International +Symposium on Computer Architecture (ISCA), June 2010. + + - S. Che, M. Boyer, J. Meng, D. Tarjan, J. W. Sheaffer, Sang-Ha Lee and K. Skadron. +Rodinia: A Benchmark Suite for Heterogeneous Computing. IEEE International Symposium +on Workload Characterization, Oct 2009. + +- J. Meng and K. Skadron. "Performance Modeling and Automatic Ghost Zone Optimization +for Iterative Stencil Loops on GPUs." In Proceedings of the 23rd Annual ACM International +Conference on Supercomputing (ICS), June 2009. + +- L.G. Szafaryn, K. Skadron and J. Saucerman. "Experiences Accelerating MATLAB Systems +Biology Applications." in Workshop on Biomedicine in Computing (BiC) at the International +Symposium on Computer Architecture (ISCA), June 2009. + +- M. Boyer, D. Tarjan, S. T. Acton, and K. Skadron. "Accelerating Leukocyte Tracking using CUDA: +A Case Study in Leveraging Manycore Coprocessors." In Proceedings of the International Parallel +and Distributed Processing Symposium (IPDPS), May 2009. + +- S. Che, M. Boyer, J. Meng, D. Tarjan, J. W. Sheaffer, and K. Skadron. "A Performance +Study of General Purpose Applications on Graphics Processors using CUDA" Journal of +Parallel and Distributed Computing, Elsevier, June 2008. + +-------------------------------------------------------------------------------- +Other names and brands may be claimed as the property of others. + +-------------------------------------------------------------------------------- \ No newline at end of file diff --git a/sf-filter/CMakeLists.txt b/sf-filter/CMakeLists.txt new file mode 100644 index 0000000..a870aa6 --- /dev/null +++ b/sf-filter/CMakeLists.txt @@ -0,0 +1,18 @@ +set(CMAKE_CXX_COMPILER "dpcpp") +if(WIN32) + set(CMAKE_C_COMPILER "clang-cl") +endif() + +cmake_minimum_required (VERSION 2.8) + +project(SFFilter) + +set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) +set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) +set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}) + +# copy image files +file(COPY ${CMAKE_CURRENT_SOURCE_DIR}/Images/ DESTINATION ${CMAKE_CURRENT_BINARY_DIR}/Images/) + +add_subdirectory (src) +add_subdirectory (src/Utils) \ No newline at end of file diff --git a/sf-filter/Images/cat-face.bmp b/sf-filter/Images/cat-face.bmp new file mode 100644 index 0000000..6ccd75c Binary files /dev/null and b/sf-filter/Images/cat-face.bmp differ diff --git a/sf-filter/Images/cat.bmp b/sf-filter/Images/cat.bmp new file mode 100644 index 0000000..64b4512 Binary files /dev/null and b/sf-filter/Images/cat.bmp differ diff --git a/sf-filter/Makefile b/sf-filter/Makefile new file mode 100644 index 0000000..1f63e54 --- /dev/null +++ b/sf-filter/Makefile @@ -0,0 +1,23 @@ +CXX = dpcpp +CXXFLAGS = -O2 -g -std=c++17 +INCDIR = ../image-conv/src/Utils + +BUFFER_EXE_NAME = sf-filter +BUFFER_SOURCES = src/sf-filter.cpp ../image-conv/src/Utils/utils.c ../image-conv/src/Utils/gold.c ../image-conv/src/Utils/bmp-utils.c + +all: build_buffers + +build_buffers: + $(CXX) $(CXXFLAGS) -I$(INCDIR) -o $(BUFFER_EXE_NAME) $(BUFFER_SOURCES) + +build_usm: + $(CXX) $(CXXFLAGS) -o $(USM_EXE_NAME) $(USM_SOURCES) + +run: + ./$(BUFFER_EXE_NAME) + +run_usm: + ./$(USM_EXE_NAME) + +clean: + rm -rf $(BUFFER_EXE_NAME) $(USM_EXE_NAME) diff --git a/sf-filter/Makefile.fpga b/sf-filter/Makefile.fpga new file mode 100644 index 0000000..afb046e --- /dev/null +++ b/sf-filter/Makefile.fpga @@ -0,0 +1,42 @@ +CXX := dpcpp +CXXFLAGS = -O2 -g -std=c++17 +CXXINC = -Isrc/Utils -Isrc/PipeArray + +SRC := src/sf-filter.cpp src/Utils/bmp-utils.c src/Utils/gold.c src/Utils/utils.c + +.PHONY: fpga_emu run_emu clean + +fpga_emu: sf-filter.fpga_emu + +hw: sf-filter.fpga + +report: sf-filter_report.a + +profile: sf-filter.fpga_profile + +sf-filter.fpga_profile: $(SRC) + $(CXX) $(CXXFLAGS) $(CXXINC) -fintelfpga $^ -o $@ -Xshardware -Xsprofile -Xsprofile-shared-counters -DFPGA=1 + +sf-filter.fpga_emu: $(SRC) + $(CXX) $(CXXFLAGS) $(CXXINC) -fintelfpga $^ -o $@ -DFPGA_EMULATOR=1 + +a.o: $(SRC) + $(CXX) $(CXXFLAGS) -fintelfpga -c $^ -o $@ -DFPGA=1 + +sf-filter.fpga: $(SRC) + $(CXX) $(CXXFLAGS) $(CXXINC) -fintelfpga $^ -o $@ -Xshardware -DFPGA=1 + +run_emu: sf-filter.fpga_emu + ./sf-filter.fpga_emu + +run_hw: sf-filter.fpga + ./sf-filter.fpga + +dev.o: $(SRC) + $(CXX) $(CXXFLAGS) -fintelfpga -c $^ -o $@ -DFPGA=1 + +sf-filter_report.a: dev.o + $(CXX) $(CXXFLAGS) -fintelfpga -fsycl-link $^ -o $@ -Xshardware + +clean: + rm -rf *.o *.d *.out *.mon *.emu *.aocr *.aoco *.prj *.fpga_emu *.fpga_emu_buffers sf-filter.fpga *.a diff --git a/sf-filter/README.md b/sf-filter/README.md new file mode 100644 index 0000000..6e13ac1 --- /dev/null +++ b/sf-filter/README.md @@ -0,0 +1,109 @@ +# `sf-filter` Sample + +The Sobel_Feldman operator runs two image convolution processes on an input image to detect the gradients in both the horizontal and vertical directions. This application is a very simple algorithm used to detect vertical and horizontal lines in imagery. + +| Optimized for | Description +|:--- |:--- +| OS | Linux* Ubuntu* 18.04 +| Hardware | Skylake with GEN9 or newer, Intel(R) Programmable Acceleration Card with Intel(R) Arria(R) 10 GX FPGA +| Software | Intel® oneAPI DPC++ Compiler (beta) + +## Purpose + +The Sobel-Feldman operator was designed as an introduction to programming in DPC++. Two different implementations of the Sobel-Feldman operator are designed in this application. The first uses the Image-Conv sample in this repo. The other implementation has two kernels running in parallel to compute the gradients for each pixel, and the forwards the pixel values, via pipes, to a consumer task. The consumer task calculates the Root Mean Sum for each pixel, and then sends the output image back to the host. + +## Key Implementation Details + +This 'sf-filter' example uses DPC++ buffers to store image data and a kernel function to perform convolution operation on the image data, which then gets piped to a consumer kernel function. The output of this kernel function is sent back to the host, where it gets written to a bmp file. A few utility functions, such as reading image data from bmp files or writing new image data to bmp files, are provided. + +<<<<<<< HEAD +## Future Work + +The current implementation of this application only builds for the FPGA Emulator target. The current use of the PipeArray uses about 10x more ALUTs than any FPGA on Devcloud can support, however limiting this number causes the gradient producer kernels to deadlock the program. Some work should be done in the future to implement a better PipeArray. + +## License +This code sample is licensed under MIT license. +======= +We include two implementations of image-conv kernels: buffer objects based and image objects based. Buffer objects are general memory objects that can be used to store arbitrary data structure, whereas image objects are opaque data types specifically for image data and related image processing operations. + +## License +This code sample is licensed under MIT license. +>>>>>>> origin/master + +### On a Linux* System + +**The project uses CMake. Perform the following steps to build different targets.** + +``` + mkdir build + cd build + cmake .. + make + make report + make fpga + make fpga_profile +``` +* make : by default, the emulation executables are built. +* make report : generate static report on the FPGA resource utilization of the design. +* make fpga : generate FPGA binary files for the designs. Will take a couple of hours. +* make fpga_profile : generate FPGA binary to be used in run-time profiling. Will take a couple of hours. + +*A Makefile is still maintained in the directory, however, the usage of it is disencouraged.* + +## Running the Sample + +The executables (for emulation or for FPGA hardware) can be found in the build directory. Use the file name(s) to executable the samples. For example + ``` + ./sf-filter.fpga_emu + ``` +or + ``` + ./sf-filter.fpga + ``` +or + ``` + ./sf-filter.fpga_profile + ``` + +### Application Parameters +There are no command line parameters for this sample. The input images are provided in "./Images" directory. + +### Example of Output +
+$ ./sf-filter.fpga_emu
+1 found >>
+Platform: Intel(R) FPGA Emulation Platform for OpenCL(TM)
+Device: Intel(R) FPGA Emulation Device
+2 found >>
+Platform: Intel(R) FPGA SDK for OpenCL(TM)
+Device: pac_s10 : Intel PAC Platform (pac_ec00000)
+3 found >>
+Platform: Intel(R) OpenCL
+Device: Intel(R) Xeon(R) Platinum 8256 CPU @ 3.80GHz
+4 found >>Platform: SYCL host platform
+Device: SYCL host device
+Reading input image from ./Images/cat.bmp
+offset = 1078
+width = 1080
+height = 720
+bits per pixel = 8
+imageRows=720, imageCols=1080
+Running on device: Intel(R) FPGA Emulation Device                                                                                                                                                                                            Enqueuing producer 0...
+Enqueuing producer 1...
+Enqueuing consumer...
+Horizontal kernel compute time:  93.7039 ms.
+Vertical kernel compute time:  55.4169 ms.
+Consumer kernel compute time:  9.0035 ms.
+Total compute time:  197.006 ms.
+Total Image_Conv compute time:  393.862 ms.
+Output image saved as ./Images/filtered_cat.bmp.
+
+ +## Recorded Lectures + +A series of recorded lectures are provided to introduce the important concepts in this image-conv example for FPGAs. The videos can be found at the [DPC++ Tutorial playlist](https://youtube.com/playlist?list=PLZ9YeF_1_vF8RqYPNpHToklJcDRoVocU4) on Youtube and are linked individually below. + +[Introduction to Image Convolution](https://youtu.be/O_-sXNy23mw) + +[Image Objects in DPC++](https://youtu.be/Kb46lMLsve0) + diff --git a/sf-filter/src/CMakeLists.txt b/sf-filter/src/CMakeLists.txt new file mode 100644 index 0000000..f974d8e --- /dev/null +++ b/sf-filter/src/CMakeLists.txt @@ -0,0 +1,131 @@ +set(SOURCE_FILE sf-filter.cpp) +set(TARGET_NAME sf-filter) +set(EMULATOR_TARGET ${TARGET_NAME}.fpga_emu) +set(FPGA_TARGET ${TARGET_NAME}.fpga) +set(FPGA_PROFILE_TARGET ${TARGET_NAME}.fpga_profile) +set(SOURCE_FILE_IMAGE sf-filter.cpp) + +# FPGA board selection +set(A10_PAC_BOARD_NAME "intel_a10gx_pac:pac_a10") +set(S10_PAC_BOARD_NAME "intel_s10sx_pac:pac_s10") +set(SELECTED_BOARD ${A10_PAC_BOARD_NAME}) +if (NOT DEFINED FPGA_BOARD) + message(STATUS "\tFPGA_BOARD was not specified. Configuring the design to run on the Intel(R) Programmable Acceleration Card (PAC) with Intel Arria(R) 10 GX FPGA. Please refer to the README for information on board selection.") +elseif(FPGA_BOARD STREQUAL ${A10_PAC_BOARD_NAME}) + message(STATUS "\tConfiguring the design to run on the Intel(R) Programmable Acceleration Card (PAC) with Intel Arria(R) 10 GX FPGA.") +elseif(FPGA_BOARD STREQUAL ${S10_PAC_BOARD_NAME}) + message(STATUS "\tConfiguring the design to run on the Intel(R) Programmable Acceleration Card (PAC) D5005 (with Intel Stratix(R) 10 SX FPGA).") + set(SELECTED_BOARD ${S10_PAC_BOARD_NAME}) +else() + message(STATUS "\tAn invalid board name was passed in using the FPGA_BOARD flag. Configuring the design to run on the Intel(R) Programmable Acceleration Card (PAC) with Intel Arria(R) 10 GX FPGA. Please refer to the README for the list of valid board names.") +endif() + +# Flags +set(EMULATOR_COMPILE_FLAGS "-fintelfpga -DFPGA_EMULATOR") +set(EMULATOR_LINK_FLAGS "-fintelfpga") +set(HARDWARE_COMPILE_FLAGS "-fintelfpga -DFPGA") +set(HARDWARE_LINK_FLAGS "-fintelfpga -Xshardware -Xsboard=${SELECTED_BOARD} ${USER_HARDWARE_FLAGS}") +# use cmake -D USER_HARDWARE_FLAGS= to set extra flags for FPGA backend compilation +set(HARDWARE_PROFILE_COMPILE_FLAGS "-fintelfpga -DFPGA_PROFILE") +set(HARDWARE_PROFILE_LINK_FLAGS "-fintelfpga -Xshardware -Xsprofile -Xsprofile-shared-counters -Xsboard=${SELECTED_BOARD} ${USER_HARDWARE_FLAGS}") + + +# FPGA emulator +if(WIN32) + set(WIN_EMULATOR_TARGET ${EMULATOR_TARGET}.exe) + add_custom_target(fpga_emu DEPENDS ${WIN_EMULATOR_TARGET}) + separate_arguments(WIN_EMULATOR_COMPILE_FLAGS WINDOWS_COMMAND "${EMULATOR_COMPILE_FLAGS}") + add_custom_command(OUTPUT ${WIN_EMULATOR_TARGET} + COMMAND ${CMAKE_CXX_COMPILER} /EHsc ${WIN_EMULATOR_COMPILE_FLAGS} ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_FILE} -o ${CMAKE_BINARY_DIR}/${WIN_EMULATOR_TARGET} + DEPENDS ${SOURCE_FILE}) +else() + add_executable(${EMULATOR_TARGET} ${SOURCE_FILE}) + add_custom_target(fpga_emu DEPENDS ${EMULATOR_TARGET}) + set_target_properties(${EMULATOR_TARGET} PROPERTIES COMPILE_FLAGS ${EMULATOR_COMPILE_FLAGS}) + set_target_properties(${EMULATOR_TARGET} PROPERTIES LINK_FLAGS ${EMULATOR_LINK_FLAGS}) + # to link with functions in Utils library (definded in Utils/ folder) + target_link_libraries (${EMULATOR_TARGET} LINK_PUBLIC Utils) +endif() + +# CPU or GPU that supports "image" APIs +if(WIN32) + add_custom_target(nofpga COMMAND echo "Non-FPGA emulation on Windows. not tested. See README for details.") + set(WIN_EMULATOR_TARGET ${EMULATOR_TARGET}.exe) + add_custom_target(nofpga DEPENDS ${WIN_EMULATOR_TARGET}) + separate_arguments(WIN_EMULATOR_COMPILE_FLAGS WINDOWS_COMMAND "${EMULATOR_COMPILE_FLAGS}") + add_custom_command(OUTPUT ${WIN_EMULATOR_TARGET} + COMMAND ${CMAKE_CXX_COMPILER} /EHsc ${WIN_EMULATOR_COMPILE_FLAGS} ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_FILE} -o ${CMAKE_BINARY_DIR}/${WIN_EMULATOR_TARGET} + DEPENDS ${SOURCE_FILE}) +else() + add_executable(${TARGET_NAME} ${SOURCE_FILE_IMAGE}) + add_custom_target(nofpga DEPENDS ${TARGET_NAME}) + #set_target_properties(${EMULATOR_TARGET} PROPERTIES COMPILE_FLAGS ${EMULATOR_COMPILE_FLAGS}) + #set_target_properties(${EMULATOR_TARGET} PROPERTIES LINK_FLAGS ${EMULATOR_LINK_FLAGS}) + # to link with functions in Utils library (definded in Utils/ folder) + target_link_libraries (${TARGET_NAME} LINK_PUBLIC Utils) +endif() + +# FPGA hardware +if(WIN32) + add_custom_target(fpga COMMAND echo "An FPGA hardware target is not provided on Windows. See README for details.") +else() + add_executable(${FPGA_TARGET} EXCLUDE_FROM_ALL ${SOURCE_FILE}) + add_custom_target(fpga DEPENDS ${FPGA_TARGET}) + target_include_directories(${FPGA_TARGET} PUBLIC + $) + set_target_properties(${FPGA_TARGET} PROPERTIES COMPILE_FLAGS ${HARDWARE_COMPILE_FLAGS}) + set_target_properties(${FPGA_TARGET} PROPERTIES LINK_FLAGS ${HARDWARE_LINK_FLAGS}) + # to link with functions in Utils library (definded in Utils/ folder) + target_link_libraries (${FPGA_TARGET} LINK_PUBLIC Utils) +endif() + +# Generate report +if(WIN32) + set(DEVICE_OBJ_FILE ${TARGET_NAME}_report.a) + add_custom_target(report DEPENDS ${DEVICE_OBJ_FILE}) + separate_arguments(HARDWARE_LINK_FLAGS_LIST WINDOWS_COMMAND "${HARDWARE_LINK_FLAGS}") + add_custom_command(OUTPUT ${DEVICE_OBJ_FILE} + COMMAND ${CMAKE_CXX_COMPILER} /EHsc ${HARDWARE_LINK_FLAGS_LIST} -fsycl-link ${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_FILE} -o ${CMAKE_BINARY_DIR}/${DEVICE_OBJ_FILE} + DEPENDS ${SOURCE_FILE}) +else() +#[[ foreach( testsourcefile ${SOURCE_FILES} ) + # replace file suffix to create executable file name + string( REPLACE ".cpp" "_report.a" devobjfile ${testsourcefile} ) + get_filename_component(barename ${testsourcefile} NAME_WE) + add_custom_target(${barename}.report DEPENDS ${devobjfile}) + list(APPEND reportlist ${barename}.report) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${testsourcefile} ${testsourcefile} COPYONLY) + separate_arguments(HARDWARE_LINK_FLAGS_LIST UNIX_COMMAND "${HARDWARE_LINK_FLAGS}") + separate_arguments(CMAKE_CXX_FLAGS_LIST UNIX_COMMAND "${CMAKE_CXX_FLAGS}") + add_custom_command(OUTPUT ${devobjfile} + COMMAND ${CMAKE_CXX_COMPILER} ${CMAKE_CXX_FLAGS_LIST} ${HARDWARE_LINK_FLAGS_LIST} -fsycl-link ${testsourcefile} -o ${CMAKE_BINARY_DIR}/${devobjfile} + DEPENDS ${testsourcefile}) + endforeach( testsourcefile ${SOURCE_FILES} ) + add_custom_target(report DEPENDS ${reportlist}) +]] + + set(DEVICE_OBJ_FILE ${TARGET_NAME}_report.a) + add_custom_target(report DEPENDS ${DEVICE_OBJ_FILE}) + configure_file(${CMAKE_CURRENT_SOURCE_DIR}/${SOURCE_FILE} ${SOURCE_FILE} COPYONLY) + separate_arguments(HARDWARE_LINK_FLAGS_LIST UNIX_COMMAND "${HARDWARE_LINK_FLAGS}") + separate_arguments(CMAKE_CXX_FLAGS_LIST UNIX_COMMAND "${CMAKE_CXX_FLAGS}") + list(APPEND CMAKE_CXX_FLAGS_LIST -I${CMAKE_CURRENT_SOURCE_DIR} -I${CMAKE_CURRENT_SOURCE_DIR}/Utils/) + add_custom_command(OUTPUT ${DEVICE_OBJ_FILE} + COMMAND ${CMAKE_CXX_COMPILER} ${CMAKE_CXX_FLAGS_LIST} ${HARDWARE_LINK_FLAGS_LIST} -fsycl-link ${SOURCE_FILE} -o ${CMAKE_BINARY_DIR}/${DEVICE_OBJ_FILE} + DEPENDS ${SOURCE_FILE}) +endif() + +# FPGA hardware profiling +if(WIN32) + add_custom_target(fpga_profile COMMAND echo "An FPGA hardware target is not provided on Windows. See README for details.") +else() + add_executable(${FPGA_PROFILE_TARGET} EXCLUDE_FROM_ALL ${SOURCE_FILE}) + add_custom_target(fpga_profile DEPENDS ${FPGA_PROFILE_TARGET}) + target_include_directories(${FPGA_PROFILE_TARGET} PUBLIC + $) + set_target_properties(${FPGA_PROFILE_TARGET} PROPERTIES COMPILE_FLAGS ${HARDWARE_PROFILE_COMPILE_FLAGS}) + set_target_properties(${FPGA_PROFILE_TARGET} PROPERTIES LINK_FLAGS ${HARDWARE_PROFILE_LINK_FLAGS}) + # to link with functions in Utils library (definded in Utils/ folder) + target_link_libraries (${FPGA_PROFILE_TARGET} LINK_PUBLIC Utils) +endif() + diff --git a/sf-filter/src/Utils b/sf-filter/src/Utils new file mode 120000 index 0000000..1e196aa --- /dev/null +++ b/sf-filter/src/Utils @@ -0,0 +1 @@ +../../image-conv/src/Utils \ No newline at end of file diff --git a/sf-filter/src/pipe_array.hpp b/sf-filter/src/pipe_array.hpp new file mode 100644 index 0000000..0f1ca1f --- /dev/null +++ b/sf-filter/src/pipe_array.hpp @@ -0,0 +1,33 @@ +//============================================================== +// Copyright Intel Corporation +// +// SPDX-License-Identifier: MIT +// ============================================================= +#include +#include +#include + +#include "pipe_array_internal.hpp" + +template +struct PipeArray { + PipeArray() = delete; + + template + struct StructId; + + template + struct VerifyIndices { + static_assert(sizeof...(idxs) == sizeof...(dims), + "Indexing into a PipeArray requires as many indices as " + "dimensions of the PipeArray."); + static_assert(VerifierDimLayer::template VerifierIdxLayer< + idxs...>::IsValid(), + "Index out of bounds"); + using VerifiedPipe = + cl::sycl::INTEL::pipe, BaseTy, depth>; + }; + + template + using PipeAt = typename VerifyIndices::VerifiedPipe; +}; diff --git a/sf-filter/src/pipe_array_internal.hpp b/sf-filter/src/pipe_array_internal.hpp new file mode 100644 index 0000000..1b62f66 --- /dev/null +++ b/sf-filter/src/pipe_array_internal.hpp @@ -0,0 +1,26 @@ +//============================================================== +// Copyright Intel Corporation +// +// SPDX-License-Identifier: MIT +// ============================================================= + +namespace { +template +struct VerifierDimLayer { + template + struct VerifierIdxLayer { + static constexpr bool IsValid() { + return idx1 < dim1 && + (VerifierDimLayer::template VerifierIdxLayer< + idxs...>::IsValid()); + } + }; +}; +template +struct VerifierDimLayer { + template + struct VerifierIdxLayer { + static constexpr bool IsValid() { return idx < dim; } + }; +}; +} // namespace diff --git a/sf-filter/src/sf-filter.cpp b/sf-filter/src/sf-filter.cpp new file mode 100644 index 0000000..eec0d20 --- /dev/null +++ b/sf-filter/src/sf-filter.cpp @@ -0,0 +1,433 @@ +//*************************************************************************************** +// DPC++ Example +// +// Sobel-Feldman Filter with DPC++ +// +// Author: Matthew Cloutier +// +//*************************************************************************************** +#include +#include +#include +#include +#include "dpc_common.hpp" +#if FPGA || FPGA_EMULATOR || FPGA_PROFILE +#include +#endif + +using namespace sycl; + +// useful headers for image processing +#include "bmp-utils.h" + +// pipe_array headers +#include "pipe_array.hpp" + +// Debug flag for print statements +static const bool debug = true; + +static const int sobelFilterWidth = 3; +static float verticalSobelFilter[9] = { + 3.0f, 0.0f, -3.0f, + 10.0f, 0.0f, -10.0f, + 3.0f, 0.0f, -3.0f +}; + +static float horizontalSobelFilter[9] = { + 3.0f, 10.0f, 3.0f, + 0.0f, 0.0f, 0.0f, + -3.0f, -10.0f, -3.0f +}; + +// +// Image Info +// +static const char* inputImagePath = "./Images/cat.bmp"; +static const char* output_filename = "./Images/filtered_cat.bmp"; + +#define IMAGE_SIZE (720*1080) +constexpr size_t array_size = IMAGE_SIZE; +typedef std::array FloatArray; + +// 2-D pipe matrix. One pipe for each pixel, for each image filter. +constexpr size_t kDepth = 2; +constexpr size_t kNumProducers = 2; +constexpr size_t kNumRows = 720; +constexpr size_t kNumCols = 1080; + +using ProducerPipeArray = PipeArray; // number of pipes to create + +template class ProducerKernel; +class ConsumerKernel; + +template +sycl::event Producer(queue &q, float *image_in, + float *filter_in, const size_t FilterWidth, + const size_t ImageRows, const size_t ImageCols) +{ + std::cout << "Enqueuing producer " << producer_id << "..." << std::endl; + + buffer image_in_buf(image_in, range<1>(ImageRows * ImageCols)); + buffer filter_buf(filter_in, range<1>(FilterWidth*FilterWidth)); + + // + // Compute the filter width (intentionally truncate) + // + int halfFilterWidth = (int)FilterWidth/2; + + auto e = q.submit([&](handler &h) + { + auto sourcePtr = image_in_buf.get_access(h); + auto filterPtr = filter_buf.get_access(h); + + // + // Apply filter to each pixel in the image. + // Need to do this as single task to ensure the order + // of pixels being processed by the consumer kernel; + // + h.single_task>([=]() + { + for (int row = 0; row < ImageRows; row++) + { + for (int col = 0; col < ImageCols; col++) + { + // Each work item iterates around its local are based on the size of the filter. + float sum = 0.0f; + + // + // Apply the filter to the pixel neigborhood + // + for (int k = -halfFilterWidth; k <= halfFilterWidth; k++) + { + for (int l = -halfFilterWidth; l <= halfFilterWidth; l++) + { + // Indices used to access the image + int r = row + k; + int c = col + l; + + // Handle out of bounds locations by clamping to the border pixel + r = (r < 0) ? 0 : r; + c = (c < 0) ? 0 : c; + r = (r >= ImageRows) ? ImageRows - 1 : r; + c = (c >= ImageCols) ? ImageCols - 1 : c; + + sum += sourcePtr[r * ImageCols + c] * + filterPtr[(k + halfFilterWidth) * FilterWidth + (l + halfFilterWidth)]; + + } + } + // + // Write the value to the pipe for the given filter, row, col values + // + ProducerPipeArray::PipeAt::write(sum); + + } + } + }); + }); + + // + // Get the elapsed kernel time, mostly just to see difference between + // single_task and parallel_for execution schema + // + + return e; + //return (end_time_ns - start_time_ns); +} + +sycl::event Consumer(queue &q, float *image_out, const size_t ImageRows, const size_t ImageCols) +{ + std::cout << "Enqueuing consumer..." << std::endl; + + // + // Create buffers for the input and output data + // + buffer image_out_buf(image_out, range<1>(ImageRows * ImageCols)); + + auto e = q.submit([&](handler &h) + { + auto destPtr = image_out_buf.get_access(h); + + // + // Loop through each pixel and combine the results from the two producer kernels. + // Must be done as a single task to ensure the image is processed in the correct order. + // + h.single_task([=]() + { + for (int row= 0; row < ImageRows; row++) + { + for (int col = 0; col < ImageCols; col++) + { + // + // Read the results from the two data pipes and calculate the RMS + // + float result_0 = ProducerPipeArray::PipeAt<0>::read(); + float result_1 = ProducerPipeArray::PipeAt<1>::read(); + + float result = sqrt((result_0 * result_0) + (result_1 * result_1)); + + destPtr[row * ImageCols + col] = result; + } + } + }); + }); + + // + // Get the elapsed kernel time, mostly just to see difference between + // single_task and parallel_for execution schema + // + + return e; +} + +// ************************************************************************************** +// Image Convolution in DPC++ on device: +// Copied from image-conv.cpp and modified for specific Sobel-Feldman filtering +// ************************************************************************************** +void ImageConv(queue &q, float *image_in, float *image_out, + float *filter_in, const size_t FilterWidth, + const size_t ImageRows, const size_t ImageCols) +{ + // + // Create buffers for the input and output data + // + buffer image_in_buf(image_in, range<1>(ImageRows * ImageCols)); + buffer image_out_buf(image_out, range<1>(ImageRows * ImageCols)); + + // + // Create the range object for the pixel data + // + range<2> num_items {ImageRows, ImageCols}; + + // + // Create buffers that hold the filter shared between the host and the devices. + // + buffer filter_buf(filter_in, range<1>(FilterWidth*FilterWidth)); + + // + // Compute the filter width (intentionally truncate) + // + int halfFilterWidth = (int)FilterWidth/2; + + // + // Submit a command group to the queue by a lambda function that contains the data access + // permission and device computation (kernel) + // + q.submit([&](handler &h) + { + // + // Create an accessor to the buffers with access permission. + // + auto sourcePtr = image_in_buf.get_access(h); + auto destPtr = image_out_buf.get_access(h); + auto filterPtr = filter_buf.get_access(h); + + // + // Use parallel_for to run the image convolution in parallel on device. + // This executes the kernel. + // 1st parameter is the number of work items. + // 2nd parameter is the kernel, the lambda to specify what to do per work item. + // The parameter of the lambda is the work item id. + h.parallel_for(num_items, [=](id<2> item) + { + // + // Get the cow and col of the pixel assigned to this work item + // + int row = item[0]; + int col = item[1]; + + // Each work item iterates around its local are based on the size of the filter. + float sum = 0.0f; + + // + // Apply the filter to the pixel neigborhood + // + for (int k = -halfFilterWidth; k <= halfFilterWidth; k++) + { + for (int l = -halfFilterWidth; l <= halfFilterWidth; l++) + { + // Indices used to access the image + int r = row + k; + int c = col + l; + + // Handle out of bounds locations by clamping to the border pixel + r = (r < 0) ? 0 : r; + c = (c < 0) ? 0 : c; + r = (r >= ImageRows) ? ImageRows - 1 : r; + c = (c >= ImageCols) ? ImageCols - 1 : c; + + sum += sourcePtr[r * ImageCols + c] * + filterPtr[(k + halfFilterWidth) * FilterWidth + (l + halfFilterWidth)]; + + } + } + // + // Save off the new pixel value + // + destPtr[row * ImageCols + col] = sum; + + }); + }); + +} + +int main() +{ + // Create device selector for the device of your interest. + #if FPGA_EMULATOR + // DPC++ extension: FPGA emulator selector on systems without FPGA card. + INTEL::fpga_emulator_selector d_selector; + #elif FPGA || FPGA_PROFILE + // DPC++ extension: FPGA selector on systems with FPGA card. + INTEL::fpga_selector d_selector; + #else + // The default device selector will select the most performant device. + default_selector d_selector; + #endif + + float *hInputImage; + float *hOutputImage; + float *vOutputImage; + float *outputImage; + + int imageRows; + int imageCols; + + float *filter; + + #ifndef FPGA_PROFILE + // Query some platform information + unsigned number = 0; + auto myPlatforms = platform::get_platforms(); + + //loop through the platforms to poke + for (auto &onePlatform: myPlatforms) + { + std::cout << ++number << " found >>" << std::endl; + std::cout << "Platform: " << onePlatform.get_info() << std::endl; + + // Loop through devices + auto myDevices = onePlatform.get_devices(); + for (auto &oneDevice : myDevices) + { + std::cout << "Device: " << oneDevice.get_info() << std::endl; + } + } + + std::cout << std::endl; + #endif + + // + // Read in the image + // + hInputImage = readBmpFloat(inputImagePath, &imageRows, &imageCols); + if (debug) + { + printf("imageRows=%d, imageCols=%d\n", imageRows, imageCols); + } + + // + // Allocate space for the resulting images + // + hOutputImage = (float*) malloc(imageRows * imageCols * sizeof(float)); + vOutputImage = (float*) malloc(imageRows * imageCols * sizeof(float)); + outputImage = (float*) malloc(imageRows * imageCols * sizeof(float)); + + // + // Start the timer + // + dpc_common::TimeInterval exec_time; + double total_pipe_path_time_s = 0; + + try + { + auto prop_list = property_list{property::queue::enable_profiling()}; + + queue q(d_selector, dpc_common::exception_handler, prop_list); + + // + // Print device information used in the kernel + // + std::cout << "Running on device: " << + q.get_device().get_info() << std::endl; + + // + // Enqueue producers + // + double start_time = exec_time.Elapsed(); + + auto horzEvent = Producer<0>(q, hInputImage, horizontalSobelFilter, sobelFilterWidth, imageRows, imageCols); + + auto vertEvent = Producer<1>(q, hInputImage, verticalSobelFilter, sobelFilterWidth, imageRows, imageCols); + + // + // Enqueue Consumer + // + auto consEvent = Consumer(q, outputImage, imageRows, imageCols); + + q.wait(); + + double total_pipe_path_time_s = exec_time.Elapsed() - start_time; + + + auto horzKernelTime = (horzEvent.get_profiling_info() - + horzEvent.get_profiling_info()); + auto vertKernelTime = (vertEvent.get_profiling_info() - + vertEvent.get_profiling_info()); + auto consKernelTime = (consEvent.get_profiling_info() - + consEvent.get_profiling_info()); + + std::cout << "Horizontal kernel compute time: " << horzKernelTime * 1e-6 << " ms." << std::endl; + std::cout << "Vertical kernel compute time: " << vertKernelTime * 1e-6 << " ms." << std::endl; + std::cout << "Consumer kernel compute time: " << consKernelTime * 1e-6 << " ms." << std::endl; + std::cout << "Total compute time: " << total_pipe_path_time_s * 1e3 << " ms." << std::endl; + + + // + // Run the horizontal line filter, then the vertical line filter + // + filter = horizontalSobelFilter; + ImageConv(q, hInputImage, hOutputImage, filter, sobelFilterWidth, imageRows, imageCols); + + filter = verticalSobelFilter; + ImageConv(q, hInputImage, vOutputImage, filter, sobelFilterWidth, imageRows, imageCols); + } + catch(const std::exception& e) + { + std::cerr << "Caught the following error executing ImageConv:" << std::endl; + std::cerr << e.what() << std::endl; + std::terminate(); + } + + // + // No errors runing kernel, so combine the two filtered images + // + for (int r = 0; r < imageRows; r++) + { + for (int c = 0; c < imageCols; c++) + { + outputImage[r * imageCols + c] = std::sqrt( + std::pow(hOutputImage[r * imageCols + c], 2) + + std::pow(vOutputImage[r * imageCols + c], 2)); + } + } + + + // + // Stop the timer and report execution times + // + double total_img_conv_time_s = exec_time.Elapsed() - total_pipe_path_time_s; + std::cout << "Total Image_Conv compute time: " << total_img_conv_time_s * 1e3 << " ms." << std::endl; + + // + // Save the output bmp + // + printf("Output image saved as %s.\n", output_filename); + writeBmpFloat(outputImage, output_filename, imageRows, imageCols, inputImagePath); + + return 0; + +} \ No newline at end of file