Merge branch 'main' of https://github.com/ashvardanian/ParallelReductionsBenchmark

Author: ashvardanian

.vscode/settings.json

@@ -1,105 +1,115 @@
 {
+    "cSpell.words": [
+        "ashvardanian",
+        "blas",
+        "CCCL",
+        "constexpr",
+        "cublas",
+        "CUDA",
+        "Kahan",
+        "METALLIB",
+        "openmp",
+        "shfl",
+        "SPIR",
+        "STREQUAL",
+        "threadgroup",
+        "Vardanian",
+        "wmma"
+    ],
     "files.associations": {
+        "*.cl": "cpp",
+        "*.ipp": "cpp",
+        "*.metal": "cpp",
+        "*.tcc": "cpp",
+        "__debug": "cpp",
+        "__nullptr": "cpp",
+        "algorithm": "cpp",
+        "any": "cpp",
         "array": "cpp",
         "atomic": "cpp",
+        "barrier": "cpp",
         "bit": "cpp",
-        "*.tcc": "cpp",
+        "bitset": "cpp",
         "cctype": "cpp",
+        "cfenv": "cpp",
+        "charconv": "cpp",
+        "chrono": "cpp",
+        "cinttypes": "cpp",
         "clocale": "cpp",
         "cmath": "cpp",
+        "codecvt": "cpp",
         "compare": "cpp",
+        "complex": "cpp",
         "concepts": "cpp",
+        "condition_variable": "cpp",
+        "csetjmp": "cpp",
+        "csignal": "cpp",
         "cstdarg": "cpp",
         "cstddef": "cpp",
         "cstdint": "cpp",
         "cstdio": "cpp",
         "cstdlib": "cpp",
+        "cstring": "cpp",
+        "ctime": "cpp",
+        "cuchar": "cpp",
         "cwchar": "cpp",
         "cwctype": "cpp",
-        "map": "cpp",
-        "unordered_map": "cpp",
-        "vector": "cpp",
+        "deque": "cpp",
         "exception": "cpp",
-        "algorithm": "cpp",
-        "functional": "cpp",
-        "iterator": "cpp",
-        "memory": "cpp",
-        "memory_resource": "cpp",
-        "numeric": "cpp",
-        "optional": "cpp",
-        "random": "cpp",
-        "string": "cpp",
-        "string_view": "cpp",
-        "system_error": "cpp",
-        "tuple": "cpp",
-        "type_traits": "cpp",
-        "utility": "cpp",
+        "expected": "cpp",
+        "forward_list": "cpp",
         "fstream": "cpp",
+        "functional": "cpp",
+        "future": "cpp",
+        "hash_map": "cpp",
+        "hash_set": "cpp",
         "initializer_list": "cpp",
+        "iomanip": "cpp",
         "iosfwd": "cpp",
         "iostream": "cpp",
         "istream": "cpp",
+        "iterator": "cpp",
+        "latch": "cpp",
         "limits": "cpp",
+        "list": "cpp",
+        "locale": "cpp",
+        "map": "cpp",
+        "memory": "cpp",
+        "memory_resource": "cpp",
+        "mutex": "cpp",
         "new": "cpp",
         "numbers": "cpp",
+        "numeric": "cpp",
+        "optional": "cpp",
         "ostream": "cpp",
+        "propagate_const": "cpp",
+        "random": "cpp",
         "ranges": "cpp",
-        "sstream": "cpp",
-        "stdexcept": "cpp",
-        "streambuf": "cpp",
-        "cinttypes": "cpp",
-        "typeinfo": "cpp",
-        "deque": "cpp",
+        "ratio": "cpp",
         "regex": "cpp",
-        "forward_list": "cpp",
-        "list": "cpp",
-        "valarray": "cpp",
-        "cstring": "cpp",
-        "ctime": "cpp",
-        "any": "cpp",
-        "bitset": "cpp",
-        "chrono": "cpp",
-        "codecvt": "cpp",
-        "complex": "cpp",
-        "condition_variable": "cpp",
+        "scoped_allocator": "cpp",
+        "semaphore": "cpp",
         "set": "cpp",
-        "unordered_set": "cpp",
-        "ratio": "cpp",
-        "future": "cpp",
-        "iomanip": "cpp",
-        "mutex": "cpp",
         "shared_mutex": "cpp",
+        "span": "cpp",
+        "sstream": "cpp",
+        "stdexcept": "cpp",
         "stop_token": "cpp",
+        "streambuf": "cpp",
+        "string": "cpp",
+        "string_view": "cpp",
+        "strstream": "cpp",
+        "system_error": "cpp",
         "thread": "cpp",
+        "tuple": "cpp",
+        "type_traits": "cpp",
         "typeindex": "cpp",
+        "typeinfo": "cpp",
+        "unordered_map": "cpp",
+        "unordered_set": "cpp",
+        "utility": "cpp",
+        "valarray": "cpp",
         "variant": "cpp",
-        "csetjmp": "cpp",
-        "csignal": "cpp",
-        "strstream": "cpp",
-        "scoped_allocator": "cpp",
-        "cfenv": "cpp",
-        "hash_map": "cpp",
-        "*.ipp": "cpp",
-        "__debug": "cpp",
-        "barrier": "cpp",
-        "charconv": "cpp",
-        "propagate_const": "cpp",
-        "semaphore": "cpp",
-        "span": "cpp",
-        "expected": "cpp",
-        "locale": "cpp",
-        "__nullptr": "cpp",
-        "cuchar": "cpp",
-        "hash_set": "cpp",
-        "latch": "cpp"
-    },
-    "cSpell.words": [
-        "ashvardanian",
-        "CCCL",
-        "CUDA",
-        "Kahan",
-        "shfl",
-        "SPIR",
-        "STREQUAL"
-    ]
+        "vector": "cpp"
+    }
 }

CMakeLists.txt

@@ -104,6 +104,50 @@ if(CMAKE_CUDA_COMPILER_ID STREQUAL "NVIDIA" OR CMAKE_CUDA_COMPILER_ID STREQUAL "
     set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --extended-lambda")
     set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode arch=compute_86,code=sm_86")
 
+elseif(CMAKE_CXX_COMPILER_ID STREQUAL "AppleClang" AND 0)
+    set(METAL_SRC ${CMAKE_CURRENT_SOURCE_DIR}/reduce_metal.msl)
+    set(METALLIB_OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/reduce_metal.metallib)
+
+    add_custom_command(
+        OUTPUT ${METALLIB_OUTPUT}
+        COMMAND xcrun -sdk macosx metal -c ${METAL_SRC} -o reduce_metal.air
+        COMMAND xcrun -sdk macosx metallib reduce_metal.air -o ${METALLIB_OUTPUT}
+        DEPENDS ${METAL_SRC}
+        COMMENT "Compiling Metal shader: ${METAL_SRC} -> ${METALLIB_OUTPUT}"
+        VERBATIM
+    )
+
+    # Create a pseudo target to build the metallib
+    add_custom_target(MetalLibBuild ALL
+        DEPENDS ${METALLIB_OUTPUT}
+    )
+
+    enable_language(OBJCXX)
+    set_property(SOURCE reduce_metal.hpp PROPERTY LANGUAGE OBJCXX)
+    set_property(SOURCE reduce_bench.cpp PROPERTY LANGUAGE OBJCXX)
+    target_link_libraries(reduce_bench
+        PRIVATE
+        "-framework Metal"
+        "-framework Foundation"
+    )
+    set_source_files_properties(
+        reduce_bench.cpp
+        PROPERTIES
+        COMPILE_FLAGS "-x objective-c++ -fobjc-arc"
+    )
+
+    # Make sure reduce_bench depends on MetalLibBuild so the .metallib is built first
+    add_dependencies(reduce_bench MetalLibBuild)
+
+    # Copy the `metallib` to the same folder as reduce_bench
+    add_custom_command(
+        TARGET reduce_bench
+        POST_BUILD
+        COMMAND ${CMAKE_COMMAND} -E copy_if_different
+        ${METALLIB_OUTPUT}
+        $<TARGET_FILE_DIR:reduce_bench>
+    )
+
 elseif(CMAKE_CXX_COMPILER_ID STREQUAL "Clang")
     message("-- Detected Clang Compiler")
 elseif(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")

README.md

@@ -14,14 +14,14 @@ This repository contains several educational examples showcasing the performance
   - 🔜 NEON and SVE on Arm.
 - OpenMP `reduction` clause.
 - Thrust with its `thrust::reduce`.
-- CUDA kernels with warp-reductions.
+- CUDA kernels with and w/out warp-reductions.
 - OpenCL kernels, eight of them.
 - Parallel STL `<algorithm>` in GCC with Intel oneTBB.
 
 Previously, it also compared ArrayFire, Halide, and Vulkan queues for SPIR-V kernels and SyCL.
 Examples were collected from early 2010s until 2019 and later updated in 2022.
 
-- [Lecture Slides](blob/main/presentation/slides.pdf) from 2019.
+- [Lecture Slides](https://drive.google.com/file/d/16AicAl99t3ZZFnza04Wnw_Vuem0w8lc7/view?usp=sharing) from 2019.
 - [CppRussia Talk](https://youtu.be/AA4RI6o0h1U) in Russia in 2019.
 - [JetBrains Talk](https://youtu.be/BUtHOftDm_Y) in Germany & Russia in 2019.
 
@@ -35,9 +35,9 @@ You are expected to build this on an x86 machine with CUDA drivers installed.
 ```sh
 cmake -B build_release
 cmake --build build_release --config Release
-build_release/reduce_bench # To run all available benchmarks on default array size
-build_release/reduce_bench --benchmark_filter="" # Control Google Benchmark params
-PARALLEL_REDUCTIONS_LENGTH=1000 build_release/reduce_bench # Try different array size
+build_release/reduce_bench                                 # Run all benchmarks
+build_release/reduce_bench --benchmark_filter="cuda"       # Only CUDA-related
+PARALLEL_REDUCTIONS_LENGTH=1000 build_release/reduce_bench # Set a different input size
 ```
 
 Need a more fine-grained control to run only CUDA-based backends?