diff --git a/Makefile b/Makefile index 44b98aa9b998..7f93df612170 100644 --- a/Makefile +++ b/Makefile @@ -2399,7 +2399,6 @@ install: $(LIB_DIR)/libHalide.a $(BIN_DIR)/libHalide.$(SHARED_EXT) $(INCLUDE_DIR cp $(ROOT_DIR)/tutorial/figures/*.jpg $(PREFIX)/share/halide/tutorial/figures cp $(ROOT_DIR)/tutorial/figures/*.mp4 $(PREFIX)/share/halide/tutorial/figures cp $(ROOT_DIR)/tutorial/*.cpp $(PREFIX)/share/halide/tutorial - cp $(ROOT_DIR)/tutorial/*.h $(PREFIX)/share/halide/tutorial cp $(ROOT_DIR)/tutorial/*.sh $(PREFIX)/share/halide/tutorial cp $(ROOT_DIR)/tools/GenGen.cpp $(PREFIX)/share/halide/tools cp $(ROOT_DIR)/tools/RunGen.h $(PREFIX)/share/halide/tools @@ -2477,7 +2476,6 @@ $(DISTRIB_DIR)/lib/libHalide.$(SHARED_EXT): \ cp $(ROOT_DIR)/tutorial/figures/*.jpg $(DISTRIB_DIR)/tutorial/figures cp $(ROOT_DIR)/tutorial/figures/*.mp4 $(DISTRIB_DIR)/tutorial/figures cp $(ROOT_DIR)/tutorial/*.cpp $(DISTRIB_DIR)/tutorial - cp $(ROOT_DIR)/tutorial/*.h $(DISTRIB_DIR)/tutorial cp $(ROOT_DIR)/tutorial/*.sh $(DISTRIB_DIR)/tutorial cp $(ROOT_DIR)/tools/GenGen.cpp $(DISTRIB_DIR)/tools cp $(ROOT_DIR)/tools/RunGen.h $(DISTRIB_DIR)/tools diff --git a/python_bindings/tutorial/lesson_01_basics.py b/python_bindings/tutorial/lesson_01_basics.py index f82baf029716..26529b3bcffe 100755 --- a/python_bindings/tutorial/lesson_01_basics.py +++ b/python_bindings/tutorial/lesson_01_basics.py @@ -4,10 +4,6 @@ # This lesson demonstrates basic usage of Halide as a JIT compiler for imaging. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_01_basics -# in a shell with the current directory at python_bindings/ - import halide as hl diff --git a/python_bindings/tutorial/lesson_02_input_image.py b/python_bindings/tutorial/lesson_02_input_image.py index b80e9befa259..bffd53743975 100644 --- a/python_bindings/tutorial/lesson_02_input_image.py +++ b/python_bindings/tutorial/lesson_02_input_image.py @@ -2,10 +2,6 @@ # This lesson demonstrates how to pass in input images. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_02_input_image -# in a shell with the current directory at python_bindings/ - import halide as hl import numpy as np import halide.imageio diff --git a/python_bindings/tutorial/lesson_03_debugging_1.py b/python_bindings/tutorial/lesson_03_debugging_1.py index 0fe6ee29b862..b1ee563ffbdf 100755 --- a/python_bindings/tutorial/lesson_03_debugging_1.py +++ b/python_bindings/tutorial/lesson_03_debugging_1.py @@ -5,10 +5,6 @@ # This lesson demonstrates how to inspect what the Halide compiler is # producing. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_03_debugging_1 -# in a shell with the current directory at python_bindings/ - import halide as hl diff --git a/python_bindings/tutorial/lesson_04_debugging_2.py b/python_bindings/tutorial/lesson_04_debugging_2.py index 6e9384d84f1a..a85ceae288a1 100755 --- a/python_bindings/tutorial/lesson_04_debugging_2.py +++ b/python_bindings/tutorial/lesson_04_debugging_2.py @@ -4,10 +4,6 @@ # This lesson demonstrates how to follow what Halide is doing at runtime. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_04_debugging_2 -# in a shell with the current directory at python_bindings/ - import halide as hl @@ -44,16 +40,15 @@ def main(): # that describes the algorithm. # Now we tell Halide to use a parallel for loop over the y - # coordinate. On linux we run this using a thread pool and a task - # queue. On os x we call into grand central dispatch, which does - # the same thing for us. + # coordinate. Halide's runtime maintains its own pool of worker + # threads and a task queue. parallel_gradient.parallel(y) # This time the printfs should come out of order, because each # scanline is potentially being processed in a different # thread. The number of threads should adapt to your system, but - # on linux you can control it manually using the environment - # variable HL_NUMTHREADS. + # you can control it manually using the environment variable + # HL_NUM_THREADS. print("\nEvaluating parallel_gradient") parallel_gradient.realize([8, 8]) diff --git a/python_bindings/tutorial/lesson_05_scheduling_1.py b/python_bindings/tutorial/lesson_05_scheduling_1.py index 5c6b6e732200..b642e4ed2b46 100755 --- a/python_bindings/tutorial/lesson_05_scheduling_1.py +++ b/python_bindings/tutorial/lesson_05_scheduling_1.py @@ -6,9 +6,6 @@ # evaluate pixels in a hl.Func, including vectorization, # parallelization, unrolling, and tiling. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_05_scheduling_1 -# in a shell with the current directory at python_bindings/ import halide as hl diff --git a/python_bindings/tutorial/lesson_06_realizing_over_shifted_domains.py b/python_bindings/tutorial/lesson_06_realizing_over_shifted_domains.py index 80f7e195e16a..5e2b0359b8c3 100755 --- a/python_bindings/tutorial/lesson_06_realizing_over_shifted_domains.py +++ b/python_bindings/tutorial/lesson_06_realizing_over_shifted_domains.py @@ -5,10 +5,6 @@ # This lesson demonstrates how to evaluate a hl.Func over a domain that # does not start at (0, 0). -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_06_realizing_over_shifted_domains -# in a shell with the current directory at python_bindings/ - import halide as hl @@ -78,7 +74,8 @@ def main(): # fact read out-of-bounds and probably crash. # What if we want to evaluate our hl.Func over some region that - # isn't rectangular? Too bad. Halide only does rectangles :) + # isn't rectangular? Unfortunately, we can't. Halide only supports + # rectangular domains. print("Success!") return 0 diff --git a/python_bindings/tutorial/lesson_07_multi_stage_pipelines.py b/python_bindings/tutorial/lesson_07_multi_stage_pipelines.py index 1e95813a2832..1b37f8e2fc9e 100755 --- a/python_bindings/tutorial/lesson_07_multi_stage_pipelines.py +++ b/python_bindings/tutorial/lesson_07_multi_stage_pipelines.py @@ -4,10 +4,6 @@ # This lesson demonstrates how express multi-stage pipelines. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_07_multi_stage_pipelines -# in a shell with the current directory at python_bindings/ - import halide as hl import halide.imageio @@ -92,7 +88,7 @@ def main(): print("Created blurry_parrot_1.png") # This is usually the fastest way to deal with boundaries: - # don't write code that reads out of bounds :) The more + # avoid reading out of bounds in the first place. The more # general solution is our next example. # The same pipeline, with a boundary condition on the input. diff --git a/python_bindings/tutorial/lesson_08_scheduling_2.py b/python_bindings/tutorial/lesson_08_scheduling_2.py index 45565698dd60..17861768ca66 100755 --- a/python_bindings/tutorial/lesson_08_scheduling_2.py +++ b/python_bindings/tutorial/lesson_08_scheduling_2.py @@ -4,10 +4,6 @@ # This lesson demonstrates how schedule multi-stage pipelines. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_08_scheduling_2 -# in a shell with the current directory at python_bindings/ - import halide as hl import numpy as np import math @@ -441,7 +437,7 @@ def main(): # iterations are lying around for us to reuse. This assumes that # previous values of x or y happened earlier in time and have # finished. This is not true if you parallelize or vectorize - # either loop. Darn. If you parallelize, Halide won't inject the + # either loop. If you parallelize, Halide won't inject the # optimizations that skip work already done if there's a parallel # loop in between the store_at level and the compute_at level, # and won't fold the storage down into a circular buffer either, @@ -659,8 +655,6 @@ def main(): consumer.print_loop_nest() print() - # Look on my code, ye mighty, and despair! - # Let's check the C result against the Halide result. Doing # this I found several bugs in my C implementation, which # should tell you something. @@ -675,7 +669,7 @@ def main(): # This stuff is hard. We ended up in a three-way trade-off # between memory bandwidth, redundant work, and # parallelism. Halide can't make the correct choice for you - # automatically (sorry). Instead it tries to make it easier for + # automatically. Instead it tries to make it easier for # you to explore various options, without messing up your # program. In fact, Halide promises that scheduling calls like # compute_root won't change the meaning of your algorithm -- you diff --git a/python_bindings/tutorial/lesson_09_update_definitions.py b/python_bindings/tutorial/lesson_09_update_definitions.py index 56801c858f22..ba10a1328fde 100755 --- a/python_bindings/tutorial/lesson_09_update_definitions.py +++ b/python_bindings/tutorial/lesson_09_update_definitions.py @@ -5,9 +5,6 @@ # This lesson demonstrates how to define a hl.Func in multiple passes, # including scattering. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_09_update_definitions -# in a shell with the current directory at python_bindings/ import halide as hl import halide.imageio diff --git a/python_bindings/tutorial/lesson_10_aot_compilation_generate.py b/python_bindings/tutorial/lesson_10_aot_compilation_generate.py index 81ba0ec275d2..91a38a76efbb 100755 --- a/python_bindings/tutorial/lesson_10_aot_compilation_generate.py +++ b/python_bindings/tutorial/lesson_10_aot_compilation_generate.py @@ -11,15 +11,6 @@ # uses that object file to actually run the pipeline. This means that # compiling this code is a multi-step process. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_10_aot_compilation_generate -# in a shell with the current directory at python_bindings/ - -# This will generate a file lesson_10_halide.py.cpp that still needs -# to be compiled. Use -# make test_tutorial_lesson_10_aot_compilation_run -# to generate and run a Python module called "lesson_10_halide". - # The benefits of this approach are that the final program: # - Doesn't do any jit compilation at runtime, so it's fast. # - Doesn't depend on libHalide at all, so it's a small, easy-to-deploy binary. diff --git a/python_bindings/tutorial/lesson_11_cross_compilation.py b/python_bindings/tutorial/lesson_11_cross_compilation.py index aae8fd8bc004..5dc598b98e7f 100755 --- a/python_bindings/tutorial/lesson_11_cross_compilation.py +++ b/python_bindings/tutorial/lesson_11_cross_compilation.py @@ -4,10 +4,6 @@ # This lesson demonstrates how to use Halide as a cross-compiler. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_11_cross_compilation -# in a shell with the current directory at python_bindings/ - import halide as hl from struct import unpack diff --git a/python_bindings/tutorial/lesson_12_using_the_gpu.py b/python_bindings/tutorial/lesson_12_using_the_gpu.py index e6e30ab55255..f743dd769e68 100755 --- a/python_bindings/tutorial/lesson_12_using_the_gpu.py +++ b/python_bindings/tutorial/lesson_12_using_the_gpu.py @@ -4,10 +4,6 @@ # This lesson demonstrates how to use Halide to run code on a GPU. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_12_using_the_gpu -# in a shell with the current directory at python_bindings/ - import halide as hl import halide.imageio @@ -285,8 +281,6 @@ def find_gpu_target(): features_to_try.append(hl.TargetFeature.D3D12Compute) features_to_try.append(hl.TargetFeature.OpenCL) elif target.os == hl.TargetOS.OSX: - # OS X doesn't update its OpenCL drivers, so they tend to be broken. - # CUDA would also be a fine choice on machines with NVidia GPUs. features_to_try.append(hl.TargetFeature.Metal) else: features_to_try.append(hl.TargetFeature.OpenCL) diff --git a/python_bindings/tutorial/lesson_13_tuples.py b/python_bindings/tutorial/lesson_13_tuples.py index 89b239b1aa79..fad5c4616c5b 100755 --- a/python_bindings/tutorial/lesson_13_tuples.py +++ b/python_bindings/tutorial/lesson_13_tuples.py @@ -5,10 +5,6 @@ # This lesson describes how to write Funcs that evaluate to multiple # values. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_13_tuples -# in a shell with the current directory at python_bindings/ - import math import halide as hl diff --git a/python_bindings/tutorial/lesson_14_types.py b/python_bindings/tutorial/lesson_14_types.py index b5ea52df021e..1f3570c9e9d9 100755 --- a/python_bindings/tutorial/lesson_14_types.py +++ b/python_bindings/tutorial/lesson_14_types.py @@ -4,11 +4,6 @@ # This lesson more precisely describes Halide's type system. -# This lesson can be built by invoking the command: -# make test_tutorial_lesson_14_types -# in a shell with the current directory at python_bindings/ - - import halide as hl diff --git a/tools/halide_thread_pool.h b/tools/halide_thread_pool.h index 739b1444b66c..c4be162abb56 100644 --- a/tools/halide_thread_pool.h +++ b/tools/halide_thread_pool.h @@ -93,7 +93,7 @@ class ThreadPool { // we'll use assert, not internal_assert. assert(desired_num_threads > 0); - std::lock_guard lock(mutex); + std::scoped_lock lock(mutex); // Create all the threads. for (size_t i = 0; i < desired_num_threads; ++i) { @@ -104,7 +104,7 @@ class ThreadPool { ~ThreadPool() { // Wake everyone up and tell them the party's over and it's time to go home { - std::lock_guard lock(mutex); + std::scoped_lock lock(mutex); shutting_down = true; wakeup_threads.notify_all(); } @@ -117,7 +117,7 @@ class ThreadPool { template std::future async(Func func, Args... args) { - std::lock_guard lock(mutex); + std::scoped_lock lock(mutex); Job job; // Don't use std::forward here: we never want args passed by reference, diff --git a/tutorial/CMakeLists.txt b/tutorial/CMakeLists.txt index 88759a07c2d4..779c43d50882 100644 --- a/tutorial/CMakeLists.txt +++ b/tutorial/CMakeLists.txt @@ -1,10 +1,8 @@ -find_package(OpenMP) - configure_file(images/gray.png images/gray.png COPYONLY) configure_file(images/rgb.png images/rgb.png COPYONLY) function(add_tutorial source_file) - set(options WITH_IMAGE_IO WITH_OPENMP) + set(options WITH_IMAGE_IO) set(oneValueArgs) set(multiValueArgs SRCS GROUPS) cmake_parse_arguments(args "${options}" "${oneValueArgs}" "${multiValueArgs}" ${ARGN}) @@ -34,20 +32,6 @@ function(add_tutorial source_file) if (args_WITH_IMAGE_IO) target_link_libraries(${name} PRIVATE Halide::ImageIO) endif () - - if (args_WITH_OPENMP) - if (TARGET OpenMP::OpenMP_CXX) - target_link_libraries(${name} PRIVATE OpenMP::OpenMP_CXX) - else () - # Compile anyway but suppress warnings about unrecognised pragmas - target_compile_options( - "${name}" - PRIVATE - $<$:/Wd4068> - $<$>:-Wno-unknown-pragmas> - ) - endif () - endif () endfunction() # Declare tutorials @@ -73,8 +57,8 @@ endif () add_tutorial(lesson_06_realizing_over_shifted_domains.cpp) add_tutorial(lesson_07_multi_stage_pipelines.cpp WITH_IMAGE_IO) -add_tutorial(lesson_08_scheduling_2.cpp WITH_IMAGE_IO WITH_OPENMP GROUPS multithreaded) -add_tutorial(lesson_09_update_definitions.cpp WITH_IMAGE_IO WITH_OPENMP GROUPS multithreaded) +add_tutorial(lesson_08_scheduling_2.cpp WITH_IMAGE_IO GROUPS multithreaded) +add_tutorial(lesson_09_update_definitions.cpp WITH_IMAGE_IO GROUPS multithreaded) if ("NVPTX" IN_LIST Halide_LLVM_COMPONENTS) if (Halide_TARGET MATCHES "wasm") diff --git a/tutorial/clock.h b/tutorial/clock.h deleted file mode 100644 index 3132f0745de5..000000000000 --- a/tutorial/clock.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef HALIDE_TUTORIAL_CLOCK_H -#define HALIDE_TUTORIAL_CLOCK_H - -// A current_time function for use in the tests. Returns time in -// milliseconds. - -#include "halide_benchmark.h" - -inline double current_time() { - static auto start_time = Halide::Tools::benchmark_now().time_since_epoch(); - - auto now = Halide::Tools::benchmark_now().time_since_epoch() - start_time; - return std::chrono::duration_cast(now).count() / 1e3; -} - -#endif // HALIDE_TUTORIAL_CLOCK_H diff --git a/tutorial/lesson_01_basics.cpp b/tutorial/lesson_01_basics.cpp index e6a407b1e6b1..f00912efa004 100644 --- a/tutorial/lesson_01_basics.cpp +++ b/tutorial/lesson_01_basics.cpp @@ -3,18 +3,12 @@ // This lesson demonstrates basic usage of Halide as a JIT compiler for imaging. // On linux, you can compile and run it like so: -// g++ lesson_01*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_01 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_01 - -// On os x: -// g++ lesson_01*.cpp -g -I -L -lHalide -o lesson_01 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_01 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_01_basics -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_01*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_01 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_01 + +// On macOS: +// g++ lesson_01*.cpp -g -I -L -lHalide -o lesson_01 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_01 // The only Halide header file you need is Halide.h. It includes all of Halide. #include "Halide.h" diff --git a/tutorial/lesson_02_input_image.cpp b/tutorial/lesson_02_input_image.cpp index 7895659e04d1..6e34787b5a10 100644 --- a/tutorial/lesson_02_input_image.cpp +++ b/tutorial/lesson_02_input_image.cpp @@ -4,18 +4,12 @@ // them. // On linux, you can compile and run it like so: -// g++ lesson_02*.cpp -g -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -lpthread -ldl -o lesson_02 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_02 - -// On os x: -// g++ lesson_02*.cpp -g -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -o lesson_02 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_02 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_02_input_image -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_02*.cpp -g -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -lpthread -ldl -o lesson_02 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_02 + +// On macOS: +// g++ lesson_02*.cpp -g -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -o lesson_02 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_02 // The only Halide header file you need is Halide.h. It includes all of Halide. #include "Halide.h" diff --git a/tutorial/lesson_03_debugging_1.cpp b/tutorial/lesson_03_debugging_1.cpp index 0a253415b781..2fc7115b68f3 100644 --- a/tutorial/lesson_03_debugging_1.cpp +++ b/tutorial/lesson_03_debugging_1.cpp @@ -3,18 +3,12 @@ // This lesson demonstrates how to inspect what the Halide compiler is producing. // On linux, you can compile and run it like so: -// g++ lesson_03*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_03 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_03 +// g++ lesson_03*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_03 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_03 -// On os x: -// g++ lesson_03*.cpp -g -I -L -lHalide -o lesson_03 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_03 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_03_debugging_1 -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_03*.cpp -g -I -L -lHalide -o lesson_03 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_03 #include "Halide.h" #include diff --git a/tutorial/lesson_04_debugging_2.cpp b/tutorial/lesson_04_debugging_2.cpp index b801e9afeab1..c2bafdac6d13 100644 --- a/tutorial/lesson_04_debugging_2.cpp +++ b/tutorial/lesson_04_debugging_2.cpp @@ -3,18 +3,12 @@ // This lesson demonstrates how to follow what Halide is doing at runtime. // On linux, you can compile and run it like so: -// g++ lesson_04*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_04 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_04 +// g++ lesson_04*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_04 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_04 -// On os x: -// g++ lesson_04*.cpp -g -I -L -lHalide -o lesson_04 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_04 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_04_debugging_2 -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_04*.cpp -g -I -L -lHalide -o lesson_04 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_04 #include "Halide.h" #include @@ -56,16 +50,15 @@ int main() { // that describes the algorithm. // Now we tell Halide to use a parallel for loop over the y - // coordinate. On Linux we run this using a thread pool and a task - // queue. On OS X we call into grand central dispatch, which does - // the same thing for us. + // coordinate. Halide's runtime maintains its own pool of worker + // threads and a task queue. parallel_gradient.parallel(y); // This time the printfs should come out of order, because each // scanline is potentially being processed in a different // thread. The number of threads should adapt to your system, but - // on linux you can control it manually using the environment - // variable HL_NUM_THREADS. + // you can control it manually using the environment variable + // HL_NUM_THREADS. printf("\nEvaluating parallel_gradient\n"); parallel_gradient.realize({8, 8}); } diff --git a/tutorial/lesson_05_scheduling_1.cpp b/tutorial/lesson_05_scheduling_1.cpp index 6e97d392f822..d697e336f915 100644 --- a/tutorial/lesson_05_scheduling_1.cpp +++ b/tutorial/lesson_05_scheduling_1.cpp @@ -5,18 +5,12 @@ // parallelization, unrolling, and tiling. // On linux, you can compile and run it like so: -// g++ lesson_05*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_05 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_05 - -// On os x: -// g++ lesson_05*.cpp -g -I -L -lHalide -o lesson_05 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_05 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_05_scheduling_1 -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_05*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_05 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_05 + +// On macOS: +// g++ lesson_05*.cpp -g -I -L -lHalide -o lesson_05 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_05 #include "Halide.h" #include diff --git a/tutorial/lesson_06_realizing_over_shifted_domains.cpp b/tutorial/lesson_06_realizing_over_shifted_domains.cpp index 234cb9c9861a..793e7bfab8e6 100644 --- a/tutorial/lesson_06_realizing_over_shifted_domains.cpp +++ b/tutorial/lesson_06_realizing_over_shifted_domains.cpp @@ -4,18 +4,12 @@ // does not start at (0, 0). // On linux, you can compile and run it like so: -// g++ lesson_06*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_06 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_06 +// g++ lesson_06*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_06 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_06 -// On os x: -// g++ lesson_06*.cpp -g -I -L -lHalide -o lesson_06 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_06 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_06_realizing_over_shifted_domains -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_06*.cpp -g -I -L -lHalide -o lesson_06 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_06 #include "Halide.h" #include @@ -97,7 +91,8 @@ int main() { // fact read out-of-bounds and probably crash. // What if we want to evaluate our Func over some region that - // isn't rectangular? Too bad. Halide only does rectangles :) + // isn't rectangular? Unfortunately, we can't. Halide only supports + // rectangular domains. printf("Success!\n"); return 0; diff --git a/tutorial/lesson_07_multi_stage_pipelines.cpp b/tutorial/lesson_07_multi_stage_pipelines.cpp index 4c40b85fde41..3d5844c0b986 100644 --- a/tutorial/lesson_07_multi_stage_pipelines.cpp +++ b/tutorial/lesson_07_multi_stage_pipelines.cpp @@ -1,18 +1,12 @@ // Halide tutorial lesson 7: Multi-stage pipelines // On linux, you can compile and run it like so: -// g++ lesson_07*.cpp -g -std=c++17 -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -lpthread -ldl -o lesson_07 -// LD_LIBRARY_PATH= ./lesson_07 +// g++ lesson_07*.cpp -g -std=c++17 -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -lpthread -ldl -o lesson_07 +// LD_LIBRARY_PATH= ./lesson_07 -// On os x: -// g++ lesson_07*.cpp -g -std=c++17 -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -o lesson_07 -// DYLD_LIBRARY_PATH= ./lesson_07 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_07_multi_stage_pipelines -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_07*.cpp -g -std=c++17 -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -o lesson_07 +// DYLD_LIBRARY_PATH= ./lesson_07 #include "Halide.h" #include @@ -93,7 +87,7 @@ int main() { save_image(result, "blurry_parrot_1.png"); // This is usually the fastest way to deal with boundaries: - // don't write code that reads out of bounds :) The more + // avoid reading out of bounds in the first place. The more // general solution is our next example. } diff --git a/tutorial/lesson_08_scheduling_2.cpp b/tutorial/lesson_08_scheduling_2.cpp index 79ea5764bec5..4ab870e4a6ec 100644 --- a/tutorial/lesson_08_scheduling_2.cpp +++ b/tutorial/lesson_08_scheduling_2.cpp @@ -1,18 +1,12 @@ // Halide tutorial lesson 8: Scheduling multi-stage pipelines // On linux, you can compile and run it like so: -// g++ lesson_08*.cpp -g -std=c++17 -I -L -lHalide -lpthread -ldl -o lesson_08 -// LD_LIBRARY_PATH= ./lesson_08 +// g++ lesson_08*.cpp -g -std=c++17 -I -L -lHalide -lpthread -ldl -o lesson_08 +// LD_LIBRARY_PATH= ./lesson_08 -// On os x: -// g++ lesson_08*.cpp -g -std=c++17 -I -L -lHalide -o lesson_08 -// DYLD_LIBRARY_PATH= ./lesson_08 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_08_scheduling_2 -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_08*.cpp -g -std=c++17 -I -L -lHalide -o lesson_08 +// DYLD_LIBRARY_PATH= ./lesson_08 #include "Halide.h" #include @@ -459,7 +453,7 @@ int main() { // iterations are lying around for us to reuse. This assumes that // previous values of x or y happened earlier in time and have // finished. This is not true if you parallelize or vectorize - // either loop. Darn. If you parallelize, Halide won't inject the + // either loop. If you parallelize, Halide won't inject the // optimizations that skip work already done if there's a parallel // loop in between the store_at level and the compute_at level, // and won't fold the storage down into a circular buffer either, @@ -665,8 +659,6 @@ int main() { consumer.print_loop_nest(); printf("\n"); - // Look on my code, ye mighty, and despair! - // Let's check the C result against the Halide result. Doing // this I found several bugs in my C implementation, which // should tell you something. @@ -686,7 +678,7 @@ int main() { // This stuff is hard. We ended up in a three-way trade-off // between memory bandwidth, redundant work, and // parallelism. Halide can't make the correct choice for you - // automatically (sorry). Instead it tries to make it easier for + // automatically. Instead it tries to make it easier for // you to explore various options, without messing up your // program. In fact, Halide promises that scheduling calls like // compute_root won't change the meaning of your algorithm -- you diff --git a/tutorial/lesson_09_update_definitions.cpp b/tutorial/lesson_09_update_definitions.cpp index 39a0b0cc7389..422671849c22 100644 --- a/tutorial/lesson_09_update_definitions.cpp +++ b/tutorial/lesson_09_update_definitions.cpp @@ -1,31 +1,30 @@ // Halide tutorial lesson 9: Multi-pass Funcs, update definitions, and reductions // On linux, you can compile and run it like so: -// g++ lesson_09*.cpp -g -std=c++17 -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -lpthread -ldl -fopenmp -o lesson_09 -// LD_LIBRARY_PATH= ./lesson_09 +// g++ lesson_09*.cpp -g -std=c++17 -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -lpthread -ldl -o lesson_09 +// LD_LIBRARY_PATH= ./lesson_09 -// On os x (will only work if you actually have g++, not Apple's pretend g++ which is actually clang): -// g++ lesson_09*.cpp -g -std=c++17 -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -fopenmp -o lesson_09 -// DYLD_LIBRARY_PATH= ./lesson_09 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_09_update_definitions -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_09*.cpp -g -std=c++17 -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -o lesson_09 +// DYLD_LIBRARY_PATH= ./lesson_09 #include "Halide.h" +#include #include #include +#include #include -// We're going to be using x86 SSE intrinsics later on in this lesson. -#ifdef __SSE2__ -#include +// We're going to be using ARM Neon intrinsics later on in this lesson. +#ifdef __ARM_NEON +#include #endif -// We'll also need a clock to do performance testing at the end. -#include "clock.h" +// We'll also need a way to time performance at the end. +#include "halide_benchmark.h" + +// And a thread pool to run the C equivalent's parallel for loop. +#include "halide_thread_pool.h" using namespace Halide; @@ -787,114 +786,118 @@ int main() { Buffer halide_result = spread.realize({input.width(), input.height()}); -// The C equivalent is almost too horrible to contemplate (and -// took me a long time to debug). This time I want to time -// both the Halide version and the C version, so I'll use sse -// intrinsics for the vectorization, and openmp to do the -// parallel for loop (you'll need to compile with -fopenmp or -// similar to get correct timing). -#ifdef __SSE2__ +// The C equivalent is fairly involved. This time I want to time +// both the Halide version and the C version, so I'll use ARM +// Neon intrinsics for the vectorization, and Halide::Tools::ThreadPool +// to do the parallel for loop. +#ifdef __ARM_NEON // Don't include the time required to allocate the output buffer. Buffer c_result(input.width(), input.height()); -#ifdef _OPENMP - double t1 = current_time(); -#endif + int y_tiles = (input.height() + 31) / 32; - // Run this one hundred times so we can average the timing results. - for (int iters = 0; iters < 100; iters++) { + // The body of the parallel for loop over y tiles, run once per + // (iteration, tile) pair by whichever pool thread picks up that job. + auto run_tile = [&](int yo) { + int y_base = std::min(yo * 32, input.height() - 32); -#pragma omp parallel for - for (int yo = 0; yo < (input.height() + 31) / 32; yo++) { - int y_base = std::min(yo * 32, input.height() - 32); + // Compute clamped in a circular buffer of size 8 + // (smallest power of two greater than 5). Each thread + // needs its own allocation, so it must occur here. - // Compute clamped in a circular buffer of size 8 - // (smallest power of two greater than 5). Each thread - // needs its own allocation, so it must occur here. + size_t clamped_width = input.width() + 4; + std::vector clamped_storage(clamped_width * 8); - size_t clamped_width = input.width() + 4; - std::vector clamped_storage(clamped_width * 8); + for (int yi = 0; yi < 32; yi++) { + int y = y_base + yi; - for (int yi = 0; yi < 32; yi++) { - int y = y_base + yi; + uint8_t *output_row = &c_result(0, y); - uint8_t *output_row = &c_result(0, y); + // Compute clamped for this scanline, skipping rows + // already computed within this slice. + int min_y_clamped = (yi == 0) ? (y - 2) : (y + 2); + int max_y_clamped = (y + 2); + for (int cy = min_y_clamped; cy <= max_y_clamped; cy++) { + // Figure out which row of the circular buffer + // we're filling in using bitmasking: + uint8_t *clamped_row = + &clamped_storage[(cy & 7) * clamped_width]; - // Compute clamped for this scanline, skipping rows - // already computed within this slice. - int min_y_clamped = (yi == 0) ? (y - 2) : (y + 2); - int max_y_clamped = (y + 2); - for (int cy = min_y_clamped; cy <= max_y_clamped; cy++) { - // Figure out which row of the circular buffer - // we're filling in using bitmasking: - uint8_t *clamped_row = - &clamped_storage[(cy & 7) * clamped_width]; + // Figure out which row of the input we're reading + // from by clamping the y coordinate: + int clamped_y = std::min(std::max(cy, 0), input.height() - 1); + uint8_t *input_row = &input(0, clamped_y); - // Figure out which row of the input we're reading - // from by clamping the y coordinate: - int clamped_y = std::min(std::max(cy, 0), input.height() - 1); - uint8_t *input_row = &input(0, clamped_y); + // Fill it in with the padding. + for (int x = -2; x < input.width() + 2; x++) { + int clamped_x = std::min(std::max(x, 0), input.width() - 1); + *clamped_row++ = input_row[clamped_x]; + } + } + + // Now iterate over vectors of x for the pure step of the output. + for (int x_vec = 0; x_vec < (input.width() + 15) / 16; x_vec++) { + int x_base = std::min(x_vec * 16, input.width() - 16); - // Fill it in with the padding. - for (int x = -2; x < input.width() + 2; x++) { - int clamped_x = std::min(std::max(x, 0), input.width() - 1); - *clamped_row++ = input_row[clamped_x]; + // Allocate storage for the minimum and maximum + // helpers. One vector is enough. + uint8x16_t minimum_storage, maximum_storage; + + // The pure step for the maximum is a vector of zeros + maximum_storage = vdupq_n_u8(0); + + // The update step for maximum + for (int max_y = y - 2; max_y <= y + 2; max_y++) { + uint8_t *clamped_row = + &clamped_storage[(max_y & 7) * clamped_width]; + for (int max_x = x_base - 2; max_x <= x_base + 2; max_x++) { + uint8x16_t v = vld1q_u8(clamped_row + max_x + 2); + maximum_storage = vmaxq_u8(maximum_storage, v); } } - // Now iterate over vectors of x for the pure step of the output. - for (int x_vec = 0; x_vec < (input.width() + 15) / 16; x_vec++) { - int x_base = std::min(x_vec * 16, input.width() - 16); - - // Allocate storage for the minimum and maximum - // helpers. One vector is enough. - __m128i minimum_storage, maximum_storage; - - // The pure step for the maximum is a vector of zeros - maximum_storage = _mm_setzero_si128(); - - // The update step for maximum - for (int max_y = y - 2; max_y <= y + 2; max_y++) { - uint8_t *clamped_row = - &clamped_storage[(max_y & 7) * clamped_width]; - for (int max_x = x_base - 2; max_x <= x_base + 2; max_x++) { - __m128i v = _mm_loadu_si128( - (__m128i const *)(clamped_row + max_x + 2)); - maximum_storage = _mm_max_epu8(maximum_storage, v); - } - } + // The pure step for the minimum is a vector of ones. + minimum_storage = vdupq_n_u8(0xff); - // The pure step for the minimum is a vector of - // ones. Create it by comparing something to - // itself. - minimum_storage = _mm_cmpeq_epi32(_mm_setzero_si128(), - _mm_setzero_si128()); - - // The update step for minimum. - for (int min_y = y - 2; min_y <= y + 2; min_y++) { - uint8_t *clamped_row = - &clamped_storage[(min_y & 7) * clamped_width]; - for (int min_x = x_base - 2; min_x <= x_base + 2; min_x++) { - __m128i v = _mm_loadu_si128( - (__m128i const *)(clamped_row + min_x + 2)); - minimum_storage = _mm_min_epu8(minimum_storage, v); - } + // The update step for minimum. + for (int min_y = y - 2; min_y <= y + 2; min_y++) { + uint8_t *clamped_row = + &clamped_storage[(min_y & 7) * clamped_width]; + for (int min_x = x_base - 2; min_x <= x_base + 2; min_x++) { + uint8x16_t v = vld1q_u8(clamped_row + min_x + 2); + minimum_storage = vminq_u8(minimum_storage, v); } + } - // Now compute the spread. - __m128i spread = _mm_sub_epi8(maximum_storage, minimum_storage); + // Now compute the spread. + uint8x16_t spread = vsubq_u8(maximum_storage, minimum_storage); - // Store it. - _mm_storeu_si128((__m128i *)(output_row + x_base), spread); - } + // Store it. + vst1q_u8(output_row + x_base, spread); } } + }; + + // The pool's worker threads are created here, once, so that + // spawning them isn't included in the timing below. + ThreadPool pool; + std::vector> futures; + futures.reserve(y_tiles); + + auto t1 = benchmark_now(); + + // Run this one hundred times so we can average the timing results. + for (int iters = 0; iters < 100; iters++) { + for (int yo = 0; yo < y_tiles; yo++) { + futures.push_back(pool.async(run_tile, yo)); + } + for (auto &f : futures) { + f.get(); + } + futures.clear(); } -// Skip the timing comparison if we don't have openmp -// enabled. Otherwise it's unfair to C. -#ifdef _OPENMP - double t2 = current_time(); + auto t2 = benchmark_now(); // Now run the Halide version again without the // jit-compilation overhead. Also run it one hundred times. @@ -902,7 +905,7 @@ int main() { spread.realize(halide_result); } - double t3 = current_time(); + auto t3 = benchmark_now(); // Report the timings. On my machine they both take about 3ms // for the 4-megapixel input (fast!), which makes sense, @@ -910,9 +913,8 @@ int main() { // parallelization strategy. However I find the Halide easier // to read, write, debug, modify, and port. printf("Halide spread took %f ms. C equivalent took %f ms\n", - (t3 - t2) / 100, (t2 - t1) / 100); - -#endif // _OPENMP + 1000 * benchmark_duration_seconds(t2, t3) / 100, + 1000 * benchmark_duration_seconds(t1, t2) / 100); // Check the results match: for (int y = 0; y < input.height(); y++) { @@ -925,7 +927,7 @@ int main() { } } -#endif // __SSE2__ +#endif // __ARM_NEON } printf("Success!\n"); diff --git a/tutorial/lesson_10_aot_compilation_generate.cpp b/tutorial/lesson_10_aot_compilation_generate.cpp index 0165d95be031..b060dba770b1 100644 --- a/tutorial/lesson_10_aot_compilation_generate.cpp +++ b/tutorial/lesson_10_aot_compilation_generate.cpp @@ -1,4 +1,4 @@ -// Halide tutorial lesson 10: AOT compilation part 1 +// Halide tutorial lesson 10: AOT compilation: compiling the pipeline // This lesson demonstrates how to use Halide as an more traditional // ahead-of-time (AOT) compiler. @@ -10,27 +10,21 @@ // compiling this code is a multi-step process. // On linux, you can compile and run it like so: -// g++ lesson_10*generate.cpp -g -std=c++17 -I -L -lHalide -lpthread -ldl -o lesson_10_generate -// LD_LIBRARY_PATH= ./lesson_10_generate -// g++ lesson_10*run.cpp lesson_10_halide.a -std=c++17 -I -lpthread -ldl -o lesson_10_run +// g++ lesson_10*generate.cpp -g -std=c++17 -I -L -lHalide -lpthread -ldl -o lesson_10_generate +// LD_LIBRARY_PATH= ./lesson_10_generate +// g++ lesson_10*run.cpp lesson_10_halide.a -std=c++17 -I -lpthread -ldl -o lesson_10_run // ./lesson_10_run -// On os x: -// g++ lesson_10*generate.cpp -g -std=c++17 -I -L -lHalide -o lesson_10_generate -// DYLD_LIBRARY_PATH= ./lesson_10_generate -// g++ lesson_10*run.cpp lesson_10_halide.a -o lesson_10_run -I +// On macOS: +// g++ lesson_10*generate.cpp -g -std=c++17 -I -L -lHalide -o lesson_10_generate +// DYLD_LIBRARY_PATH= ./lesson_10_generate +// g++ lesson_10*run.cpp lesson_10_halide.a -std=c++17 -o lesson_10_run -I // ./lesson_10_run // The benefits of this approach are that the final program: // - Doesn't do any jit compilation at runtime, so it's fast. // - Doesn't depend on libHalide at all, so it's a small, easy-to-deploy binary. -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_10_aot_compilation_run -// in a shell with the current directory at the top of the halide -// source tree. - #include "Halide.h" #include using namespace Halide; diff --git a/tutorial/lesson_10_aot_compilation_run.cpp b/tutorial/lesson_10_aot_compilation_run.cpp index 03c840f79531..c1d4f3832ec0 100644 --- a/tutorial/lesson_10_aot_compilation_run.cpp +++ b/tutorial/lesson_10_aot_compilation_run.cpp @@ -1,4 +1,4 @@ -// Halide tutorial lesson 10: AOT compilation part 2 +// Halide tutorial lesson 10: AOT compilation: running the compiled pipeline // Before reading this file, see lesson_10_aot_compilation_generate.cpp diff --git a/tutorial/lesson_11_cross_compilation.cpp b/tutorial/lesson_11_cross_compilation.cpp index 5215206c5223..4c296560c8b3 100644 --- a/tutorial/lesson_11_cross_compilation.cpp +++ b/tutorial/lesson_11_cross_compilation.cpp @@ -4,18 +4,12 @@ // generate code for any platform from any platform. // On linux, you can compile and run it like so: -// g++ lesson_11*.cpp -g -std=c++17 -I -L -lHalide -lpthread -ldl -o lesson_11 -// LD_LIBRARY_PATH= ./lesson_11 - -// On os x: -// g++ lesson_11*.cpp -g -std=c++17 -I -L -lHalide -o lesson_11 -// DYLD_LIBRARY_PATH= ./lesson_11 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_11_cross_compilation -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_11*.cpp -g -std=c++17 -I -L -lHalide -lpthread -ldl -o lesson_11 +// LD_LIBRARY_PATH= ./lesson_11 + +// On macOS: +// g++ lesson_11*.cpp -g -std=c++17 -I -L -lHalide -o lesson_11 +// DYLD_LIBRARY_PATH= ./lesson_11 #include "Halide.h" #include diff --git a/tutorial/lesson_12_using_the_gpu.cpp b/tutorial/lesson_12_using_the_gpu.cpp index 0c399a537f3d..72827c73fa1f 100644 --- a/tutorial/lesson_12_using_the_gpu.cpp +++ b/tutorial/lesson_12_using_the_gpu.cpp @@ -3,25 +3,19 @@ // This lesson demonstrates how to use Halide to run code on a GPU using OpenCL. // On linux, you can compile and run it like so: -// g++ lesson_12*.cpp -g -std=c++17 -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -lpthread -ldl -o lesson_12 -// LD_LIBRARY_PATH= ./lesson_12 +// g++ lesson_12*.cpp -g -std=c++17 -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -lpthread -ldl -o lesson_12 +// LD_LIBRARY_PATH= ./lesson_12 -// On os x: -// g++ lesson_12*.cpp -g -std=c++17 -I -I -L -lHalide `libpng-config --cflags --ldflags` -ljpeg -o lesson_12 -// DYLD_LIBRARY_PATH= ./lesson_12 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_12_using_the_gpu -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_12*.cpp -g -std=c++17 -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -o lesson_12 +// DYLD_LIBRARY_PATH= ./lesson_12 #include "Halide.h" #include -// Include a clock to do performance testing. -#include "clock.h" +// Include halide_benchmark.h to do performance testing. +#include "halide_benchmark.h" // Include some support code for loading pngs. #include "halide_image_io.h" @@ -205,7 +199,7 @@ class MyPipeline { double best_time = 0.0; for (int i = 0; i < 3; i++) { - double t1 = current_time(); + auto t1 = benchmark_now(); // Run the filter 100 times. for (int j = 0; j < 100; j++) { @@ -215,9 +209,9 @@ class MyPipeline { // Force any GPU code to finish by copying the buffer back to the CPU. output.copy_to_host(); - double t2 = current_time(); + auto t2 = benchmark_now(); - double elapsed = (t2 - t1) / 100; + double elapsed = 1000 * benchmark_duration_seconds(t1, t2) / 100; if (i == 0 || elapsed < best_time) { best_time = elapsed; } @@ -296,8 +290,6 @@ Target find_gpu_target() { } features_to_try.push_back(Target::OpenCL); } else if (target.os == Target::OSX) { - // OS X doesn't update its OpenCL drivers, so they tend to be broken. - // CUDA would also be a fine choice on machines with NVidia GPUs. features_to_try.push_back(Target::Metal); } else { features_to_try.push_back(Target::OpenCL); diff --git a/tutorial/lesson_13_tuples.cpp b/tutorial/lesson_13_tuples.cpp index 9a4c616637cd..ba13bb9b64dc 100644 --- a/tutorial/lesson_13_tuples.cpp +++ b/tutorial/lesson_13_tuples.cpp @@ -4,18 +4,12 @@ // values. // On linux, you can compile and run it like so: -// g++ lesson_13*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_13 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_13 - -// On os x: -// g++ lesson_13*.cpp -g -I -L -lHalide -o lesson_13 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_13 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_13_tuples -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_13*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_13 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_13 + +// On macOS: +// g++ lesson_13*.cpp -g -I -L -lHalide -o lesson_13 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_13 #include "Halide.h" #include diff --git a/tutorial/lesson_14_types.cpp b/tutorial/lesson_14_types.cpp index aaca8cb18d5d..ad53e3373169 100644 --- a/tutorial/lesson_14_types.cpp +++ b/tutorial/lesson_14_types.cpp @@ -3,18 +3,12 @@ // This lesson more precisely describes Halide's type system. // On linux, you can compile and run it like so: -// g++ lesson_14*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_14 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_14 - -// On os x: -// g++ lesson_14*.cpp -g -I -L -lHalide -o lesson_14 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_14 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_14_types -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_14*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_14 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_14 + +// On macOS: +// g++ lesson_14*.cpp -g -I -L -lHalide -o lesson_14 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_14 #include "Halide.h" #include diff --git a/tutorial/lesson_15_generators.cpp b/tutorial/lesson_15_generators.cpp index 00ecf3fe57bf..14cae0b11cdb 100644 --- a/tutorial/lesson_15_generators.cpp +++ b/tutorial/lesson_15_generators.cpp @@ -1,22 +1,16 @@ -// Halide tutorial lesson 15: Generators part 1 +// Halide tutorial lesson 15: Generators: writing a generator // This lesson demonstrates how to encapsulate Halide pipelines into // reusable components called generators. // On linux, you can compile and run it like so: -// g++ lesson_15*.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -lpthread -ldl -o lesson_15_generate +// g++ lesson_15*.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -lpthread -ldl -o lesson_15_generate // bash lesson_15_generators_usage.sh -// On os x: -// g++ lesson_15*.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -o lesson_15_generate +// On macOS: +// g++ lesson_15*.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -o lesson_15_generate // bash lesson_15_generators_usage.sh -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_15_generators -// in a shell with the current directory at the top of the halide -// source tree. - #include "Halide.h" #include diff --git a/tutorial/lesson_15_generators_usage.sh b/tutorial/lesson_15_generators_usage.sh index b597ef23e8e8..581930594f28 100755 --- a/tutorial/lesson_15_generators_usage.sh +++ b/tutorial/lesson_15_generators_usage.sh @@ -1,6 +1,6 @@ #!/usr/bin/env bash -# Halide tutorial lesson 15: Generators part 2 +# Halide tutorial lesson 15: Generators: command-line usage # This shell script demonstrates how to use a binary containing # Generators from the command line. Normally you'd call these binaries @@ -14,6 +14,19 @@ # To run this script: # bash lesson_15_generators_usage.sh +# Later on, we'll cross-compile object files for other targets (e.g. Linux and +# Windows, from macOS). The system nm/ar on macOS only understand Mach-O, so on +# macOS we prefer llvm-nm/llvm-ar (from Homebrew LLVM, if installed) since they +# can read object files for any platform. +NM="nm" +AR="ar" +if [ "$(uname)" = "Darwin" ] && command -v brew >/dev/null && LLVM_PREFIX=$(brew --prefix llvm 2>/dev/null); then + if [ -x "$LLVM_PREFIX/bin/llvm-nm" ] && [ -x "$LLVM_PREFIX/bin/llvm-ar" ]; then + NM="$LLVM_PREFIX/bin/llvm-nm" + AR="$LLVM_PREFIX/bin/llvm-ar" + fi +fi + # First we define a helper function that checks that a file exists check_file_exists() { FILE=$1 @@ -27,7 +40,7 @@ check_file_exists() { check_symbol() { FILE=$1 SYM=$2 - if ! nm "$FILE" | grep "$SYM" >/dev/null; then + if ! "$NM" "$FILE" | grep "$SYM" >/dev/null; then echo "$SYM not found in $FILE" exit 1 fi @@ -152,18 +165,18 @@ check_symbol my_second_generator_3.a my_second_generator_3 # files. echo "The halide runtime:" -nm my_second_generator_1.a | grep "[SWT] _\?halide_" +"$NM" my_second_generator_1.a | grep "[SWT] _\?halide_" # Let's define some functions to check that the runtime exists in a file. check_runtime() { - if ! (nm "$1" | grep "[TSW] _\?halide_" >/dev/null); then + if ! ("$NM" "$1" | grep "[TSW] _\?halide_" >/dev/null); then echo "Halide runtime not found in $1" exit 1 fi } check_no_runtime() { - if nm "$1" | grep "[TSW] _\?halide_" >/dev/null; then + if "$NM" "$1" | grep "[TSW] _\?halide_" >/dev/null; then echo "Halide runtime found in $1" exit 1 fi @@ -224,7 +237,7 @@ check_runtime halide_runtime_x86.o # gives us three versions of the pipeline for varying levels of x86, # combined with a single runtime that will work on nearly all x86 # processors. -ar q my_first_generator_multi.a \ +"$AR" q my_first_generator_multi.a \ my_first_generator_basic.o \ my_first_generator_sse41.o \ my_first_generator_avx.o \ diff --git a/tutorial/lesson_16_rgb_generate.cpp b/tutorial/lesson_16_rgb_generate.cpp index fb8983f1c4ca..c07b75e96bec 100644 --- a/tutorial/lesson_16_rgb_generate.cpp +++ b/tutorial/lesson_16_rgb_generate.cpp @@ -1,27 +1,21 @@ -// Halide tutorial lesson 16: RGB images and memory layouts part 1 +// Halide tutorial lesson 16: RGB images and memory layouts: compiling the pipeline // This lesson demonstrates how to feed Halide RGB images in // interleaved or planar format, and how to write code optimized for // each case. -// On linux or os x, you can compile and run it like so: +// On linux or macOS, you can compile and run it like so: -// g++ lesson_16_rgb_generate.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -lpthread -ldl -o lesson_16_generate -// export LD_LIBRARY_PATH= # For linux -// export DYLD_LIBRARY_PATH= # For OS X +// g++ lesson_16_rgb_generate.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -lpthread -ldl -o lesson_16_generate +// export LD_LIBRARY_PATH= # For linux +// export DYLD_LIBRARY_PATH= # For macOS // ./lesson_16_generate -g brighten -o . -f brighten_planar target=host layout=planar // ./lesson_16_generate -g brighten -o . -f brighten_interleaved target=host layout=interleaved // ./lesson_16_generate -g brighten -o . -f brighten_either target=host layout=either // ./lesson_16_generate -g brighten -o . -f brighten_specialized target=host layout=specialized -// g++ lesson_16_rgb_run.cpp brighten_*.o -ldl -lpthread -o lesson_16_run +// g++ lesson_16_rgb_run.cpp brighten_*.a -std=c++17 -I -I -ldl -lpthread -o lesson_16_run // ./lesson_16_run -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_16_rgb_run -// in a shell with the current directory at the top of the halide -// source tree. - #include "Halide.h" #include diff --git a/tutorial/lesson_16_rgb_run.cpp b/tutorial/lesson_16_rgb_run.cpp index 0caec3e31803..3906a111f8ba 100644 --- a/tutorial/lesson_16_rgb_run.cpp +++ b/tutorial/lesson_16_rgb_run.cpp @@ -1,4 +1,4 @@ -// Halide tutorial lesson 16: RGB images and memory layouts part 2 +// Halide tutorial lesson 16: RGB images and memory layouts: running the compiled pipeline // Before reading this file, see lesson_16_rgb_generate.cpp diff --git a/tutorial/lesson_17_predicated_rdom.cpp b/tutorial/lesson_17_predicated_rdom.cpp index 0b9fd0a06f9f..cd6ed9c9648b 100644 --- a/tutorial/lesson_17_predicated_rdom.cpp +++ b/tutorial/lesson_17_predicated_rdom.cpp @@ -4,18 +4,12 @@ // subsets of a reduction domain using predicates. // On linux, you can compile and run it like so: -// g++ lesson_17*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_17 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_17 - -// On os x: -// g++ lesson_17*.cpp -g -I -L -lHalide -o lesson_17 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_17 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_17_predicated_rdom -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_17*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_17 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_17 + +// On macOS: +// g++ lesson_17*.cpp -g -I -L -lHalide -o lesson_17 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_17 #include "Halide.h" #include diff --git a/tutorial/lesson_18_parallel_associative_reductions.cpp b/tutorial/lesson_18_parallel_associative_reductions.cpp index 7123f076d32d..9730a6a84d2d 100644 --- a/tutorial/lesson_18_parallel_associative_reductions.cpp +++ b/tutorial/lesson_18_parallel_associative_reductions.cpp @@ -4,18 +4,12 @@ // reduction using the scheduling directive 'rfactor'. // On linux, you can compile and run it like so: -// g++ lesson_18*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_18 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_18 - -// On os x: -// g++ lesson_18*.cpp -g -I -L -lHalide -o lesson_18 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_18 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_18_parallel_associative_reductions -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_18*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_18 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_18 + +// On macOS: +// g++ lesson_18*.cpp -g -I -L -lHalide -o lesson_18 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_18 #include "Halide.h" #include diff --git a/tutorial/lesson_19_wrapper_funcs.cpp b/tutorial/lesson_19_wrapper_funcs.cpp index 869b5ef012b3..d56e6159c910 100644 --- a/tutorial/lesson_19_wrapper_funcs.cpp +++ b/tutorial/lesson_19_wrapper_funcs.cpp @@ -5,17 +5,12 @@ // from a Func or an ImageParam. // On linux, you can compile and run it like so: -// g++ lesson_19*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_19 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_19 +// g++ lesson_19*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_19 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_19 -// On os x: -// g++ lesson_19*.cpp -g -I -L -lHalide -o lesson_19 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_19 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_19_wrapper_funcs -// in a shell at the top of the halide source tree. +// On macOS: +// g++ lesson_19*.cpp -g -I -L -lHalide -o lesson_19 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_19 // The only Halide header file you need is Halide.h. It includes all of Halide. #include "Halide.h" diff --git a/tutorial/lesson_20_cloning_funcs.cpp b/tutorial/lesson_20_cloning_funcs.cpp index b8290d60fabe..c442f5cb72e7 100644 --- a/tutorial/lesson_20_cloning_funcs.cpp +++ b/tutorial/lesson_20_cloning_funcs.cpp @@ -4,17 +4,12 @@ // a Func. // On linux, you can compile and run it like so: -// g++ lesson_20*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_20 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_20 +// g++ lesson_20*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_20 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_20 -// On os x: -// g++ lesson_20*.cpp -g -I -L -lHalide -o lesson_20 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_20 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_20_cloning_funcs -// in a shell at the top of the halide source tree. +// On macOS: +// g++ lesson_20*.cpp -g -I -L -lHalide -o lesson_20 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_20 // The only Halide header file you need is Halide.h. It includes all of Halide. #include "Halide.h" diff --git a/tutorial/lesson_21_auto_scheduler_generate.cpp b/tutorial/lesson_21_auto_scheduler_generate.cpp index 87dfb4511243..b78a26edd044 100644 --- a/tutorial/lesson_21_auto_scheduler_generate.cpp +++ b/tutorial/lesson_21_auto_scheduler_generate.cpp @@ -1,26 +1,20 @@ -// Halide tutorial lesson 21: Auto-Scheduler +// Halide tutorial lesson 21: Auto-Scheduler: generating a schedule // So far we have written Halide schedules by hand, but it is also possible to // ask Halide to suggest a reasonable schedule. We call this auto-scheduling. // This lesson demonstrates how to use the autoscheduler to generate a // copy-pasteable CPU schedule that can be subsequently improved upon. -// On linux or os x, you can compile and run it like so: +// On linux or macOS, you can compile and run it like so: -// g++ lesson_21_auto_scheduler_generate.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -lpthread -ldl -o lesson_21_generate -// export LD_LIBRARY_PATH= # For linux -// export DYLD_LIBRARY_PATH= # For OS X -// ./lesson_21_generate -o . -g auto_schedule_gen -f auto_schedule_false -e static_library,h,schedule target=host auto_schedule=false -// ./lesson_21_generate -o . -g auto_schedule_gen -f auto_schedule_true -e static_library,h,schedule -p -S Mullapudi2016 target=host autoscheduler=Mullapudi2016 autoscheduler.parallelism=32 autoscheduler.last_level_cache_size=16777216 autoscheduler.balance=40 -// g++ lesson_21_auto_scheduler_run.cpp -std=c++17 -I -I auto_schedule_false.a auto_schedule_true.a -ldl -lpthread -o lesson_21_run +// g++ lesson_21_auto_scheduler_generate.cpp /GenGen.cpp -g -std=c++17 -fno-rtti -I -L -lHalide -lpthread -ldl -o lesson_21_generate +// export LD_LIBRARY_PATH= # For linux +// export DYLD_LIBRARY_PATH= # For macOS +// ./lesson_21_generate -o . -g auto_schedule_gen -f auto_schedule_false -e static_library,h,schedule target=host +// ./lesson_21_generate -o . -g auto_schedule_gen -f auto_schedule_true -e static_library,h,schedule -p target=host autoscheduler=Mullapudi2016 autoscheduler.parallelism=32 autoscheduler.last_level_cache_size=16777216 autoscheduler.balance=40 +// g++ lesson_21_auto_scheduler_run.cpp -std=c++17 -I -I auto_schedule_false.a auto_schedule_true.a -ldl -lpthread -o lesson_21_run // ./lesson_21_run -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_21_auto_scheduler_run -// in a shell with the current directory at the top of the halide -// source tree. - #include "Halide.h" #include diff --git a/tutorial/lesson_21_auto_scheduler_run.cpp b/tutorial/lesson_21_auto_scheduler_run.cpp index 242ab7df9739..79de6fcd1493 100644 --- a/tutorial/lesson_21_auto_scheduler_run.cpp +++ b/tutorial/lesson_21_auto_scheduler_run.cpp @@ -1,4 +1,4 @@ -// Halide tutorial lesson 21: Auto-Scheduler +// Halide tutorial lesson 21: Auto-Scheduler: running the compiled pipeline // Before reading this file, see lesson_21_auto_scheduler_generate.cpp diff --git a/tutorial/lesson_22_jit_performance.cpp b/tutorial/lesson_22_jit_performance.cpp index ed20ac2160cb..39c66fde2577 100644 --- a/tutorial/lesson_22_jit_performance.cpp +++ b/tutorial/lesson_22_jit_performance.cpp @@ -4,17 +4,12 @@ // various Halide methods of doing "Just-In-Time" compilation. // On linux, you can compile and run it like so: -// g++ lesson_22*.cpp -g -I -I -L -lHalide -lpthread -ldl -o lesson_22 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_22 +// g++ lesson_22*.cpp -g -I -I -L -lHalide -lpthread -ldl -o lesson_22 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_22 -// On os x: -// g++ lesson_22*.cpp -g -I -I -L -lHalide -o lesson_22 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_22 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_22_jit_performance -// in a shell at the top of the halide source tree. +// On macOS: +// g++ lesson_22*.cpp -g -I -I -L -lHalide -o lesson_22 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_22 #include "Halide.h" #include "halide_benchmark.h" diff --git a/tutorial/lesson_23_serialization.cpp b/tutorial/lesson_23_serialization.cpp index 4b7057e00d14..dc0a0edb2ab6 100644 --- a/tutorial/lesson_23_serialization.cpp +++ b/tutorial/lesson_23_serialization.cpp @@ -12,18 +12,12 @@ // change; we recommend that you avoid relying on it for production work at this time. // On linux, you can compile this tutorial and run it like so: -// g++ lesson_23*.cpp -g -I -I -L -lHalide -lpthread -ldl -o lesson_23 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_23 - -// On os x: -// g++ lesson_23*.cpp -g -I -I -L -lHalide -o lesson_23 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_23 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_23_serialization -// in a shell with the current directory at the top of the halide -// source tree. +// g++ lesson_23*.cpp -g -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -lpthread -ldl -o lesson_23 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_23 + +// On macOS: +// g++ lesson_23*.cpp -g -I -I -L -lHalide $(pkg-config --cflags --libs libpng libjpeg) -o lesson_23 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_23 #include "Halide.h" #include diff --git a/tutorial/lesson_24_async.cpp b/tutorial/lesson_24_async.cpp index 648be7be77a2..c48fac0d6479 100644 --- a/tutorial/lesson_24_async.cpp +++ b/tutorial/lesson_24_async.cpp @@ -4,18 +4,12 @@ // using scheduling directives 'async' and 'ring_buffer'. // On linux, you can compile and run it like so: -// g++ lesson_24*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_24 -std=c++17 -// LD_LIBRARY_PATH= ./lesson_24 +// g++ lesson_24*.cpp -g -I -L -lHalide -lpthread -ldl -o lesson_24 -std=c++17 +// LD_LIBRARY_PATH= ./lesson_24 -// On os x: -// g++ lesson_24*.cpp -g -I -L -lHalide -o lesson_24 -std=c++17 -// DYLD_LIBRARY_PATH= ./lesson_24 - -// If you have the entire Halide source tree, you can also build it by -// running: -// make tutorial_lesson_24_async -// in a shell with the current directory at the top of the halide -// source tree. +// On macOS: +// g++ lesson_24*.cpp -g -I -L -lHalide -o lesson_24 -std=c++17 +// DYLD_LIBRARY_PATH= ./lesson_24 #include "Halide.h" #include