From 6366150020b1ee73e4d0172236201065a58279ca Mon Sep 17 00:00:00 2001 From: stijn Date: Mon, 10 Aug 2026 09:58:16 +0200 Subject: [PATCH 01/58] Initial commit of new API --- .gitignore | 3 + .gitmodules | 3 + CMakeLists.txt | 161 +- LICENSE | 1 + Makefile | 9 +- README.md => README | 14 +- examples/stencil.cu => docs/.gitkeep | 0 docs/Doxyfile | 762 ++++------ examples/CMakeLists.txt | 115 -- examples/cpp_threads_vector_add.cu | 107 -- examples/histogram.cu | 93 -- examples/matrix_multiply.cu | 136 -- examples/point_in_poly.cu | 112 -- examples/reduction.cu | 167 -- examples/structs.cu | 44 - examples/vector_add.cu | 96 -- external/unordered_dense | 1 + include/kmm/api/access.hpp | 130 -- include/kmm/api/argument.hpp | 92 -- include/kmm/api/array.hpp | 612 +++----- include/kmm/api/array_base.hpp | 36 + include/kmm/api/array_instance.hpp | 35 - include/kmm/api/buffer.hpp | 46 + include/kmm/api/context.hpp | 104 ++ include/kmm/api/device_context.hpp | 20 + include/kmm/api/launch_arg.hpp | 85 ++ include/kmm/api/launcher.hpp | 79 - include/kmm/api/mapper.hpp | 268 ---- include/kmm/api/parallel_submit.hpp | 127 -- include/kmm/api/reduce.hpp | 140 ++ include/kmm/api/runtime_handle.hpp | 261 ---- include/kmm/api/scope.hpp | 73 + include/kmm/api/struct_argument.hpp | 94 -- include/kmm/api/task_group.hpp | 51 - include/kmm/api/view_argument.hpp | 67 - include/kmm/core/backends.hpp | 438 ------ include/kmm/core/bounds.hpp | 300 ++++ include/kmm/core/buffer.hpp | 79 - include/kmm/core/checked_compare.hpp | 570 +++++++ include/kmm/core/checked_math.hpp | 314 ++++ include/kmm/core/commands.hpp | 82 - include/kmm/core/config.hpp | 66 - include/kmm/core/const_value.hpp | 93 ++ include/kmm/core/data_type.hpp | 121 -- include/kmm/core/distribution.hpp | 82 - include/kmm/core/domain.hpp | 69 - include/kmm/core/identifiers.hpp | 303 ---- include/kmm/core/integer_fun.hpp | 77 + include/kmm/core/key_value.hpp | 87 ++ include/kmm/core/layout.hpp | 947 ++++++++++++ include/kmm/{utils => core}/macros.hpp | 85 +- include/kmm/core/panic.hpp | 65 + include/kmm/core/point.hpp | 143 ++ include/kmm/core/range.hpp | 299 ++++ include/kmm/core/reduction.hpp | 38 - include/kmm/core/resource.hpp | 254 ---- include/kmm/core/shape.hpp | 212 +++ include/kmm/core/strides.hpp | 347 +++++ include/kmm/core/type_utils.hpp | 180 +++ include/kmm/core/vec.hpp | 238 +++ include/kmm/core/view.hpp | 1173 ++++---------- include/kmm/kmm.hpp | 11 - include/kmm/memops/gpu_copy.hpp | 43 - include/kmm/memops/gpu_fill.hpp | 10 - include/kmm/memops/gpu_reduction.hpp | 19 - include/kmm/memops/host_copy.hpp | 9 - include/kmm/memops/host_fill.hpp | 9 - include/kmm/memops/host_reduction.hpp | 12 - include/kmm/memops/types.hpp | 73 - include/kmm/planner/array_descriptor.hpp | 50 - include/kmm/planner/read_planner.hpp | 32 - include/kmm/planner/reduction_planner.hpp | 59 - include/kmm/planner/write_planner.hpp | 31 - include/kmm/runtime/allocators/arena.hpp | 83 + include/kmm/runtime/allocators/base.hpp | 112 +- include/kmm/runtime/allocators/block.hpp | 56 - include/kmm/runtime/allocators/caching.hpp | 39 - include/kmm/runtime/allocators/device.hpp | 65 +- .../kmm/runtime/allocators/device_pool.hpp | 59 + include/kmm/runtime/allocators/managed.hpp | 21 + include/kmm/runtime/allocators/pinned.hpp | 18 + include/kmm/runtime/allocators/system.hpp | 12 +- include/kmm/runtime/buffer.hpp | 45 + include/kmm/runtime/buffer_registry.hpp | 57 - include/kmm/runtime/data_interface.hpp | 147 ++ include/kmm/runtime/device_event.hpp | 146 ++ include/kmm/runtime/device_resources.hpp | 54 - include/kmm/runtime/device_stream.hpp | 94 ++ include/kmm/runtime/device_stream_pool.hpp | 70 + .../kmm/runtime/device_stream_registry.hpp | 40 + include/kmm/runtime/executor.hpp | 89 -- include/kmm/runtime/identifiers.hpp | 142 ++ include/kmm/runtime/memops/copy.hpp | 85 ++ include/kmm/runtime/memops/fill.hpp | 108 ++ include/kmm/runtime/memops/reduction.hpp | 117 ++ include/kmm/runtime/memops/types.hpp | 85 ++ include/kmm/runtime/memory_buffer.hpp | 332 ++++ include/kmm/runtime/memory_manager.hpp | 127 +- include/kmm/runtime/memory_system.hpp | 155 +- include/kmm/runtime/reduction_manager.hpp | 65 + include/kmm/runtime/requisition.hpp | 71 + include/kmm/runtime/runtime.hpp | 232 ++- include/kmm/runtime/scheduler.hpp | 70 - include/kmm/runtime/stream_manager.hpp | 247 --- include/kmm/{core => runtime}/system_info.hpp | 93 +- include/kmm/runtime/task_graph.hpp | 73 - include/kmm/utils/checked_math.hpp | 392 ----- include/kmm/utils/fixed_array.hpp | 275 ---- include/kmm/utils/function_ref.hpp | 52 + include/kmm/utils/geometry.hpp | 669 -------- include/kmm/utils/gpu_utils.hpp | 155 +- include/kmm/utils/hash_utils.hpp | 45 +- include/kmm/utils/integer_fun.hpp | 104 -- include/kmm/utils/key_value.hpp | 43 - include/kmm/utils/lru_cache.hpp | 116 ++ include/kmm/utils/notify.hpp | 94 +- include/kmm/utils/panic.hpp | 41 - include/kmm/utils/poll.hpp | 7 +- include/kmm/utils/range.hpp | 173 --- include/kmm/utils/refcnt_ptr.hpp | 213 +++ include/kmm/utils/small_vector.hpp | 151 +- src/api/array.cpp | 3 - src/api/array_instance.cpp | 76 - src/api/buffer.cpp | 95 ++ src/api/mapper.cpp | 154 -- src/api/runtime_handle.cpp | 89 -- src/core/backends.cpp | 310 ---- src/core/checked_compare.cpp | 11 + src/core/config.cpp | 118 -- src/core/data_type.cpp | 139 -- src/core/distribution.cpp | 204 --- src/core/domain.cpp | 66 - src/core/identifiers.cpp | 73 - src/{utils => core}/panic.cpp | 22 +- src/core/reduction.cpp | 28 - src/core/resource.cpp | 72 - src/core/system_info.cpp | 152 -- src/core/task.cpp | 1 - src/memops/gpu_copy.cpp | 215 --- src/memops/gpu_fill.cu | 110 -- src/memops/gpu_operators.cuh | 126 -- src/memops/gpu_reduction.cu | 311 ---- src/memops/host_copy.cpp | 60 - src/memops/host_fill.cpp | 19 - src/memops/host_operators.hpp | 139 -- src/memops/host_reduction.cpp | 163 -- src/memops/types.cpp | 247 --- src/planner/array_descriptor.cpp | 224 --- src/planner/read_planner.cpp | 63 - src/planner/reduction_planner.cpp | 249 --- src/planner/write_planner.cpp | 88 -- src/runtime/allocators/arena.cpp | 217 +++ src/runtime/allocators/base.cpp | 97 +- src/runtime/allocators/block.cpp | 284 ---- src/runtime/allocators/caching.cpp | 191 --- src/runtime/allocators/device.cpp | 202 +-- src/runtime/allocators/device_pool.cpp | 191 +++ src/runtime/allocators/managed.cpp | 32 + src/runtime/allocators/pinned.cpp | 34 + src/runtime/allocators/system.cpp | 9 +- src/runtime/buffer_registry.cpp | 151 -- src/runtime/data_interface.cpp | 186 +++ src/runtime/device_event.cpp | 151 ++ src/runtime/device_resources.cpp | 155 -- src/runtime/device_stream.cpp | 474 ++++++ src/runtime/device_stream_pool.cpp | 93 ++ src/runtime/device_stream_registry.cpp | 82 + src/runtime/executor.cpp | 716 --------- src/runtime/memops/copy.cpp | 141 ++ src/runtime/memops/fill.cpp | 165 ++ src/runtime/memops/reduction.cpp | 170 +++ src/runtime/memops/reduction.cu | 387 +++++ src/runtime/memops/simplify_dims.hpp | 69 + src/runtime/memops/types.cpp | 73 + src/runtime/memory_buffer.cpp | 381 +++++ src/runtime/memory_manager.cpp | 1346 +++++------------ src/runtime/memory_system.cpp | 298 ++-- src/runtime/reduction_manager.cpp | 470 ++++++ src/runtime/requisition.cpp | 39 + src/runtime/runtime.cpp | 531 ++++--- src/runtime/scheduler.cpp | 213 --- src/runtime/stream_manager.cpp | 583 ------- src/runtime/system_info.cpp | 144 ++ src/runtime/task_graph.cpp | 83 - src/utils/checked_math.cpp | 11 - src/utils/gpu_utils.cpp | 148 +- src/utils/notify.cpp | 17 +- src/utils/small_vector.cpp | 7 + test/CMakeLists.txt | 9 +- test/core/test_bounds.cpp | 179 +++ test/core/test_checked_compare.cpp | 338 +++++ test/core/test_checked_math.cpp | 327 ++++ test/core/test_const_value.cpp | 40 + test/core/test_integer_fun.cpp | 176 +++ test/core/test_key_value.cpp | 57 + test/core/test_layout.cpp | 512 +++++++ test/core/test_panic.cpp | 15 + test/core/test_point.cpp | 131 ++ test/core/test_range.cpp | 189 +++ test/core/test_shape.cpp | 167 ++ test/core/test_strides.cpp | 309 ++++ test/core/test_type_utils.cpp | 91 ++ test/core/test_vec.cpp | 137 ++ test/core/test_view.cpp | 272 ++++ test/dag/test_data_distribution.cpp | 104 -- test/runtime/test_memory_manager.cpp | 17 + test/runtime/test_stream_manager.cpp | 9 + test/utils/test_checked_math.cpp | 250 --- test/utils/test_function_ref.cpp | 32 + test/utils/test_geometry.cpp | 262 ---- test/utils/test_integer_fun.cpp | 138 -- test/utils/test_intrusive_ptr.cpp | 5 + test/utils/test_lru_cache.cpp | 104 ++ test/utils/test_range.cpp | 213 --- test/utils/test_small_vector.cpp | 349 +++-- test/utils/test_view.cpp | 370 ----- 216 files changed, 15965 insertions(+), 17908 deletions(-) rename README.md => README (76%) rename examples/stencil.cu => docs/.gitkeep (100%) delete mode 100644 examples/CMakeLists.txt delete mode 100644 examples/cpp_threads_vector_add.cu delete mode 100644 examples/histogram.cu delete mode 100644 examples/matrix_multiply.cu delete mode 100644 examples/point_in_poly.cu delete mode 100644 examples/reduction.cu delete mode 100644 examples/structs.cu delete mode 100644 examples/vector_add.cu create mode 160000 external/unordered_dense delete mode 100644 include/kmm/api/access.hpp delete mode 100644 include/kmm/api/argument.hpp create mode 100644 include/kmm/api/array_base.hpp delete mode 100644 include/kmm/api/array_instance.hpp create mode 100644 include/kmm/api/buffer.hpp create mode 100644 include/kmm/api/context.hpp create mode 100644 include/kmm/api/device_context.hpp create mode 100644 include/kmm/api/launch_arg.hpp delete mode 100644 include/kmm/api/launcher.hpp delete mode 100644 include/kmm/api/mapper.hpp delete mode 100644 include/kmm/api/parallel_submit.hpp create mode 100644 include/kmm/api/reduce.hpp delete mode 100644 include/kmm/api/runtime_handle.hpp create mode 100644 include/kmm/api/scope.hpp delete mode 100644 include/kmm/api/struct_argument.hpp delete mode 100644 include/kmm/api/task_group.hpp delete mode 100644 include/kmm/api/view_argument.hpp delete mode 100644 include/kmm/core/backends.hpp create mode 100644 include/kmm/core/bounds.hpp delete mode 100644 include/kmm/core/buffer.hpp create mode 100644 include/kmm/core/checked_compare.hpp create mode 100644 include/kmm/core/checked_math.hpp delete mode 100644 include/kmm/core/commands.hpp delete mode 100644 include/kmm/core/config.hpp create mode 100644 include/kmm/core/const_value.hpp delete mode 100644 include/kmm/core/data_type.hpp delete mode 100644 include/kmm/core/distribution.hpp delete mode 100644 include/kmm/core/domain.hpp delete mode 100644 include/kmm/core/identifiers.hpp create mode 100644 include/kmm/core/integer_fun.hpp create mode 100644 include/kmm/core/key_value.hpp create mode 100644 include/kmm/core/layout.hpp rename include/kmm/{utils => core}/macros.hpp (67%) create mode 100644 include/kmm/core/panic.hpp create mode 100644 include/kmm/core/point.hpp create mode 100644 include/kmm/core/range.hpp delete mode 100644 include/kmm/core/reduction.hpp delete mode 100644 include/kmm/core/resource.hpp create mode 100644 include/kmm/core/shape.hpp create mode 100644 include/kmm/core/strides.hpp create mode 100644 include/kmm/core/type_utils.hpp create mode 100644 include/kmm/core/vec.hpp delete mode 100644 include/kmm/kmm.hpp delete mode 100644 include/kmm/memops/gpu_copy.hpp delete mode 100644 include/kmm/memops/gpu_fill.hpp delete mode 100644 include/kmm/memops/gpu_reduction.hpp delete mode 100644 include/kmm/memops/host_copy.hpp delete mode 100644 include/kmm/memops/host_fill.hpp delete mode 100644 include/kmm/memops/host_reduction.hpp delete mode 100644 include/kmm/memops/types.hpp delete mode 100644 include/kmm/planner/array_descriptor.hpp delete mode 100644 include/kmm/planner/read_planner.hpp delete mode 100644 include/kmm/planner/reduction_planner.hpp delete mode 100644 include/kmm/planner/write_planner.hpp create mode 100644 include/kmm/runtime/allocators/arena.hpp delete mode 100644 include/kmm/runtime/allocators/block.hpp delete mode 100644 include/kmm/runtime/allocators/caching.hpp create mode 100644 include/kmm/runtime/allocators/device_pool.hpp create mode 100644 include/kmm/runtime/allocators/managed.hpp create mode 100644 include/kmm/runtime/allocators/pinned.hpp create mode 100644 include/kmm/runtime/buffer.hpp delete mode 100644 include/kmm/runtime/buffer_registry.hpp create mode 100644 include/kmm/runtime/data_interface.hpp create mode 100644 include/kmm/runtime/device_event.hpp delete mode 100644 include/kmm/runtime/device_resources.hpp create mode 100644 include/kmm/runtime/device_stream.hpp create mode 100644 include/kmm/runtime/device_stream_pool.hpp create mode 100644 include/kmm/runtime/device_stream_registry.hpp delete mode 100644 include/kmm/runtime/executor.hpp create mode 100644 include/kmm/runtime/identifiers.hpp create mode 100644 include/kmm/runtime/memops/copy.hpp create mode 100644 include/kmm/runtime/memops/fill.hpp create mode 100644 include/kmm/runtime/memops/reduction.hpp create mode 100644 include/kmm/runtime/memops/types.hpp create mode 100644 include/kmm/runtime/memory_buffer.hpp create mode 100644 include/kmm/runtime/reduction_manager.hpp create mode 100644 include/kmm/runtime/requisition.hpp delete mode 100644 include/kmm/runtime/scheduler.hpp delete mode 100644 include/kmm/runtime/stream_manager.hpp rename include/kmm/{core => runtime}/system_info.hpp (51%) delete mode 100644 include/kmm/runtime/task_graph.hpp delete mode 100644 include/kmm/utils/checked_math.hpp delete mode 100644 include/kmm/utils/fixed_array.hpp create mode 100644 include/kmm/utils/function_ref.hpp delete mode 100644 include/kmm/utils/geometry.hpp delete mode 100644 include/kmm/utils/integer_fun.hpp delete mode 100644 include/kmm/utils/key_value.hpp create mode 100644 include/kmm/utils/lru_cache.hpp delete mode 100644 include/kmm/utils/panic.hpp delete mode 100644 include/kmm/utils/range.hpp create mode 100644 include/kmm/utils/refcnt_ptr.hpp delete mode 100644 src/api/array.cpp delete mode 100644 src/api/array_instance.cpp create mode 100644 src/api/buffer.cpp delete mode 100644 src/api/mapper.cpp delete mode 100644 src/api/runtime_handle.cpp delete mode 100644 src/core/backends.cpp create mode 100644 src/core/checked_compare.cpp delete mode 100644 src/core/config.cpp delete mode 100644 src/core/data_type.cpp delete mode 100644 src/core/distribution.cpp delete mode 100644 src/core/domain.cpp delete mode 100644 src/core/identifiers.cpp rename src/{utils => core}/panic.cpp (63%) delete mode 100644 src/core/reduction.cpp delete mode 100644 src/core/resource.cpp delete mode 100644 src/core/system_info.cpp delete mode 100644 src/core/task.cpp delete mode 100644 src/memops/gpu_copy.cpp delete mode 100644 src/memops/gpu_fill.cu delete mode 100644 src/memops/gpu_operators.cuh delete mode 100644 src/memops/gpu_reduction.cu delete mode 100644 src/memops/host_copy.cpp delete mode 100644 src/memops/host_fill.cpp delete mode 100644 src/memops/host_operators.hpp delete mode 100644 src/memops/host_reduction.cpp delete mode 100644 src/memops/types.cpp delete mode 100644 src/planner/array_descriptor.cpp delete mode 100644 src/planner/read_planner.cpp delete mode 100644 src/planner/reduction_planner.cpp delete mode 100644 src/planner/write_planner.cpp create mode 100644 src/runtime/allocators/arena.cpp delete mode 100644 src/runtime/allocators/block.cpp delete mode 100644 src/runtime/allocators/caching.cpp create mode 100644 src/runtime/allocators/device_pool.cpp create mode 100644 src/runtime/allocators/managed.cpp create mode 100644 src/runtime/allocators/pinned.cpp delete mode 100644 src/runtime/buffer_registry.cpp create mode 100644 src/runtime/data_interface.cpp create mode 100644 src/runtime/device_event.cpp delete mode 100644 src/runtime/device_resources.cpp create mode 100644 src/runtime/device_stream.cpp create mode 100644 src/runtime/device_stream_pool.cpp create mode 100644 src/runtime/device_stream_registry.cpp delete mode 100644 src/runtime/executor.cpp create mode 100644 src/runtime/memops/copy.cpp create mode 100644 src/runtime/memops/fill.cpp create mode 100644 src/runtime/memops/reduction.cpp create mode 100644 src/runtime/memops/reduction.cu create mode 100644 src/runtime/memops/simplify_dims.hpp create mode 100644 src/runtime/memops/types.cpp create mode 100644 src/runtime/memory_buffer.cpp create mode 100644 src/runtime/reduction_manager.cpp create mode 100644 src/runtime/requisition.cpp delete mode 100644 src/runtime/scheduler.cpp delete mode 100644 src/runtime/stream_manager.cpp create mode 100644 src/runtime/system_info.cpp delete mode 100644 src/runtime/task_graph.cpp delete mode 100644 src/utils/checked_math.cpp create mode 100644 src/utils/small_vector.cpp create mode 100644 test/core/test_bounds.cpp create mode 100644 test/core/test_checked_compare.cpp create mode 100644 test/core/test_checked_math.cpp create mode 100644 test/core/test_const_value.cpp create mode 100644 test/core/test_integer_fun.cpp create mode 100644 test/core/test_key_value.cpp create mode 100644 test/core/test_layout.cpp create mode 100644 test/core/test_panic.cpp create mode 100644 test/core/test_point.cpp create mode 100644 test/core/test_range.cpp create mode 100644 test/core/test_shape.cpp create mode 100644 test/core/test_strides.cpp create mode 100644 test/core/test_type_utils.cpp create mode 100644 test/core/test_vec.cpp create mode 100644 test/core/test_view.cpp delete mode 100644 test/dag/test_data_distribution.cpp create mode 100644 test/runtime/test_memory_manager.cpp create mode 100644 test/runtime/test_stream_manager.cpp delete mode 100644 test/utils/test_checked_math.cpp create mode 100644 test/utils/test_function_ref.cpp delete mode 100644 test/utils/test_geometry.cpp delete mode 100644 test/utils/test_integer_fun.cpp create mode 100644 test/utils/test_intrusive_ptr.cpp create mode 100644 test/utils/test_lru_cache.cpp delete mode 100644 test/utils/test_range.cpp delete mode 100644 test/utils/test_view.cpp diff --git a/.gitignore b/.gitignore index a8500b20..a1e172af 100644 --- a/.gitignore +++ b/.gitignore @@ -5,3 +5,6 @@ cmake-build* # Editors .idea/ .vscode/ + +# Generated docs +docs/html/ diff --git a/.gitmodules b/.gitmodules index 01d96300..8040d975 100644 --- a/.gitmodules +++ b/.gitmodules @@ -7,3 +7,6 @@ [submodule "external/Catch2"] path = external/Catch2 url = https://github.com/catchorg/Catch2.git +[submodule "external/unordered_dense"] + path = external/unordered_dense + url = https://github.com/martinus/unordered_dense diff --git a/CMakeLists.txt b/CMakeLists.txt index c5cebf03..a028dbe3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,59 +1,41 @@ -cmake_minimum_required(VERSION 3.10) - -# Project setup -set(PROJECT_NAME "kmm") -project(${PROJECT_NAME} LANGUAGES CXX VERSION 0.3) - -# User options (features) -option(KMM_STATIC "Build a static library" OFF) -option(KMM_USE_CUDA "Build the CUDA backend" OFF) -option(KMM_USE_HIP "Build the HIP backend" OFF) - -# User options (development) -option(KMM_ENABLE_LINTER "Enable clang-tidy linter" OFF) -option(KMM_BUILD_TESTS "Build tests" OFF) -option(KMM_BUILD_EXAMPLES "Build examples" OFF) -option(KMM_BUILD_BENCHMARKS "Build benchmarks" OFF) - -# Check options -if(KMM_USE_CUDA AND KMM_USE_HIP) - message(FATAL_ERROR "CUDA and HIP backend are mutually exclusive.") -endif() +cmake_minimum_required(VERSION 3.18) -# Enable C++17 support -set(CMAKE_CXX_STANDARD 17) -set(CMAKE_CXX_STANDARD_REQUIRED ON) +project(kmm VERSION 0.1.0 LANGUAGES CXX) -file(GLOB_RECURSE sources - "${PROJECT_SOURCE_DIR}/src/*.cpp" - "${PROJECT_SOURCE_DIR}/src/*.cu" -) +option(KMM_BUILD_TESTS "Build kmm tests" ON) +option(KMM_USE_CUDA "Enable CUDA support" OFF) -# Create library -if(KMM_STATIC) - add_library(${PROJECT_NAME} STATIC ${sources}) -else() - add_library(${PROJECT_NAME} SHARED ${sources}) -endif() +file(GLOB_RECURSE sources src/*.cpp) + +add_library(kmm ${sources}) -target_include_directories(${PROJECT_NAME} PUBLIC "${PROJECT_SOURCE_DIR}/include") +target_include_directories(kmm PUBLIC ${PROJECT_SOURCE_DIR}/include) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_EXTENSIONS OFF) +target_compile_features(kmm PUBLIC cxx_std_17) +target_compile_options(kmm PRIVATE + -Wall -Wextra -Wconversion -Wno-unused-parameter + $<$:-Wpedantic> +) if(KMM_USE_CUDA) - enable_language(CUDA) - SET(CUDA_SEPARABLE_COMPILATION ON) - set(CMAKE_CUDA_SEPARABLE_COMPILATION ON) -elseif(KMM_USE_HIP) - enable_language(HIP) + enable_language(CUDA) + find_package(CUDAToolkit REQUIRED) + set_target_properties(kmm PROPERTIES + CUDA_STANDARD 17 + CUDA_STANDARD_REQUIRED ON + CUDA_EXTENSIONS OFF + ) + + file(GLOB_RECURSE cuda_sources src/*.cu) + target_sources(kmm PRIVATE ${cuda_sources}) + + target_link_libraries(kmm PUBLIC CUDA::cudart CUDA::cuda_driver) + target_compile_definitions(kmm PUBLIC KMM_USE_CUDA) endif() -# CXX flags -target_compile_options(${PROJECT_NAME} - PUBLIC - $<$:-forward-unknown-to-host-compiler> - -Wall -Wextra -Wconversion -Wno-unused-parameter - #$<$:-Xcompiler=-Werror> -) -target_compile_options(${PROJECT_NAME} PUBLIC ${CXXFLAGS}) # Enable PIC set(CMAKE_POSITION_INDEPENDENT_CODE ON) @@ -70,89 +52,12 @@ set(SPDLOG_BUILD_PIC ON) add_subdirectory(external/spdlog) target_link_libraries(${PROJECT_NAME} PUBLIC spdlog) -# Install -include(GNUInstallDirs) -set_target_properties( - ${PROJECT_NAME} - PROPERTIES - VERSION ${PROJECT_VERSION} - SOVERSION 1 -) +# Dependencies: unordered_dense +add_subdirectory(external/unordered_dense) +target_link_libraries(${PROJECT_NAME} PUBLIC unordered_dense::unordered_dense) -install( - TARGETS ${PROJECT_NAME} - LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR} - ARCHIVE DESTINATION ${CMAKE_INSTALL_LIBDIR} - RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} -) - -install( - DIRECTORY "${PROJECT_SOURCE_DIR}/include/kmm" - DESTINATION ${CMAKE_INSTALL_INCLUDEDIR} - FILES_MATCHING PATTERN "*.hpp" -) - -if(KMM_ENABLE_LINTER) - set(PROJECT_CLANG_TIDY clang-tidy) - set_target_properties(${PROJECT_NAME} PROPERTIES CXX_CLANG_TIDY "${PROJECT_CLANG_TIDY}") -endif() - -if(KMM_USE_CUDA) - find_package(CUDAToolkit REQUIRED) - target_link_libraries(${PROJECT_NAME} - PUBLIC - CUDA::cudart_static - CUDA::cuda_driver - CUDA::cublas - CUDA::nvrtc - ) - - set_target_properties( - ${PROJECT_NAME} - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) - - # Define `KMM_USE_CUDA` macro so that headers can detect CUDA usage - target_compile_definitions(${PROJECT_NAME} PUBLIC KMM_USE_CUDA=1) -elseif(KMM_USE_HIP) - if(NOT DEFINED HIP_PATH) - if(NOT DEFINED ENV{HIP_PATH}) - set(HIP_PATH "/opt/rocm/hip" CACHE PATH "Path to which HIP has been installed") - else() - set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed") - endif() - endif() - set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH}) - find_package(HIP REQUIRED) - find_package(ROCBLAS REQUIRED) - set_source_files_properties(${sources} PROPERTIES LANGUAGE HIP) - target_link_libraries(${PROJECT_NAME} - PUBLIC - hip::host - hip::device - roc::rocblas - ) - - # Define `KMM_USE_HIP` macro so that headers can detect HIP usage - target_compile_definitions(${PROJECT_NAME} PUBLIC KMM_USE_HIP=1) -endif() - -# Compile unit tests if(KMM_BUILD_TESTS) - include(CTest) add_subdirectory(external/Catch2) add_subdirectory(test) endif() -# Compile examples -if(KMM_BUILD_EXAMPLES) - add_subdirectory(examples) -endif() - -# Compile benchmarks -if(KMM_BUILD_BENCHMARKS) - add_subdirectory(benchmarks) -endif() \ No newline at end of file diff --git a/LICENSE b/LICENSE index 261eeb9e..57bc88a1 100644 --- a/LICENSE +++ b/LICENSE @@ -199,3 +199,4 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the License for the specific language governing permissions and limitations under the License. + diff --git a/Makefile b/Makefile index 3040c8e8..f7951c40 100644 --- a/Makefile +++ b/Makefile @@ -1,12 +1,15 @@ CLANG_FORMAT=clang-format-16 --verbose pretty: - ${CLANG_FORMAT} -i include/kmm/*.hpp include/kmm/*/*.hpp - $(CLANG_FORMAT) -i src/*/*.cpp src/*/*.cu src/*/*.cuh + ${CLANG_FORMAT} -i include/kmm/*/*.hpp ${CLANG_FORMAT} -i test/*/*.cpp + $(CLANG_FORMAT) -i src/*/*.cpp src/*/*.cu src/*/*.cuh ${CLANG_FORMAT} -i examples/*.cu ${CLANG_FORMAT} -i benchmarks/*.cu +docs: + cd docs && doxygen Doxyfile + all: pretty -.PHONY : pretty +.PHONY : pretty docs diff --git a/README.md b/README similarity index 76% rename from README.md rename to README index aaebd514..c7aac52d 100644 --- a/README.md +++ b/README @@ -1,11 +1,10 @@ +![KMM: Kernel Memory Manager](https://raw.githubusercontent.com/NLeSC/kmm/refs/heads/main/docs/_static/kmm-logo.png) -![KMM Logo](https://raw.githubusercontent.com/NLeSC-COMPAS/kmm/refs/heads/main/docs/_static/kmm-logo.png) -# KMM: Kernel Memory Manager - -[![CPU Build Status](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-multi-compiler.yml/badge.svg)](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-multi-compiler.yml) -[![CUDA Build Status](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-cuda-multi-compiler.yml/badge.svg)](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-cuda-multi-compiler.yml) -[![HIP Build Status](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-hip.yml/badge.svg)](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-hip.yml) +# +[![CPU Build Status](https://github.com/NLeSC/kmm/actions/workflows/cmake-multi-compiler.yml/badge.svg)](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-multi-compiler.yml) +[![CUDA Build Status](https://github.com/NLeSC/kmm/actions/workflows/cmake-cuda-multi-compiler.yml/badge.svg)](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-cuda-multi-compiler.yml) +[![HIP Build Status](https://github.com/NLeSC/kmm/actions/workflows/cmake-hip.yml/badge.svg)](https://github.com/NLeSC-COMPAS/kmm/actions/workflows/cmake-hip.yml) The **Kernel Memory Manager** (KMM) is a lightweight, high-performance framework designed for parallel dataflow execution and efficient memory management on multi-GPU platforms. @@ -26,7 +25,7 @@ Unlike frameworks that require a specific programming model, KMM integrates exis ## Resources -* [Full documentation](https://nlesc-compas.github.io/kmm) +* [Full documentation](https://nlesc.github.io/kmm) ## Example @@ -86,3 +85,4 @@ int main() { ## License KMM is made available under the terms of the Apache License version 2.0, see the file LICENSE for details. + diff --git a/examples/stencil.cu b/docs/.gitkeep similarity index 100% rename from examples/stencil.cu rename to docs/.gitkeep diff --git a/docs/Doxyfile b/docs/Doxyfile index 7cd416fe..9aced25a 100644 --- a/docs/Doxyfile +++ b/docs/Doxyfile @@ -1,4 +1,4 @@ -# Doxyfile 1.10.0 +# Doxyfile 1.9.1 # This file describes the settings to be used by the documentation system # doxygen (www.doxygen.org) for a project. @@ -12,16 +12,6 @@ # For lists, items can also be appended using: # TAG += value [value, ...] # Values that contain spaces should be placed between quotes (\" \"). -# -# Note: -# -# Use doxygen to compare the used configuration file with the template -# configuration file: -# doxygen -x [configFile] -# Use doxygen to compare the used configuration file with the template -# configuration file without replacing the environment variables or CMake type -# replacement variables: -# doxygen -x_noenv [configFile] #--------------------------------------------------------------------------- # Project related configuration options @@ -42,19 +32,19 @@ DOXYFILE_ENCODING = UTF-8 # title of most generated pages and in a few other places. # The default value is: My Project. -PROJECT_NAME = "KMM" +PROJECT_NAME = "kmm" # The PROJECT_NUMBER tag can be used to enter a project or revision number. This # could be handy for archiving the generated documentation or if some version # control system is used. -PROJECT_NUMBER = +PROJECT_NUMBER = 0.1.0 # Using the PROJECT_BRIEF tag one can provide an optional one line description # for a project that appears at the top of each page and should give viewer a # quick idea about the purpose of the project. Keep the description short. -PROJECT_BRIEF = +PROJECT_BRIEF = "Kernel launcher framework for heterogeneous systems" # With the PROJECT_LOGO tag one can specify a logo or an icon that is included # in the documentation. The maximum height of the logo should not exceed 55 @@ -63,41 +53,23 @@ PROJECT_BRIEF = PROJECT_LOGO = -# With the PROJECT_ICON tag one can specify an icon that is included in the tabs -# when the HTML document is shown. Doxygen will copy the logo to the output -# directory. - -PROJECT_ICON = - # The OUTPUT_DIRECTORY tag is used to specify the (relative or absolute) path # into which the generated documentation will be written. If a relative path is # entered, it will be relative to the location where doxygen was started. If # left blank the current directory will be used. -OUTPUT_DIRECTORY = _doxygen +OUTPUT_DIRECTORY = -# If the CREATE_SUBDIRS tag is set to YES then doxygen will create up to 4096 -# sub-directories (in 2 levels) under the output directory of each output format -# and will distribute the generated files over these directories. Enabling this +# If the CREATE_SUBDIRS tag is set to YES then doxygen will create 4096 sub- +# directories (in 2 levels) under the output directory of each output format and +# will distribute the generated files over these directories. Enabling this # option can be useful when feeding doxygen a huge amount of source files, where # putting all generated files in the same directory would otherwise causes -# performance problems for the file system. Adapt CREATE_SUBDIRS_LEVEL to -# control the number of sub-directories. +# performance problems for the file system. # The default value is: NO. CREATE_SUBDIRS = NO -# Controls the number of sub-directories that will be created when -# CREATE_SUBDIRS tag is set to YES. Level 0 represents 16 directories, and every -# level increment doubles the number of directories, resulting in 4096 -# directories at level 8 which is the default and also the maximum value. The -# sub-directories are organized in 2 levels, the first level always has a fixed -# number of 16 directories. -# Minimum value: 0, maximum value: 8, default value: 8. -# This tag requires that the tag CREATE_SUBDIRS is set to YES. - -CREATE_SUBDIRS_LEVEL = 8 - # If the ALLOW_UNICODE_NAMES tag is set to YES, doxygen will allow non-ASCII # characters to appear in the names of generated files. If set to NO, non-ASCII # characters will be escaped, for example _xE3_x81_x84 will be used for Unicode @@ -109,18 +81,26 @@ ALLOW_UNICODE_NAMES = NO # The OUTPUT_LANGUAGE tag is used to specify the language in which all # documentation generated by doxygen is written. Doxygen will use this # information to generate all constant output in the proper language. -# Possible values are: Afrikaans, Arabic, Armenian, Brazilian, Bulgarian, -# Catalan, Chinese, Chinese-Traditional, Croatian, Czech, Danish, Dutch, English -# (United States), Esperanto, Farsi (Persian), Finnish, French, German, Greek, -# Hindi, Hungarian, Indonesian, Italian, Japanese, Japanese-en (Japanese with -# English messages), Korean, Korean-en (Korean with English messages), Latvian, -# Lithuanian, Macedonian, Norwegian, Persian (Farsi), Polish, Portuguese, -# Romanian, Russian, Serbian, Serbian-Cyrillic, Slovak, Slovene, Spanish, -# Swedish, Turkish, Ukrainian and Vietnamese. +# Possible values are: Afrikaans, Arabic, Armenian, Brazilian, Catalan, Chinese, +# Chinese-Traditional, Croatian, Czech, Danish, Dutch, English (United States), +# Esperanto, Farsi (Persian), Finnish, French, German, Greek, Hungarian, +# Indonesian, Italian, Japanese, Japanese-en (Japanese with English messages), +# Korean, Korean-en (Korean with English messages), Latvian, Lithuanian, +# Macedonian, Norwegian, Persian (Farsi), Polish, Portuguese, Romanian, Russian, +# Serbian, Serbian-Cyrillic, Slovak, Slovene, Spanish, Swedish, Turkish, +# Ukrainian and Vietnamese. # The default value is: English. OUTPUT_LANGUAGE = English +# The OUTPUT_TEXT_DIRECTION tag is used to specify the direction in which all +# documentation generated by doxygen is written. Doxygen will use this +# information to generate all generated output in the proper direction. +# Possible values are: None, LTR, RTL and Context. +# The default value is: None. + +OUTPUT_TEXT_DIRECTION = None + # If the BRIEF_MEMBER_DESC tag is set to YES, doxygen will include brief member # descriptions after the members that are listed in the file and class # documentation (similar to Javadoc). Set to NO to disable this. @@ -278,16 +258,16 @@ TAB_SIZE = 4 # the documentation. An alias has the form: # name=value # For example adding -# "sideeffect=@par Side Effects:^^" +# "sideeffect=@par Side Effects:\n" # will allow you to put the command \sideeffect (or @sideeffect) in the # documentation, which will result in a user-defined paragraph with heading -# "Side Effects:". Note that you cannot put \n's in the value part of an alias -# to insert newlines (in the resulting output). You can put ^^ in the value part -# of an alias to insert a newline as if a physical newline was in the original -# file. When you need a literal { or } or , in the value part of an alias you -# have to escape them by means of a backslash (\), this can lead to conflicts -# with the commands \{ and \} for these it is advised to use the version @{ and -# @} or use a double escape (\\{ and \\}) +# "Side Effects:". You can put \n's in the value part of an alias to insert +# newlines (in the resulting output). You can put ^^ in the value part of an +# alias to insert a newline as if a physical newline was in the original file. +# When you need a literal { or } or , in the value part of an alias you have to +# escape them by means of a backslash (\), this can lead to conflicts with the +# commands \{ and \} for these it is advised to use the version @{ and @} or use +# a double escape (\\{ and \\}) ALIASES = @@ -332,8 +312,8 @@ OPTIMIZE_OUTPUT_SLICE = NO # extension. Doxygen has a built-in mapping, but you can override or extend it # using this tag. The format is ext=language, where ext is a file extension, and # language is one of the parsers supported by doxygen: IDL, Java, JavaScript, -# Csharp (C#), C, C++, Lex, D, PHP, md (Markdown), Objective-C, Python, Slice, -# VHDL, Fortran (fixed format Fortran: FortranFixed, free formatted Fortran: +# Csharp (C#), C, C++, D, PHP, md (Markdown), Objective-C, Python, Slice, VHDL, +# Fortran (fixed format Fortran: FortranFixed, free formatted Fortran: # FortranFree, unknown formatted Fortran: Fortran. In the later case the parser # tries to guess whether the code is fixed or free formatted code, this is the # default for Fortran type files). For instance to make doxygen treat .inc files @@ -369,17 +349,6 @@ MARKDOWN_SUPPORT = YES TOC_INCLUDE_HEADINGS = 5 -# The MARKDOWN_ID_STYLE tag can be used to specify the algorithm used to -# generate identifiers for the Markdown headings. Note: Every identifier is -# unique. -# Possible values are: DOXYGEN use a fixed 'autotoc_md' string followed by a -# sequence number starting at 0 and GITHUB use the lower case version of title -# with any whitespace replaced by '-' and punctuation characters removed. -# The default value is: DOXYGEN. -# This tag requires that the tag MARKDOWN_SUPPORT is set to YES. - -MARKDOWN_ID_STYLE = DOXYGEN - # When enabled doxygen tries to link words that correspond to documented # classes, or namespaces to their corresponding documentation. Such a link can # be prevented in individual cases by putting a % sign in front of the word or @@ -491,27 +460,19 @@ TYPEDEF_HIDES_STRUCT = NO LOOKUP_CACHE_SIZE = 0 -# The NUM_PROC_THREADS specifies the number of threads doxygen is allowed to use +# The NUM_PROC_THREADS specifies the number threads doxygen is allowed to use # during processing. When set to 0 doxygen will based this on the number of # cores available in the system. You can set it explicitly to a value larger # than 0 to get more control over the balance between CPU load and processing # speed. At this moment only the input processing can be done using multiple # threads. Since this is still an experimental feature the default is set to 1, -# which effectively disables parallel processing. Please report any issues you +# which efficively disables parallel processing. Please report any issues you # encounter. Generating dot graphs in parallel is controlled by the # DOT_NUM_THREADS setting. # Minimum value: 0, maximum value: 32, default value: 1. NUM_PROC_THREADS = 1 -# If the TIMESTAMP tag is set different from NO then each generated page will -# contain the date or date and time when the page was generated. Setting this to -# NO can help when comparing the output of multiple runs. -# Possible values are: YES, NO, DATETIME and DATE. -# The default value is: NO. - -TIMESTAMP = NO - #--------------------------------------------------------------------------- # Build related configuration options #--------------------------------------------------------------------------- @@ -588,16 +549,15 @@ RESOLVE_UNNAMED_PARAMS = YES # section is generated. This option has no effect if EXTRACT_ALL is enabled. # The default value is: NO. -HIDE_UNDOC_MEMBERS = NO +HIDE_UNDOC_MEMBERS = YES # If the HIDE_UNDOC_CLASSES tag is set to YES, doxygen will hide all # undocumented classes that are normally visible in the class hierarchy. If set # to NO, these classes will be included in the various overviews. This option -# will also hide undocumented C++ concepts if enabled. This option has no effect -# if EXTRACT_ALL is enabled. +# has no effect if EXTRACT_ALL is enabled. # The default value is: NO. -HIDE_UNDOC_CLASSES = NO +HIDE_UNDOC_CLASSES = YES # If the HIDE_FRIEND_COMPOUNDS tag is set to YES, doxygen will hide all friend # declarations. If set to NO, these declarations will be included in the @@ -625,17 +585,16 @@ INTERNAL_DOCS = NO # filesystem is case sensitive (i.e. it supports files in the same directory # whose names only differ in casing), the option must be set to YES to properly # deal with such files in case they appear in the input. For filesystems that -# are not case sensitive the option should be set to NO to properly deal with +# are not case sensitive the option should be be set to NO to properly deal with # output files written for symbols that only differ in casing, such as for two # classes, one named CLASS and the other named Class, and to also support # references to files without having to specify the exact matching casing. On # Windows (including Cygwin) and MacOS, users should typically set this option # to NO, whereas on Linux or other Unix flavors it should typically be set to # YES. -# Possible values are: SYSTEM, NO and YES. -# The default value is: SYSTEM. +# The default value is: system dependent. -CASE_SENSE_NAMES = SYSTEM +CASE_SENSE_NAMES = YES # If the HIDE_SCOPE_NAMES tag is set to NO then doxygen will show members with # their full class and namespace scopes in the documentation. If set to YES, the @@ -651,12 +610,6 @@ HIDE_SCOPE_NAMES = NO HIDE_COMPOUND_REFERENCE= NO -# If the SHOW_HEADERFILE tag is set to YES then the documentation for a class -# will show which file needs to be included to use the class. -# The default value is: YES. - -SHOW_HEADERFILE = YES - # If the SHOW_INCLUDE_FILES tag is set to YES then doxygen will put a list of # the files that are included by a file in the documentation of that file. # The default value is: YES. @@ -814,8 +767,7 @@ FILE_VERSION_FILTER = # output files in an output format independent way. To create the layout file # that represents doxygen's defaults, run doxygen with the -l option. You can # optionally specify a file name after the option, if omitted DoxygenLayout.xml -# will be used as the name of the layout file. See also section "Changing the -# layout of pages" for information. +# will be used as the name of the layout file. # # Note that if you run doxygen from a directory containing a file called # DoxygenLayout.xml, doxygen will parse it automatically even if the LAYOUT_FILE @@ -861,50 +813,27 @@ WARNINGS = YES WARN_IF_UNDOCUMENTED = YES # If the WARN_IF_DOC_ERROR tag is set to YES, doxygen will generate warnings for -# potential errors in the documentation, such as documenting some parameters in -# a documented function twice, or documenting parameters that don't exist or -# using markup commands wrongly. +# potential errors in the documentation, such as not documenting some parameters +# in a documented function, or documenting parameters that don't exist or using +# markup commands wrongly. # The default value is: YES. WARN_IF_DOC_ERROR = YES -# If WARN_IF_INCOMPLETE_DOC is set to YES, doxygen will warn about incomplete -# function parameter documentation. If set to NO, doxygen will accept that some -# parameters have no documentation without warning. -# The default value is: YES. - -WARN_IF_INCOMPLETE_DOC = YES - # This WARN_NO_PARAMDOC option can be enabled to get warnings for functions that # are documented, but have no documentation for their parameters or return -# value. If set to NO, doxygen will only warn about wrong parameter -# documentation, but not about the absence of documentation. If EXTRACT_ALL is -# set to YES then this flag will automatically be disabled. See also -# WARN_IF_INCOMPLETE_DOC +# value. If set to NO, doxygen will only warn about wrong or incomplete +# parameter documentation, but not about the absence of documentation. If +# EXTRACT_ALL is set to YES then this flag will automatically be disabled. # The default value is: NO. WARN_NO_PARAMDOC = NO -# If WARN_IF_UNDOC_ENUM_VAL option is set to YES, doxygen will warn about -# undocumented enumeration values. If set to NO, doxygen will accept -# undocumented enumeration values. If EXTRACT_ALL is set to YES then this flag -# will automatically be disabled. -# The default value is: NO. - -WARN_IF_UNDOC_ENUM_VAL = NO - # If the WARN_AS_ERROR tag is set to YES then doxygen will immediately stop when # a warning is encountered. If the WARN_AS_ERROR tag is set to FAIL_ON_WARNINGS # then doxygen will continue running as if WARN_AS_ERROR tag is set to NO, but # at the end of the doxygen process doxygen will return with a non-zero status. -# If the WARN_AS_ERROR tag is set to FAIL_ON_WARNINGS_PRINT then doxygen behaves -# like FAIL_ON_WARNINGS but in case no WARN_LOGFILE is defined doxygen will not -# write the warning messages in between other messages but write them at the end -# of a run, in case a WARN_LOGFILE is defined the warning messages will be -# besides being in the defined file also be shown at the end of a run, unless -# the WARN_LOGFILE is defined as - i.e. standard output (stdout) in that case -# the behavior will remain as with the setting FAIL_ON_WARNINGS. -# Possible values are: NO, YES, FAIL_ON_WARNINGS and FAIL_ON_WARNINGS_PRINT. +# Possible values are: NO, YES and FAIL_ON_WARNINGS. # The default value is: NO. WARN_AS_ERROR = NO @@ -915,27 +844,13 @@ WARN_AS_ERROR = NO # and the warning text. Optionally the format may contain $version, which will # be replaced by the version of the file (if it could be obtained via # FILE_VERSION_FILTER) -# See also: WARN_LINE_FORMAT # The default value is: $file:$line: $text. WARN_FORMAT = "$file:$line: $text" -# In the $text part of the WARN_FORMAT command it is possible that a reference -# to a more specific place is given. To make it easier to jump to this place -# (outside of doxygen) the user can define a custom "cut" / "paste" string. -# Example: -# WARN_LINE_FORMAT = "'vi $file +$line'" -# See also: WARN_FORMAT -# The default value is: at line $line of file $file. - -WARN_LINE_FORMAT = "at line $line of file $file" - # The WARN_LOGFILE tag can be used to specify a file to which warning and error # messages should be written. If left blank the output is written to standard -# error (stderr). In case the file specified cannot be opened for writing the -# warning and error messages are written to standard error. When as file - is -# specified the warning and error messages are written to standard output -# (stdout). +# error (stderr). WARN_LOGFILE = @@ -949,28 +864,18 @@ WARN_LOGFILE = # spaces. See also FILE_PATTERNS and EXTENSION_MAPPING # Note: If this tag is empty the current directory is searched. -INPUT = ../include +INPUT = ../include \ + pages # This tag can be used to specify the character encoding of the source files # that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses # libiconv (or the iconv built into libc) for the transcoding. See the libiconv # documentation (see: # https://www.gnu.org/software/libiconv/) for the list of possible encodings. -# See also: INPUT_FILE_ENCODING # The default value is: UTF-8. INPUT_ENCODING = UTF-8 -# This tag can be used to specify the character encoding of the source files -# that doxygen parses The INPUT_FILE_ENCODING tag can be used to specify -# character encoding on a per file pattern basis. Doxygen will compare the file -# name with each pattern and apply the encoding instead of the default -# INPUT_ENCODING) if there is a match. The character encodings are a list of the -# form: pattern=encoding (like *.php=ISO-8859-1). See cfg_input_encoding -# "INPUT_ENCODING" for further information on supported encodings. - -INPUT_FILE_ENCODING = - # If the value of the INPUT tag contains directories, you can use the # FILE_PATTERNS tag to specify one or more wildcard patterns (like *.cpp and # *.h) to filter out the source-files in the directories. @@ -982,22 +887,18 @@ INPUT_FILE_ENCODING = # Note the list of default checked file patterns might differ from the list of # default file extension mappings. # -# If left blank the following patterns are tested:*.c, *.cc, *.cxx, *.cxxm, -# *.cpp, *.cppm, *.ccm, *.c++, *.c++m, *.java, *.ii, *.ixx, *.ipp, *.i++, *.inl, -# *.idl, *.ddl, *.odl, *.h, *.hh, *.hxx, *.hpp, *.h++, *.ixx, *.l, *.cs, *.d, -# *.php, *.php4, *.php5, *.phtml, *.inc, *.m, *.markdown, *.md, *.mm, *.dox (to -# be provided as doxygen C comment), *.py, *.pyw, *.f90, *.f95, *.f03, *.f08, -# *.f18, *.f, *.for, *.vhd, *.vhdl, *.ucf, *.qsf and *.ice. +# If left blank the following patterns are tested:*.c, *.cc, *.cxx, *.cpp, +# *.c++, *.java, *.ii, *.ixx, *.ipp, *.i++, *.inl, *.idl, *.ddl, *.odl, *.h, +# *.hh, *.hxx, *.hpp, *.h++, *.cs, *.d, *.php, *.php4, *.php5, *.phtml, *.inc, +# *.m, *.markdown, *.md, *.mm, *.dox (to be provided as doxygen C comment), +# *.py, *.pyw, *.f90, *.f95, *.f03, *.f08, *.f18, *.f, *.for, *.vhd, *.vhdl, +# *.ucf, *.qsf and *.ice. FILE_PATTERNS = *.c \ *.cc \ *.cxx \ - *.cxxm \ *.cpp \ - *.cppm \ - *.ccm \ *.c++ \ - *.c++m \ *.java \ *.ii \ *.ixx \ @@ -1012,8 +913,6 @@ FILE_PATTERNS = *.c \ *.hxx \ *.hpp \ *.h++ \ - *.ixx \ - *.l \ *.cs \ *.d \ *.php \ @@ -1076,7 +975,10 @@ EXCLUDE_PATTERNS = # (namespaces, classes, functions, etc.) that should be excluded from the # output. The symbol name can be a fully qualified name, a word, or if the # wildcard * is used, a substring. Examples: ANamespace, AClass, -# ANamespace::AClass, ANamespace::*Test +# AClass::ANamespace, ANamespace::*Test +# +# Note that the wildcards are matched against the file with absolute path, so to +# exclude all test directories use the pattern */test/* EXCLUDE_SYMBOLS = @@ -1121,11 +1023,6 @@ IMAGE_PATH = # code is scanned, but not when the output code is generated. If lines are added # or removed, the anchors will not be placed correctly. # -# Note that doxygen will use the data processed and written to standard output -# for further processing, therefore nothing else, like debug statements or used -# commands (so in case of a Windows batch file always use @echo OFF), should be -# written to standard output. -# # Note that for custom extensions or not directly supported extensions you also # need to set EXTENSION_MAPPING for the extension otherwise the files are not # properly processed by doxygen. @@ -1167,15 +1064,6 @@ FILTER_SOURCE_PATTERNS = USE_MDFILE_AS_MAINPAGE = -# The Fortran standard specifies that for fixed formatted Fortran code all -# characters from position 72 are to be considered as comment. A common -# extension is to allow longer lines before the automatic comment starts. The -# setting FORTRAN_COMMENT_AFTER will also make it possible that longer lines can -# be processed before the automatic comment starts. -# Minimum value: 7, maximum value: 10000, default value: 72. - -FORTRAN_COMMENT_AFTER = 72 - #--------------------------------------------------------------------------- # Configuration options related to source browsing #--------------------------------------------------------------------------- @@ -1190,8 +1078,7 @@ FORTRAN_COMMENT_AFTER = 72 SOURCE_BROWSER = NO # Setting the INLINE_SOURCES tag to YES will include the body of functions, -# multi-line macros, enums or list initialized variables directly into the -# documentation. +# classes and enums directly into the documentation. # The default value is: NO. INLINE_SOURCES = NO @@ -1263,6 +1150,44 @@ USE_HTAGS = NO VERBATIM_HEADERS = YES +# If the CLANG_ASSISTED_PARSING tag is set to YES then doxygen will use the +# clang parser (see: +# http://clang.llvm.org/) for more accurate parsing at the cost of reduced +# performance. This can be particularly helpful with template rich C++ code for +# which doxygen's built-in parser lacks the necessary type information. +# Note: The availability of this option depends on whether or not doxygen was +# generated with the -Duse_libclang=ON option for CMake. +# The default value is: NO. + +CLANG_ASSISTED_PARSING = NO + +# If clang assisted parsing is enabled and the CLANG_ADD_INC_PATHS tag is set to +# YES then doxygen will add the directory of each input to the include path. +# The default value is: YES. + +CLANG_ADD_INC_PATHS = YES + +# If clang assisted parsing is enabled you can provide the compiler with command +# line options that you would normally use when invoking the compiler. Note that +# the include paths will already be set by doxygen for the files and directories +# specified with INPUT and INCLUDE_PATH. +# This tag requires that the tag CLANG_ASSISTED_PARSING is set to YES. + +CLANG_OPTIONS = + +# If clang assisted parsing is enabled you can provide the clang parser with the +# path to the directory containing a file called compile_commands.json. This +# file is the compilation database (see: +# http://clang.llvm.org/docs/HowToSetupToolingForLLVM.html) containing the +# options used when the source files were built. This is equivalent to +# specifying the -p option to a clang tool, such as clang-check. These options +# will then be passed to the parser. Any options specified with CLANG_OPTIONS +# will be added as well. +# Note: The availability of this option depends on whether or not doxygen was +# generated with the -Duse_libclang=ON option for CMake. + +CLANG_DATABASE_PATH = + #--------------------------------------------------------------------------- # Configuration options related to the alphabetical class index #--------------------------------------------------------------------------- @@ -1274,11 +1199,10 @@ VERBATIM_HEADERS = YES ALPHABETICAL_INDEX = YES -# The IGNORE_PREFIX tag can be used to specify a prefix (or a list of prefixes) -# that should be ignored while generating the index headers. The IGNORE_PREFIX -# tag works for classes, function and member names. The entity will be placed in -# the alphabetical list under the first letter of the entity name that remains -# after removing the prefix. +# In case all classes in a project start with a common prefix, all classes will +# be put under the same header in the alphabetical index. The IGNORE_PREFIX tag +# can be used to specify a prefix (or a list of prefixes) that should be ignored +# while generating the index headers. # This tag requires that the tag ALPHABETICAL_INDEX is set to YES. IGNORE_PREFIX = @@ -1290,7 +1214,7 @@ IGNORE_PREFIX = # If the GENERATE_HTML tag is set to YES, doxygen will generate HTML output # The default value is: YES. -GENERATE_HTML = NO +GENERATE_HTML = YES # The HTML_OUTPUT tag is used to specify where the HTML docs will be put. If a # relative path is entered the value of OUTPUT_DIRECTORY will be put in front of @@ -1357,15 +1281,11 @@ HTML_STYLESHEET = # Doxygen will copy the style sheet files to the output directory. # Note: The order of the extra style sheet files is of importance (e.g. the last # style sheet in the list overrules the setting of the previous ones in the -# list). -# Note: Since the styling of scrollbars can currently not be overruled in -# Webkit/Chromium, the styling will be left out of the default doxygen.css if -# one or more extra stylesheets have been specified. So if scrollbar -# customization is desired it has to be added explicitly. For an example see the -# documentation. +# list). For an example see the documentation. # This tag requires that the tag GENERATE_HTML is set to YES. -HTML_EXTRA_STYLESHEET = +HTML_EXTRA_STYLESHEET = doxygen-awesome-css/doxygen-awesome.css \ + doxygen-awesome-css/doxygen-awesome-sidebar-only.css # The HTML_EXTRA_FILES tag can be used to specify one or more extra images or # other source files which should be copied to the HTML output directory. Note @@ -1377,22 +1297,9 @@ HTML_EXTRA_STYLESHEET = HTML_EXTRA_FILES = -# The HTML_COLORSTYLE tag can be used to specify if the generated HTML output -# should be rendered with a dark or light theme. -# Possible values are: LIGHT always generate light mode output, DARK always -# generate dark mode output, AUTO_LIGHT automatically set the mode according to -# the user preference, use light mode if no preference is set (the default), -# AUTO_DARK automatically set the mode according to the user preference, use -# dark mode if no preference is set and TOGGLE allow to user to switch between -# light and dark mode via a button. -# The default value is: AUTO_LIGHT. -# This tag requires that the tag GENERATE_HTML is set to YES. - -HTML_COLORSTYLE = AUTO_LIGHT - # The HTML_COLORSTYLE_HUE tag controls the color of the HTML output. Doxygen # will adjust the colors in the style sheet and background images according to -# this color. Hue is specified as an angle on a color-wheel, see +# this color. Hue is specified as an angle on a colorwheel, see # https://en.wikipedia.org/wiki/Hue for more information. For instance the value # 0 represents red, 60 is yellow, 120 is green, 180 is cyan, 240 is blue, 300 # purple, and 360 is red again. @@ -1402,7 +1309,7 @@ HTML_COLORSTYLE = AUTO_LIGHT HTML_COLORSTYLE_HUE = 220 # The HTML_COLORSTYLE_SAT tag controls the purity (or saturation) of the colors -# in the HTML output. For a value of 0 the output will use gray-scales only. A +# in the HTML output. For a value of 0 the output will use grayscales only. A # value of 255 will produce the most vivid colors. # Minimum value: 0, maximum value: 255, default value: 100. # This tag requires that the tag GENERATE_HTML is set to YES. @@ -1420,6 +1327,15 @@ HTML_COLORSTYLE_SAT = 100 HTML_COLORSTYLE_GAMMA = 80 +# If the HTML_TIMESTAMP tag is set to YES then the footer of each generated HTML +# page will contain the date and time when the page was generated. Setting this +# to YES can help to show when doxygen was last run and thus if the +# documentation is up to date. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_TIMESTAMP = NO + # If the HTML_DYNAMIC_MENUS tag is set to YES then the generated HTML # documentation will contain a main index with vertical navigation menus that # are dynamically created via JavaScript. If disabled, the navigation index will @@ -1439,33 +1355,6 @@ HTML_DYNAMIC_MENUS = YES HTML_DYNAMIC_SECTIONS = NO -# If the HTML_CODE_FOLDING tag is set to YES then classes and functions can be -# dynamically folded and expanded in the generated HTML source code. -# The default value is: YES. -# This tag requires that the tag GENERATE_HTML is set to YES. - -HTML_CODE_FOLDING = YES - -# If the HTML_COPY_CLIPBOARD tag is set to YES then doxygen will show an icon in -# the top right corner of code and text fragments that allows the user to copy -# its content to the clipboard. Note this only works if supported by the browser -# and the web page is served via a secure context (see: -# https://www.w3.org/TR/secure-contexts/), i.e. using the https: or file: -# protocol. -# The default value is: YES. -# This tag requires that the tag GENERATE_HTML is set to YES. - -HTML_COPY_CLIPBOARD = YES - -# Doxygen stores a couple of settings persistently in the browser (via e.g. -# cookies). By default these settings apply to all HTML pages generated by -# doxygen across all projects. The HTML_PROJECT_COOKIE tag can be used to store -# the settings under a project specific key, such that the user preferences will -# be stored separately. -# This tag requires that the tag GENERATE_HTML is set to YES. - -HTML_PROJECT_COOKIE = - # With HTML_INDEX_NUM_ENTRIES one can control the preferred number of entries # shown in the various tree structured indices initially; the user can expand # and collapse entries dynamically later on. Doxygen will expand the tree to @@ -1502,13 +1391,6 @@ GENERATE_DOCSET = NO DOCSET_FEEDNAME = "Doxygen generated docs" -# This tag determines the URL of the docset feed. A documentation feed provides -# an umbrella under which multiple documentation sets from a single provider -# (such as a company or product suite) can be grouped. -# This tag requires that the tag GENERATE_DOCSET is set to YES. - -DOCSET_FEEDURL = - # This tag specifies a string that should uniquely identify the documentation # set bundle. This should be a reverse domain-name style string, e.g. # com.mycompany.MyDocSet. Doxygen will append .docset to the name. @@ -1534,12 +1416,8 @@ DOCSET_PUBLISHER_NAME = Publisher # If the GENERATE_HTMLHELP tag is set to YES then doxygen generates three # additional HTML index files: index.hhp, index.hhc, and index.hhk. The # index.hhp is a project file that can be read by Microsoft's HTML Help Workshop -# on Windows. In the beginning of 2021 Microsoft took the original page, with -# a.o. the download links, offline the HTML help workshop was already many years -# in maintenance mode). You can download the HTML help workshop from the web -# archives at Installation executable (see: -# http://web.archive.org/web/20160201063255/http://download.microsoft.com/downlo -# ad/0/A/9/0A939EF6-E31C-430F-A3DF-DFAE7960D564/htmlhelp.exe). +# (see: +# https://www.microsoft.com/en-us/download/details.aspx?id=21138) on Windows. # # The HTML Help Workshop contains a compiler that can convert all HTML output # generated by doxygen into a single compiled HTML file (.chm). Compiled HTML @@ -1596,16 +1474,6 @@ BINARY_TOC = NO TOC_EXPAND = NO -# The SITEMAP_URL tag is used to specify the full URL of the place where the -# generated documentation will be placed on the server by the user during the -# deployment of the documentation. The generated sitemap is called sitemap.xml -# and placed on the directory specified by HTML_OUTPUT. In case no SITEMAP_URL -# is specified no sitemap is generated. For information about the sitemap -# protocol see https://www.sitemaps.org -# This tag requires that the tag GENERATE_HTML is set to YES. - -SITEMAP_URL = - # If the GENERATE_QHP tag is set to YES and both QHP_NAMESPACE and # QHP_VIRTUAL_FOLDER are set, an additional index file will be generated that # can be used as input for Qt's qhelpgenerator to generate a Qt Compressed Help @@ -1699,7 +1567,7 @@ ECLIPSE_DOC_ID = org.doxygen.Project # The default value is: NO. # This tag requires that the tag GENERATE_HTML is set to YES. -DISABLE_INDEX = NO +DISABLE_INDEX = YES # The GENERATE_TREEVIEW tag is used to specify whether a tree-like index # structure should be generated to display hierarchical information. If the tag @@ -1708,27 +1576,15 @@ DISABLE_INDEX = NO # to work a browser that supports JavaScript, DHTML, CSS and frames is required # (i.e. any modern browser). Windows users are probably better off using the # HTML help feature. Via custom style sheets (see HTML_EXTRA_STYLESHEET) one can -# further fine tune the look of the index (see "Fine-tuning the output"). As an -# example, the default style sheet generated by doxygen has an example that -# shows how to put an image at the root of the tree instead of the PROJECT_NAME. -# Since the tree basically has the same information as the tab index, you could -# consider setting DISABLE_INDEX to YES when enabling this option. -# The default value is: NO. -# This tag requires that the tag GENERATE_HTML is set to YES. - -GENERATE_TREEVIEW = NO - -# When both GENERATE_TREEVIEW and DISABLE_INDEX are set to YES, then the -# FULL_SIDEBAR option determines if the side bar is limited to only the treeview -# area (value NO) or if it should extend to the full height of the window (value -# YES). Setting this to YES gives a layout similar to -# https://docs.readthedocs.io with more room for contents, but less room for the -# project logo, title, and description. If either GENERATE_TREEVIEW or -# DISABLE_INDEX is set to NO, this option has no effect. +# further fine-tune the look of the index. As an example, the default style +# sheet generated by doxygen has an example that shows how to put an image at +# the root of the tree instead of the PROJECT_NAME. Since the tree basically has +# the same information as the tab index, you could consider setting +# DISABLE_INDEX to YES when enabling this option. # The default value is: NO. # This tag requires that the tag GENERATE_HTML is set to YES. -FULL_SIDEBAR = NO +GENERATE_TREEVIEW = YES # The ENUM_VALUES_PER_LINE tag can be used to set the number of enum values that # doxygen will group on one line in the generated HTML documentation. @@ -1754,13 +1610,6 @@ TREEVIEW_WIDTH = 250 EXT_LINKS_IN_WINDOW = NO -# If the OBFUSCATE_EMAILS tag is set to YES, doxygen will obfuscate email -# addresses. -# The default value is: YES. -# This tag requires that the tag GENERATE_HTML is set to YES. - -OBFUSCATE_EMAILS = YES - # If the HTML_FORMULA_FORMAT option is set to svg, doxygen will use the pdf2svg # tool (see https://github.com/dawbarton/pdf2svg) or inkscape (see # https://inkscape.org) to generate formulas as SVG images instead of PNGs for @@ -1781,6 +1630,17 @@ HTML_FORMULA_FORMAT = png FORMULA_FONTSIZE = 10 +# Use the FORMULA_TRANSPARENT tag to determine whether or not the images +# generated for formulas are transparent PNGs. Transparent PNGs are not +# supported properly for IE 6.0, but are supported on all modern browsers. +# +# Note that when changing this option you need to delete any form_*.png files in +# the HTML output directory before the changes have effect. +# The default value is: YES. +# This tag requires that the tag GENERATE_HTML is set to YES. + +FORMULA_TRANSPARENT = YES + # The FORMULA_MACROFILE can contain LaTeX \newcommand and \renewcommand commands # to create new LaTeX commands to be used in formulas as building blocks. See # the section "Including formulas" for details. @@ -1798,29 +1658,11 @@ FORMULA_MACROFILE = USE_MATHJAX = NO -# With MATHJAX_VERSION it is possible to specify the MathJax version to be used. -# Note that the different versions of MathJax have different requirements with -# regards to the different settings, so it is possible that also other MathJax -# settings have to be changed when switching between the different MathJax -# versions. -# Possible values are: MathJax_2 and MathJax_3. -# The default value is: MathJax_2. -# This tag requires that the tag USE_MATHJAX is set to YES. - -MATHJAX_VERSION = MathJax_2 - # When MathJax is enabled you can set the default output format to be used for -# the MathJax output. For more details about the output format see MathJax -# version 2 (see: -# http://docs.mathjax.org/en/v2.7-latest/output.html) and MathJax version 3 -# (see: -# http://docs.mathjax.org/en/latest/web/components/output.html). +# the MathJax output. See the MathJax site (see: +# http://docs.mathjax.org/en/v2.7-latest/output.html) for more details. # Possible values are: HTML-CSS (which is slower, but has the best -# compatibility. This is the name for Mathjax version 2, for MathJax version 3 -# this will be translated into chtml), NativeMML (i.e. MathML. Only supported -# for NathJax 2. For MathJax version 3 chtml will be used instead.), chtml (This -# is the name for Mathjax version 3, for MathJax version 2 this will be -# translated into HTML-CSS) and SVG. +# compatibility), NativeMML (i.e. MathML) and SVG. # The default value is: HTML-CSS. # This tag requires that the tag USE_MATHJAX is set to YES. @@ -1833,21 +1675,15 @@ MATHJAX_FORMAT = HTML-CSS # MATHJAX_RELPATH should be ../mathjax. The default value points to the MathJax # Content Delivery Network so you can quickly see the result without installing # MathJax. However, it is strongly recommended to install a local copy of -# MathJax from https://www.mathjax.org before deployment. The default value is: -# - in case of MathJax version 2: https://cdn.jsdelivr.net/npm/mathjax@2 -# - in case of MathJax version 3: https://cdn.jsdelivr.net/npm/mathjax@3 +# MathJax from https://www.mathjax.org before deployment. +# The default value is: https://cdn.jsdelivr.net/npm/mathjax@2. # This tag requires that the tag USE_MATHJAX is set to YES. -MATHJAX_RELPATH = +MATHJAX_RELPATH = https://cdn.jsdelivr.net/npm/mathjax@2 # The MATHJAX_EXTENSIONS tag can be used to specify one or more MathJax # extension names that should be enabled during MathJax rendering. For example -# for MathJax version 2 (see -# https://docs.mathjax.org/en/v2.7-latest/tex.html#tex-and-latex-extensions): # MATHJAX_EXTENSIONS = TeX/AMSmath TeX/AMSsymbols -# For example for MathJax version 3 (see -# http://docs.mathjax.org/en/latest/input/tex/extensions/index.html): -# MATHJAX_EXTENSIONS = ams # This tag requires that the tag USE_MATHJAX is set to YES. MATHJAX_EXTENSIONS = @@ -2027,31 +1863,29 @@ PAPER_TYPE = a4 EXTRA_PACKAGES = -# The LATEX_HEADER tag can be used to specify a user-defined LaTeX header for -# the generated LaTeX document. The header should contain everything until the -# first chapter. If it is left blank doxygen will generate a standard header. It -# is highly recommended to start with a default header using -# doxygen -w latex new_header.tex new_footer.tex new_stylesheet.sty -# and then modify the file new_header.tex. See also section "Doxygen usage" for -# information on how to generate the default header that doxygen normally uses. +# The LATEX_HEADER tag can be used to specify a personal LaTeX header for the +# generated LaTeX document. The header should contain everything until the first +# chapter. If it is left blank doxygen will generate a standard header. See +# section "Doxygen usage" for information on how to let doxygen write the +# default header to a separate file. # -# Note: Only use a user-defined header if you know what you are doing! -# Note: The header is subject to change so you typically have to regenerate the -# default header when upgrading to a newer version of doxygen. The following -# commands have a special meaning inside the header (and footer): For a -# description of the possible markers and block names see the documentation. +# Note: Only use a user-defined header if you know what you are doing! The +# following commands have a special meaning inside the header: $title, +# $datetime, $date, $doxygenversion, $projectname, $projectnumber, +# $projectbrief, $projectlogo. Doxygen will replace $title with the empty +# string, for the replacement values of the other commands the user is referred +# to HTML_HEADER. # This tag requires that the tag GENERATE_LATEX is set to YES. LATEX_HEADER = -# The LATEX_FOOTER tag can be used to specify a user-defined LaTeX footer for -# the generated LaTeX document. The footer should contain everything after the -# last chapter. If it is left blank doxygen will generate a standard footer. See +# The LATEX_FOOTER tag can be used to specify a personal LaTeX footer for the +# generated LaTeX document. The footer should contain everything after the last +# chapter. If it is left blank doxygen will generate a standard footer. See # LATEX_HEADER for more information on how to generate a default footer and what -# special commands can be used inside the footer. See also section "Doxygen -# usage" for information on how to generate the default footer that doxygen -# normally uses. Note: Only use a user-defined footer if you know what you are -# doing! +# special commands can be used inside the footer. +# +# Note: Only use a user-defined footer if you know what you are doing! # This tag requires that the tag GENERATE_LATEX is set to YES. LATEX_FOOTER = @@ -2094,16 +1928,10 @@ PDF_HYPERLINKS = YES USE_PDFLATEX = YES -# The LATEX_BATCHMODE tag signals the behavior of LaTeX in case of an error. -# Possible values are: NO same as ERROR_STOP, YES same as BATCH, BATCH In batch -# mode nothing is printed on the terminal, errors are scrolled as if is -# hit at every error; missing files that TeX tries to input or request from -# keyboard input (\read on a not open input stream) cause the job to abort, -# NON_STOP In nonstop mode the diagnostic message will appear on the terminal, -# but there is no possibility of user interaction just like in batch mode, -# SCROLL In scroll mode, TeX will stop only for missing files to input or if -# keyboard input is necessary and ERROR_STOP In errorstop mode, TeX will stop at -# each error, asking for user intervention. +# If the LATEX_BATCHMODE tag is set to YES, doxygen will add the \batchmode +# command to the generated LaTeX files. This will instruct LaTeX to keep running +# if errors occur, instead of asking the user for help. This option is also used +# when generating formulas in HTML. # The default value is: NO. # This tag requires that the tag GENERATE_LATEX is set to YES. @@ -2116,6 +1944,16 @@ LATEX_BATCHMODE = NO LATEX_HIDE_INDICES = NO +# If the LATEX_SOURCE_CODE tag is set to YES then doxygen will include source +# code with syntax highlighting in the LaTeX output. +# +# Note that which sources are shown also depends on other settings such as +# SOURCE_BROWSER. +# The default value is: NO. +# This tag requires that the tag GENERATE_LATEX is set to YES. + +LATEX_SOURCE_CODE = NO + # The LATEX_BIB_STYLE tag can be used to specify the style to use for the # bibliography, e.g. plainnat, or ieeetr. See # https://en.wikipedia.org/wiki/BibTeX and \cite for more info. @@ -2124,6 +1962,14 @@ LATEX_HIDE_INDICES = NO LATEX_BIB_STYLE = plain +# If the LATEX_TIMESTAMP tag is set to YES then the footer of each generated +# page will contain the date and time when the page was generated. Setting this +# to NO can help when comparing the output of multiple runs. +# The default value is: NO. +# This tag requires that the tag GENERATE_LATEX is set to YES. + +LATEX_TIMESTAMP = NO + # The LATEX_EMOJI_DIRECTORY tag is used to specify the (relative or absolute) # path from which the emoji images will be read. If a relative path is entered, # it will be relative to the LATEX_OUTPUT directory. If left blank the @@ -2188,6 +2034,16 @@ RTF_STYLESHEET_FILE = RTF_EXTENSIONS_FILE = +# If the RTF_SOURCE_CODE tag is set to YES then doxygen will include source code +# with syntax highlighting in the RTF output. +# +# Note that which sources are shown also depends on other settings such as +# SOURCE_BROWSER. +# The default value is: NO. +# This tag requires that the tag GENERATE_RTF is set to YES. + +RTF_SOURCE_CODE = NO + #--------------------------------------------------------------------------- # Configuration options related to the man page output #--------------------------------------------------------------------------- @@ -2240,7 +2096,7 @@ MAN_LINKS = NO # captures the structure of the code including all documentation. # The default value is: NO. -GENERATE_XML = YES +GENERATE_XML = NO # The XML_OUTPUT tag is used to specify where the XML pages will be put. If a # relative path is entered the value of OUTPUT_DIRECTORY will be put in front of @@ -2284,44 +2140,27 @@ GENERATE_DOCBOOK = NO DOCBOOK_OUTPUT = docbook +# If the DOCBOOK_PROGRAMLISTING tag is set to YES, doxygen will include the +# program listings (including syntax highlighting and cross-referencing +# information) to the DOCBOOK output. Note that enabling this will significantly +# increase the size of the DOCBOOK output. +# The default value is: NO. +# This tag requires that the tag GENERATE_DOCBOOK is set to YES. + +DOCBOOK_PROGRAMLISTING = NO + #--------------------------------------------------------------------------- # Configuration options for the AutoGen Definitions output #--------------------------------------------------------------------------- # If the GENERATE_AUTOGEN_DEF tag is set to YES, doxygen will generate an -# AutoGen Definitions (see https://autogen.sourceforge.net/) file that captures +# AutoGen Definitions (see http://autogen.sourceforge.net/) file that captures # the structure of the code including all documentation. Note that this feature # is still experimental and incomplete at the moment. # The default value is: NO. GENERATE_AUTOGEN_DEF = NO -#--------------------------------------------------------------------------- -# Configuration options related to Sqlite3 output -#--------------------------------------------------------------------------- - -# If the GENERATE_SQLITE3 tag is set to YES doxygen will generate a Sqlite3 -# database with symbols found by doxygen stored in tables. -# The default value is: NO. - -GENERATE_SQLITE3 = NO - -# The SQLITE3_OUTPUT tag is used to specify where the Sqlite3 database will be -# put. If a relative path is entered the value of OUTPUT_DIRECTORY will be put -# in front of it. -# The default directory is: sqlite3. -# This tag requires that the tag GENERATE_SQLITE3 is set to YES. - -SQLITE3_OUTPUT = sqlite3 - -# The SQLITE3_RECREATE_DB tag is set to YES, the existing doxygen_sqlite3.db -# database file will be recreated with each doxygen run. If set to NO, doxygen -# will warn if a database file is already found and not modify it. -# The default value is: YES. -# This tag requires that the tag GENERATE_SQLITE3 is set to YES. - -SQLITE3_RECREATE_DB = YES - #--------------------------------------------------------------------------- # Configuration options related to the Perl module output #--------------------------------------------------------------------------- @@ -2396,8 +2235,7 @@ SEARCH_INCLUDES = YES # The INCLUDE_PATH tag can be used to specify one or more directories that # contain include files that are not input files but should be processed by the -# preprocessor. Note that the INCLUDE_PATH is not recursive, so the setting of -# RECURSIVE has no effect here. +# preprocessor. # This tag requires that the tag SEARCH_INCLUDES is set to YES. INCLUDE_PATH = @@ -2464,15 +2302,15 @@ TAGFILES = GENERATE_TAGFILE = -# If the ALLEXTERNALS tag is set to YES, all external classes and namespaces -# will be listed in the class and namespace index. If set to NO, only the -# inherited external classes will be listed. +# If the ALLEXTERNALS tag is set to YES, all external class will be listed in +# the class index. If set to NO, only the inherited external classes will be +# listed. # The default value is: NO. ALLEXTERNALS = NO # If the EXTERNAL_GROUPS tag is set to YES, all external groups will be listed -# in the topic index. If set to NO, only the current project's groups will be +# in the modules index. If set to NO, only the current project's groups will be # listed. # The default value is: YES. @@ -2486,9 +2324,25 @@ EXTERNAL_GROUPS = YES EXTERNAL_PAGES = YES #--------------------------------------------------------------------------- -# Configuration options related to diagram generator tools +# Configuration options related to the dot tool #--------------------------------------------------------------------------- +# If the CLASS_DIAGRAMS tag is set to YES, doxygen will generate a class diagram +# (in HTML and LaTeX) for classes with base or super classes. Setting the tag to +# NO turns the diagrams off. Note that this option also works with HAVE_DOT +# disabled, but it is recommended to install and use dot, since it yields more +# powerful graphs. +# The default value is: YES. + +CLASS_DIAGRAMS = YES + +# You can include diagrams made with dia in doxygen documentation. Doxygen will +# then run dia to produce the diagram and insert it in the documentation. The +# DIA_PATH tag allows you to specify the directory where the dia binary resides. +# If left empty dia is assumed to be found in the default search path. + +DIA_PATH = + # If set to YES the inheritance and collaboration graphs will hide inheritance # and usage relations if the target is undocumented or is not a class. # The default value is: YES. @@ -2497,12 +2351,12 @@ HIDE_UNDOC_RELATIONS = YES # If you set the HAVE_DOT tag to YES then doxygen will assume the dot tool is # available from the path. This tool is part of Graphviz (see: -# https://www.graphviz.org/), a graph visualization toolkit from AT&T and Lucent +# http://www.graphviz.org/), a graph visualization toolkit from AT&T and Lucent # Bell Labs. The other options in this section have no effect if this option is # set to NO -# The default value is: NO. +# The default value is: YES. -HAVE_DOT = NO +HAVE_DOT = YES # The DOT_NUM_THREADS specifies the number of dot invocations doxygen is allowed # to run in parallel. When set to 0 doxygen will base this on the number of @@ -2514,77 +2368,49 @@ HAVE_DOT = NO DOT_NUM_THREADS = 0 -# DOT_COMMON_ATTR is common attributes for nodes, edges and labels of -# subgraphs. When you want a differently looking font in the dot files that -# doxygen generates you can specify fontname, fontcolor and fontsize attributes. -# For details please see Node, -# Edge and Graph Attributes specification You need to make sure dot is able -# to find the font, which can be done by putting it in a standard location or by -# setting the DOTFONTPATH environment variable or by setting DOT_FONTPATH to the -# directory containing the font. Default graphviz fontsize is 14. -# The default value is: fontname=Helvetica,fontsize=10. +# When you want a differently looking font in the dot files that doxygen +# generates you can specify the font name using DOT_FONTNAME. You need to make +# sure dot is able to find the font, which can be done by putting it in a +# standard location or by setting the DOTFONTPATH environment variable or by +# setting DOT_FONTPATH to the directory containing the font. +# The default value is: Helvetica. # This tag requires that the tag HAVE_DOT is set to YES. -DOT_COMMON_ATTR = "fontname=Helvetica,fontsize=10" +DOT_FONTNAME = Helvetica -# DOT_EDGE_ATTR is concatenated with DOT_COMMON_ATTR. For elegant style you can -# add 'arrowhead=open, arrowtail=open, arrowsize=0.5'. Complete documentation about -# arrows shapes. -# The default value is: labelfontname=Helvetica,labelfontsize=10. +# The DOT_FONTSIZE tag can be used to set the size (in points) of the font of +# dot graphs. +# Minimum value: 4, maximum value: 24, default value: 10. # This tag requires that the tag HAVE_DOT is set to YES. -DOT_EDGE_ATTR = "labelfontname=Helvetica,labelfontsize=10" +DOT_FONTSIZE = 10 -# DOT_NODE_ATTR is concatenated with DOT_COMMON_ATTR. For view without boxes -# around nodes set 'shape=plain' or 'shape=plaintext' Shapes specification -# The default value is: shape=box,height=0.2,width=0.4. -# This tag requires that the tag HAVE_DOT is set to YES. - -DOT_NODE_ATTR = "shape=box,height=0.2,width=0.4" - -# You can set the path where dot can find font specified with fontname in -# DOT_COMMON_ATTR and others dot attributes. +# By default doxygen will tell dot to use the default font as specified with +# DOT_FONTNAME. If you specify a different font using DOT_FONTNAME you can set +# the path where dot can find it using this tag. # This tag requires that the tag HAVE_DOT is set to YES. DOT_FONTPATH = -# If the CLASS_GRAPH tag is set to YES or GRAPH or BUILTIN then doxygen will -# generate a graph for each documented class showing the direct and indirect -# inheritance relations. In case the CLASS_GRAPH tag is set to YES or GRAPH and -# HAVE_DOT is enabled as well, then dot will be used to draw the graph. In case -# the CLASS_GRAPH tag is set to YES and HAVE_DOT is disabled or if the -# CLASS_GRAPH tag is set to BUILTIN, then the built-in generator will be used. -# If the CLASS_GRAPH tag is set to TEXT the direct and indirect inheritance -# relations will be shown as texts / links. Explicit enabling an inheritance -# graph or choosing a different representation for an inheritance graph of a -# specific class, can be accomplished by means of the command \inheritancegraph. -# Disabling an inheritance graph can be accomplished by means of the command -# \hideinheritancegraph. -# Possible values are: NO, YES, TEXT, GRAPH and BUILTIN. +# If the CLASS_GRAPH tag is set to YES then doxygen will generate a graph for +# each documented class showing the direct and indirect inheritance relations. +# Setting this tag to YES will force the CLASS_DIAGRAMS tag to NO. # The default value is: YES. +# This tag requires that the tag HAVE_DOT is set to YES. CLASS_GRAPH = YES # If the COLLABORATION_GRAPH tag is set to YES then doxygen will generate a # graph for each documented class showing the direct and indirect implementation # dependencies (inheritance, containment, and class references variables) of the -# class with other documented classes. Explicit enabling a collaboration graph, -# when COLLABORATION_GRAPH is set to NO, can be accomplished by means of the -# command \collaborationgraph. Disabling a collaboration graph can be -# accomplished by means of the command \hidecollaborationgraph. +# class with other documented classes. # The default value is: YES. # This tag requires that the tag HAVE_DOT is set to YES. COLLABORATION_GRAPH = YES # If the GROUP_GRAPHS tag is set to YES then doxygen will generate a graph for -# groups, showing the direct groups dependencies. Explicit enabling a group -# dependency graph, when GROUP_GRAPHS is set to NO, can be accomplished by means -# of the command \groupgraph. Disabling a directory graph can be accomplished by -# means of the command \hidegroupgraph. See also the chapter Grouping in the -# manual. +# groups, showing the direct groups dependencies. # The default value is: YES. # This tag requires that the tag HAVE_DOT is set to YES. @@ -2626,8 +2452,8 @@ DOT_UML_DETAILS = NO # The DOT_WRAP_THRESHOLD tag can be used to set the maximum number of characters # to display on a single line. If the actual line length exceeds this threshold -# significantly it will be wrapped across multiple lines. Some heuristics are -# applied to avoid ugly line breaks. +# significantly it will wrapped across multiple lines. Some heuristics are apply +# to avoid ugly line breaks. # Minimum value: 0, maximum value: 1000, default value: 17. # This tag requires that the tag HAVE_DOT is set to YES. @@ -2644,9 +2470,7 @@ TEMPLATE_RELATIONS = NO # If the INCLUDE_GRAPH, ENABLE_PREPROCESSING and SEARCH_INCLUDES tags are set to # YES then doxygen will generate a graph for each documented file showing the # direct and indirect include dependencies of the file with other documented -# files. Explicit enabling an include graph, when INCLUDE_GRAPH is is set to NO, -# can be accomplished by means of the command \includegraph. Disabling an -# include graph can be accomplished by means of the command \hideincludegraph. +# files. # The default value is: YES. # This tag requires that the tag HAVE_DOT is set to YES. @@ -2655,10 +2479,7 @@ INCLUDE_GRAPH = YES # If the INCLUDED_BY_GRAPH, ENABLE_PREPROCESSING and SEARCH_INCLUDES tags are # set to YES then doxygen will generate a graph for each documented file showing # the direct and indirect include dependencies of the file with other documented -# files. Explicit enabling an included by graph, when INCLUDED_BY_GRAPH is set -# to NO, can be accomplished by means of the command \includedbygraph. Disabling -# an included by graph can be accomplished by means of the command -# \hideincludedbygraph. +# files. # The default value is: YES. # This tag requires that the tag HAVE_DOT is set to YES. @@ -2698,30 +2519,22 @@ GRAPHICAL_HIERARCHY = YES # If the DIRECTORY_GRAPH tag is set to YES then doxygen will show the # dependencies a directory has on other directories in a graphical way. The # dependency relations are determined by the #include relations between the -# files in the directories. Explicit enabling a directory graph, when -# DIRECTORY_GRAPH is set to NO, can be accomplished by means of the command -# \directorygraph. Disabling a directory graph can be accomplished by means of -# the command \hidedirectorygraph. +# files in the directories. # The default value is: YES. # This tag requires that the tag HAVE_DOT is set to YES. DIRECTORY_GRAPH = YES -# The DIR_GRAPH_MAX_DEPTH tag can be used to limit the maximum number of levels -# of child directories generated in directory dependency graphs by dot. -# Minimum value: 1, maximum value: 25, default value: 1. -# This tag requires that the tag DIRECTORY_GRAPH is set to YES. - -DIR_GRAPH_MAX_DEPTH = 1 - # The DOT_IMAGE_FORMAT tag can be used to set the image format of the images # generated by dot. For an explanation of the image formats see the section # output formats in the documentation of the dot tool (Graphviz (see: -# https://www.graphviz.org/)). +# http://www.graphviz.org/)). # Note: If you choose svg you need to set HTML_FILE_EXTENSION to xhtml in order # to make the SVG files visible in IE 9+ (other browsers do not have this # requirement). -# Possible values are: png, jpg, gif, svg, png:gd, png:gd:gd, png:cairo, +# Possible values are: png, png:cairo, png:cairo:cairo, png:cairo:gd, png:gd, +# png:gd:gd, jpg, jpg:cairo, jpg:cairo:gd, jpg:gd, jpg:gd:gd, gif, gif:cairo, +# gif:cairo:gd, gif:gd, gif:gd:gd, svg, png:gd, png:gd:gd, png:cairo, # png:cairo:gd, png:cairo:cairo, png:cairo:gdiplus, png:gdiplus and # png:gdiplus:gdiplus. # The default value is: png. @@ -2754,12 +2567,11 @@ DOT_PATH = DOTFILE_DIRS = -# You can include diagrams made with dia in doxygen documentation. Doxygen will -# then run dia to produce the diagram and insert it in the documentation. The -# DIA_PATH tag allows you to specify the directory where the dia binary resides. -# If left empty dia is assumed to be found in the default search path. +# The MSCFILE_DIRS tag can be used to specify one or more directories that +# contain msc files that are included in the documentation (see the \mscfile +# command). -DIA_PATH = +MSCFILE_DIRS = # The DIAFILE_DIRS tag can be used to specify one or more directories that # contain dia files that are included in the documentation (see the \diafile @@ -2768,10 +2580,10 @@ DIA_PATH = DIAFILE_DIRS = # When using plantuml, the PLANTUML_JAR_PATH tag should be used to specify the -# path where java can find the plantuml.jar file or to the filename of jar file -# to be used. If left blank, it is assumed PlantUML is not used or called during -# a preprocessing step. Doxygen will generate a warning when it encounters a -# \startuml command in this case and will not generate output for the diagram. +# path where java can find the plantuml.jar file. If left blank, it is assumed +# PlantUML is not used or called during a preprocessing step. Doxygen will +# generate a warning when it encounters a \startuml command in this case and +# will not generate output for the diagram. PLANTUML_JAR_PATH = @@ -2809,6 +2621,18 @@ DOT_GRAPH_MAX_NODES = 50 MAX_DOT_GRAPH_DEPTH = 0 +# Set the DOT_TRANSPARENT tag to YES to generate images with a transparent +# background. This is disabled by default, because dot on Windows does not seem +# to support this out of the box. +# +# Warning: Depending on the platform used, enabling this option may lead to +# badly anti-aliased labels on the edges of a graph (i.e. they become hard to +# read). +# The default value is: NO. +# This tag requires that the tag HAVE_DOT is set to YES. + +DOT_TRANSPARENT = NO + # Set the DOT_MULTI_TARGETS tag to YES to allow dot to generate multiple output # files in one run (i.e. multiple -o and -T options on the command line). This # makes dot run faster, but since only newer versions of dot (>1.8.10) support @@ -2821,8 +2645,6 @@ DOT_MULTI_TARGETS = NO # If the GENERATE_LEGEND tag is set to YES doxygen will generate a legend page # explaining the meaning of the various boxes and arrows in the dot generated # graphs. -# Note: This tag requires that UML_LOOK isn't set, i.e. the doxygen internal -# graphical representation for inheritance and collaboration diagrams is used. # The default value is: YES. # This tag requires that the tag HAVE_DOT is set to YES. @@ -2831,24 +2653,8 @@ GENERATE_LEGEND = YES # If the DOT_CLEANUP tag is set to YES, doxygen will remove the intermediate # files that are used to generate the various graphs. # -# Note: This setting is not only used for dot files but also for msc temporary -# files. +# Note: This setting is not only used for dot files but also for msc and +# plantuml temporary files. # The default value is: YES. DOT_CLEANUP = YES - -# You can define message sequence charts within doxygen comments using the \msc -# command. If the MSCGEN_TOOL tag is left empty (the default), then doxygen will -# use a built-in version of mscgen tool to produce the charts. Alternatively, -# the MSCGEN_TOOL tag can also specify the name an external tool. For instance, -# specifying prog as the value, doxygen will call the tool as prog -T -# -o . The external tool should support -# output file formats "png", "eps", "svg", and "ismap". - -MSCGEN_TOOL = - -# The MSCFILE_DIRS tag can be used to specify one or more directories that -# contain msc files that are included in the documentation (see the \mscfile -# command). - -MSCFILE_DIRS = diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt deleted file mode 100644 index 782e7a22..00000000 --- a/examples/CMakeLists.txt +++ /dev/null @@ -1,115 +0,0 @@ -add_executable(vector_add_example ${PROJECT_SOURCE_DIR}/examples/vector_add.cu) -if(KMM_USE_CUDA) - set_target_properties( - vector_add_example - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) -endif() -if(KMM_USE_HIP) - set_source_files_properties(${PROJECT_SOURCE_DIR}/examples/vector_add.cu PROPERTIES LANGUAGE HIP) -endif() -target_compile_features(vector_add_example PRIVATE cxx_std_17) -target_link_libraries(vector_add_example PRIVATE kmm) - -add_executable(cpp_threads_vector_add_example ${PROJECT_SOURCE_DIR}/examples/cpp_threads_vector_add.cu) -if(KMM_USE_CUDA) - set_target_properties( - cpp_threads_vector_add_example - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) -endif() -if(KMM_USE_HIP) - set_source_files_properties(${PROJECT_SOURCE_DIR}/examples/cpp_threads_vector_add.cu PROPERTIES LANGUAGE HIP) -endif() -target_compile_features(cpp_threads_vector_add_example PRIVATE cxx_std_17) -target_link_libraries(cpp_threads_vector_add_example PRIVATE kmm) - -add_executable(point_in_poly_example ${PROJECT_SOURCE_DIR}/examples/point_in_poly.cu) -if(KMM_USE_CUDA) - set_target_properties( - point_in_poly_example - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) -endif() -if(KMM_USE_HIP) - set_source_files_properties(${PROJECT_SOURCE_DIR}/examples/point_in_poly.cu PROPERTIES LANGUAGE HIP) -endif() -target_compile_features(point_in_poly_example PRIVATE cxx_std_17) -target_link_libraries(point_in_poly_example PRIVATE kmm) - - -add_executable(matrix_multiply_example ${PROJECT_SOURCE_DIR}/examples/matrix_multiply.cu) -if(KMM_USE_CUDA) - set_target_properties( - matrix_multiply_example - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) -endif() -if(KMM_USE_HIP) - set_source_files_properties(${PROJECT_SOURCE_DIR}/examples/matrix_multiply.cu PROPERTIES LANGUAGE HIP) -endif() -target_compile_features(matrix_multiply_example PRIVATE cxx_std_17) -target_link_libraries(matrix_multiply_example PRIVATE kmm) - - -add_executable(histogram_example ${PROJECT_SOURCE_DIR}/examples/histogram.cu) -if(KMM_USE_CUDA) - set_target_properties( - histogram_example - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) -endif() -if(KMM_USE_HIP) - set_source_files_properties(${PROJECT_SOURCE_DIR}/examples/histogram.cu PROPERTIES LANGUAGE HIP) -endif() -target_compile_features(histogram_example PRIVATE cxx_std_17) -target_link_libraries(histogram_example PRIVATE kmm) - - -add_executable(reduction_example ${PROJECT_SOURCE_DIR}/examples/reduction.cu) -if(KMM_USE_CUDA) - set_target_properties( - reduction_example - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) -endif() -if(KMM_USE_HIP) - set_source_files_properties(${PROJECT_SOURCE_DIR}/examples/reduction.cu PROPERTIES LANGUAGE HIP) -endif() -target_compile_features(reduction_example PRIVATE cxx_std_17) -target_link_libraries(reduction_example PRIVATE kmm) - - -add_executable(structs_example ${PROJECT_SOURCE_DIR}/examples/structs.cu) -if(KMM_USE_CUDA) - set_target_properties( - structs_example - PROPERTIES - CUDA_ARCHITECTURES "80" - CUDA_SEPARABLE_COMPILATION ON - CUDA_RESOLVE_DEVICE_SYMBOLS ON - ) -endif() -if(KMM_USE_HIP) - set_source_files_properties(${PROJECT_SOURCE_DIR}/examples/structs.cu PROPERTIES LANGUAGE HIP) -endif() -target_compile_features(structs_example PRIVATE cxx_std_17) -target_link_libraries(structs_example PRIVATE kmm) \ No newline at end of file diff --git a/examples/cpp_threads_vector_add.cu b/examples/cpp_threads_vector_add.cu deleted file mode 100644 index 4cca88d6..00000000 --- a/examples/cpp_threads_vector_add.cu +++ /dev/null @@ -1,107 +0,0 @@ -#include -#include - -#include "spdlog/spdlog.h" - -#include "kmm/kmm.hpp" - -__global__ void initialize_range(kmm::Range range, kmm::GPUSubviewMut output) { - int64_t i = blockIdx.x * blockDim.x + threadIdx.x + range.begin; - if (i >= range.end) { - return; - } - - output[i] = float(i); -} - -__global__ void fill_range( - kmm::Range range, - float value, - kmm::GPUSubviewMut output -) { - int64_t i = blockIdx.x * blockDim.x + threadIdx.x + range.begin; - if (i >= range.end) { - return; - } - - output[i] = value; -} - -__global__ void vector_add( - kmm::Range range, - kmm::GPUSubviewMut output, - kmm::GPUSubview left, - kmm::GPUSubview right -) { - int64_t i = blockIdx.x * blockDim.x + threadIdx.x + range.begin; - - if (i >= range.end) { - return; - } - - output[i] = left[i] + right[i]; -} - -void main_loop(unsigned int id, kmm::RuntimeHandle& rt, long n, long chunk_size, dim3 block_size) { - using namespace kmm::placeholders; - auto A = kmm::Array {n}; - auto B = kmm::Array {n}; - auto C = kmm::Array {n}; - auto domain = kmm::TileDomain(n, chunk_size); - - rt.parallel_submit( // - domain, - kmm::GPUKernel(initialize_range, block_size), - _x, - write(A[_x]) - ); - - rt.parallel_submit( - domain, - kmm::GPUKernel(fill_range, block_size), - _x, - float(1.0), - write(B[_x]) - ); - - rt.parallel_submit( - domain, - kmm::GPUKernel(vector_add, block_size), - _x, - write(C[_x]), - A[_x], - B[_x] - ); - - auto result = std::vector(n); - C.copy_to(result); - - // Correctness check - for (long i = 0; i < n; i++) { - if (result[i] != float(i) + 1.0F) { - std::cerr << "[THREAD " << id << "] - wrong result at " << i << " : " << result[i] - << " != " << float(i) + 1 << std::endl; - return; - } - } -} - -int main() { - auto rt = kmm::make_runtime(); - spdlog::set_level(spdlog::level::warn); - long n = 200'000'000; - long chunk_size = n / 10; - dim3 block_size = 256; - unsigned int num_threads = 16; - std::vector threads; - - for (unsigned int thread = 0; thread < num_threads; thread++) { - threads.emplace_back(main_loop, thread, std::ref(rt), n, chunk_size, block_size); - } - for (unsigned int thread = 0; thread < num_threads; thread++) { - threads.at(thread).join(); - } - - std::cout << "Correctness check completed." << std::endl; - return EXIT_SUCCESS; -} diff --git a/examples/histogram.cu b/examples/histogram.cu deleted file mode 100644 index baa3d3dc..00000000 --- a/examples/histogram.cu +++ /dev/null @@ -1,93 +0,0 @@ -#include - -#include "spdlog/spdlog.h" - -#include "kmm/kmm.hpp" - -void initialize_image( - unsigned long seed, - int width, - int height, - kmm::SubviewMut image -) { - std::mt19937 rand {seed}; - std::uniform_int_distribution dist {}; - - for (int i = 0; i < height; i++) { - for (int j = 0; j < width; j++) { - image[i][j] = dist(rand); - } - } -} - -void initialize_images( - kmm::Range subrange, - int width, - int height, - kmm::SubviewMut images -) { - for (auto i = subrange.begin; i < subrange.end; i++) { - initialize_image(i, width, height, images.drop_axis<0>(i)); - } -} - -__global__ void calculate_histogram( - kmm::Range image_ids, - int width, - int height, - kmm::GPUSubview images, - kmm::GPUSubviewMut histogram -) { - int image_id = blockIdx.z * blockDim.z + threadIdx.z + image_ids.begin; - int i = blockIdx.y * blockDim.y + threadIdx.y; - int j = blockIdx.x * blockDim.x + threadIdx.x; - - if (image_id < int(image_ids.end) && i < height && j < width) { - uint8_t value = images[image_id][i][j]; - atomicAdd(&histogram[image_id][value], 1); - } -} - -int main() { - using namespace kmm::placeholders; - spdlog::set_level(spdlog::level::trace); - - auto rt = kmm::make_runtime(); - int width = 1080; - int height = 1920; - int num_images = 2500; - int images_per_chunk = 500; - dim3 block_size = 256; - - auto histogram = kmm::Array {{256}}; - auto images = kmm::Array {{num_images, height, width}}; - - auto _imageid = kmm::Axis(2); - auto _i = kmm::Axis(1); - auto _j = kmm::Axis(0); - - rt.parallel_submit( - kmm::TileDomain(num_images, images_per_chunk), - kmm::Host(initialize_images), - _imageid, - width, - height, - write(images[_imageid][_][_]) - ); - - rt.synchronize(); - - rt.parallel_submit( - kmm::TileDomain({width, height, num_images}, {width, height, images_per_chunk}), - kmm::GPUKernel(calculate_histogram, block_size), - _imageid, - width, - height, - images[_imageid][_i][_j], - reduce(kmm::Reduction::Sum, privatize(_imageid), histogram[_]) - ); - - rt.synchronize(); - - return 0; -} diff --git a/examples/matrix_multiply.cu b/examples/matrix_multiply.cu deleted file mode 100644 index c74e1736..00000000 --- a/examples/matrix_multiply.cu +++ /dev/null @@ -1,136 +0,0 @@ -#include "spdlog/spdlog.h" - -#include "kmm/kmm.hpp" - -void fill_array(kmm::Bounds<2> region, kmm::SubviewMut array, float value) { - for (auto i = region.x.begin; i < region.x.end; i++) { - for (auto j = region.y.begin; j < region.y.end; j++) { - array[i][j] = value; - } - } -} - -void matrix_multiply( - kmm::DeviceResource& device, - kmm::Bounds<3> region, - int n, - int m, - int k, - kmm::GPUSubviewMut C, - kmm::GPUSubview A, - kmm::GPUSubview B -) { - using kmm::checked_cast; - - float alpha = 1.0; - float beta = 0.0; - - const float* A_ptr = A.data_at({region.y.begin, region.x.begin}); - const float* B_ptr = B.data_at({region.x.begin, region.z.begin}); - float* C_ptr = C.data_at({region.y.begin, region.z.begin}); - -#if __CUDA_ARCH__ - KMM_GPU_CHECK(cublasGemmEx( - device.blas(), - CUBLAS_OP_T, - CUBLAS_OP_T, - checked_cast(region.y.size()), - checked_cast(region.z.size()), - checked_cast(region.x.size()), - &alpha, - A_ptr, - CUDA_R_32F, - checked_cast(A.stride()), - B_ptr, - CUDA_R_32F, - checked_cast(B.stride()), - &beta, - C_ptr, - CUDA_R_32F, - checked_cast(C.stride()), - CUDA_R_32F, - CUBLAS_GEMM_DEFAULT - )); -#elif __HIP_DEVICE_COMPILE__ - KMM_GPU_CHECK(rocblas_gemm_ex( - device.blas(), - rocblas_operation_transpose, - rocblas_operation_transpose, - checked_cast(region.y.size()), - checked_cast(region.z.size()), - checked_cast(region.x.size()), - &alpha, - A_ptr, - rocblas_datatype_f32_r, - checked_cast(A.stride()), - B_ptr, - rocblas_datatype_f32_r, - checked_cast(B.stride()), - &beta, - C_ptr, - rocblas_datatype_f32_r, - checked_cast(C.stride()), - C_ptr, - rocblas_datatype_f32_r, - checked_cast(C.stride()), - rocblas_datatype_f32_r, - rocblas_gemm_algo_standard, - 0, - 0 - )); -#endif -} - -int main() { - using namespace kmm::placeholders; - spdlog::set_level(spdlog::level::trace); - - auto rt = kmm::make_runtime(); - int n = 50000; - int m = 50000; - int k = 50000; - int chunk_size = n / 5; - - auto A = kmm::Array {{n, k}}; - auto B = kmm::Array {{k, m}}; - auto C = kmm::Array {{n, m}}; - - rt.parallel_submit( - {n, k}, - {chunk_size, chunk_size}, - kmm::Host(fill_array), - bounds(_x, _y), - write(A[_x][_y]), - 1.0F - ); - - rt.parallel_submit( - {k, m}, - {chunk_size, chunk_size}, - kmm::Host(fill_array), - bounds(_x, _y), - write(B[_x][_y]), - 1.0F - ); - - for (size_t repeat = 0; repeat < 1; repeat++) { - C.reset(); - - rt.parallel_submit( - {k, n, m}, - {chunk_size, chunk_size, chunk_size}, - kmm::GPU(matrix_multiply), - bounds(_x, _y, _z), - n, - m, - k, - reduce(kmm::Reduction::Sum, C[_y][_z]), - A[_y][_x], - B[_x][_z] - ); - - rt.synchronize(); - } - - return EXIT_SUCCESS; -} diff --git a/examples/point_in_poly.cu b/examples/point_in_poly.cu deleted file mode 100644 index fc89065b..00000000 --- a/examples/point_in_poly.cu +++ /dev/null @@ -1,112 +0,0 @@ -#ifdef KMM_USE_CUDA - #include -#elif KMM_USE_HIP - #include -#endif - -#include "spdlog/spdlog.h" - -#include "kmm/api/launcher.hpp" -#include "kmm/api/mapper.hpp" -#include "kmm/api/runtime_handle.hpp" - -__global__ void cn_pnpoly( - kmm::Range chunk, - kmm::GPUSubviewMut bitmap, - kmm::GPUSubview points, - int nvertices, - kmm::GPUView vertices -) { - int i = blockIdx.x * blockDim.x + threadIdx.x + chunk.begin; - - if (i < chunk.end) { - int c = 0; - float2 p = points[i]; - - int k = nvertices - 1; - - for (int j = 0; j < nvertices; k = j++) { // edge from v to vp - float2 vj = vertices[j]; - float2 vk = vertices[k]; - - float slope = (vk.x - vj.x) / (vk.y - vj.y); - - if (((vj.y > p.y) != (vk.y > p.y)) && //if p is between vj and vk vertically - (p.x < slope * (p.y - vj.y) + vj.x - )) { //if p.x crosses the line vj-vk when moved in positive x-direction - c = !c; - } - } - - bitmap[i] = c; // 0 if even (out), and 1 if odd (in) - } -} - -__global__ void init_points(kmm::Range chunk, kmm::GPUSubviewMut points) { - int i = blockIdx.x * blockDim.x + threadIdx.x + chunk.begin; - - if (i < chunk.end) { -#if __CUDA_ARCH__ - curandStatePhilox4_32_10_t state; - curand_init(1234, i, 0, &state); - points[i] = {curand_normal(&state), curand_normal(&state)}; -#elif __HIP_DEVICE_COMPILE__ - rocrand_state_philox4x32_10 state; - rocrand_init(1234, i, 0, &state); - points[i] = {rocrand_normal(&state), rocrand_normal(&state)}; -#endif - } -} - -void init_polygon(kmm::Range chunk, int nvertices, kmm::ViewMut vertices) { - for (int64_t i = chunk.begin; i < chunk.end; i++) { - float angle = float(i) / float(nvertices) * float(2.0F * M_PI); - vertices[i] = {cosf(angle), sinf(angle)}; - } -} - -int main() { - using namespace kmm::placeholders; - spdlog::set_level(spdlog::level::trace); - - auto rt = kmm::make_runtime(); - int nvertices = 1000; - int npoints = 1'000'000'000; - int npoints_per_chunk = npoints / 10; - dim3 block_size = 256; - - auto vertices = kmm::Array {nvertices}; - auto points = kmm::Array {npoints}; - auto bitmap = kmm::Array {npoints}; - - rt.submit( - kmm::ResourceId::host(), - kmm::Host(init_polygon), - kmm::Range(nvertices), - nvertices, - write(vertices) - ); - - rt.parallel_submit( - {npoints}, - {npoints_per_chunk}, - kmm::GPUKernel(init_points, block_size), - _x, - write(points[_x]) - ); - - rt.parallel_submit( - {npoints}, - {npoints_per_chunk}, - kmm::GPUKernel(cn_pnpoly, block_size), - _x, - write(bitmap[_x]), - points[_x], - nvertices, - vertices - ); - - rt.synchronize(); - - return EXIT_SUCCESS; -} diff --git a/examples/reduction.cu b/examples/reduction.cu deleted file mode 100644 index 2c128ee6..00000000 --- a/examples/reduction.cu +++ /dev/null @@ -1,167 +0,0 @@ -#include - -#include "spdlog/spdlog.h" - -#include "kmm/kmm.hpp" - -__global__ void initialize_matrix_kernel( - kmm::Bounds<2, int> chunk, - kmm::GPUSubviewMut matrix -) { - int i = blockIdx.y * blockDim.y + threadIdx.y + chunk.y.begin; - int j = blockIdx.x * blockDim.x + threadIdx.x + chunk.x.begin; - - if (i < chunk.y.end && j < chunk.x.end) { - matrix[i][j] = float(i + 2 * j); - } -} - -__global__ void sum_total_kernel( - kmm::Bounds<2, int> chunk, - kmm::GPUSubview matrix, - kmm::GPUSubviewMut sum -) { - int i = blockIdx.y * blockDim.y + threadIdx.y + chunk.y.begin; - int j = blockIdx.x * blockDim.x + threadIdx.x + chunk.x.begin; - - if (i < chunk.y.end && j < chunk.x.end) { - sum[i][j] += matrix[i][j]; - } -} - -__global__ void sum_rows_kernel( - kmm::Bounds<2, int> chunk, - kmm::GPUSubview matrix, - kmm::GPUSubviewMut rows_sum -) { - int i = blockIdx.y * blockDim.y + threadIdx.y + chunk.y.begin; - int j = blockIdx.x * blockDim.x + threadIdx.x + chunk.x.begin; - - if (i < chunk.y.end && j < chunk.x.end) { - rows_sum[i][j] += matrix[i][j]; - } -} - -__global__ void sum_cols_kernel( - kmm::Bounds<2, int> chunk, - kmm::GPUSubview matrix, - kmm::GPUSubviewMut cols_sum -) { - int i = blockIdx.y * blockDim.y + threadIdx.y + chunk.y.begin; - int j = blockIdx.x * blockDim.x + threadIdx.x + chunk.x.begin; - - if (i < chunk.y.end && j < chunk.x.end) { - cols_sum[j][i] += matrix[i][j]; - } -} - -bool is_close(float expected, float gotten) { - return fabsf(expected - gotten) < fmaxf(1e-3F * fabsf(expected), 1e-9F); -} - -int run(kmm::RuntimeHandle& rt, int width, int height, int chunk_width, int chunk_height) { - using namespace kmm::placeholders; - auto domain = kmm::TileDomain({width, height}, {chunk_width, chunk_height}); - auto matrix = kmm::Array {{height, width}}; - - std::cout << "Execute for chunk size: " << chunk_width << "x" << chunk_height << "." - << std::endl; - - rt.parallel_submit( - domain, - kmm::GPUKernel(initialize_matrix_kernel, {16, 16}), - bounds(_x, _y), - write(matrix[_y][_x]) - ); - - rt.synchronize(); - - auto total_sum = kmm::Scalar(); - auto rows_sum = kmm::Array(height); - auto cols_sum = kmm::Array(width); - - rt.parallel_submit( - domain, - kmm::GPUKernel(sum_total_kernel, {16, 16}), - bounds(_x, _y), - matrix[_y][_x], - reduce(kmm::Reduction::Sum, privatize(_y, _x), total_sum) - ); - - rt.synchronize(); - - rt.parallel_submit( - domain, - kmm::GPUKernel(sum_rows_kernel, {16, 16}), - bounds(_x, _y), - matrix[_y][_x], - reduce(kmm::Reduction::Sum, privatize(_y), rows_sum[_x]) - ); - - rt.synchronize(); - - rt.parallel_submit( - domain, - kmm::GPUKernel(sum_cols_kernel, {16, 16}), - bounds(_x, _y), - matrix(_y, _x), - reduce(kmm::Reduction::Sum, privatize(_x), cols_sum[_y]) - ); - - rt.synchronize(); - - float total; - total_sum.copy_to(&total); - - if (!is_close(total, float(1.87125e+08))) { - std::cerr << "Wrong result for total_sum : " << total << " != " << float(1.87125e+08) - << std::endl; - return EXIT_FAILURE; - } - - std::vector rows; - rows_sum.copy_to(rows); - - for (int i = 0; i < height; i++) { - float expected = (float(width - 1) * 0.5F + float(2 * i)) * float(width); - - if (!is_close(rows[i], expected)) { - std::cerr << "Wrong result for rows_sum[" << i << "]: " << rows[i] << " != " << expected - << std::endl; - return EXIT_FAILURE; - } - } - - std::vector cols; - cols_sum.copy_to(cols); - - for (int i = 0; i < width; i++) { - float expected = (float(height - 1) + float(i)) * float(height); - - if (!is_close(cols[i], expected)) { - std::cerr << "Wrong result for cols_sum[" << i << "]: " << cols[i] << " != " << expected - << std::endl; - return EXIT_FAILURE; - } - } - - return EXIT_SUCCESS; -} - -int main() { - spdlog::set_level(spdlog::level::trace); - auto rt = kmm::make_runtime(); - int width = 500; - int height = 500; - - for (int nx = 1; nx <= 8; nx++) { - for (int ny = 1; ny <= 8; ny++) { - if (run(rt, width, height, width / nx, height / ny) != EXIT_SUCCESS) { - return EXIT_FAILURE; - } - } - } - - std::cout << "Correctness check completed." << std::endl; - return EXIT_SUCCESS; -} \ No newline at end of file diff --git a/examples/structs.cu b/examples/structs.cu deleted file mode 100644 index 9c0a4060..00000000 --- a/examples/structs.cu +++ /dev/null @@ -1,44 +0,0 @@ -#include - -#include "spdlog/spdlog.h" - -#include "kmm/kmm.hpp" - -/// This defines the struct for the host-side code -struct Example { - int x; - kmm::Array y; -}; - -/// This defines the struct for the device-side code -struct ExampleView { - int x; - kmm::View y; -}; - -// This defines the fields of the `Example` struct -KMM_DEFINE_STRUCT_ARGUMENT(Example, it.x, it.y) - -// This defines that the "view" of `Example` is `ExampleView` -KMM_DEFINE_STRUCT_VIEW(Example, ExampleView) - -void example(kmm::Range range, ExampleView input) { - KMM_ASSERT(input.x == 123); - KMM_ASSERT(input.y.size() == 3); - KMM_ASSERT(input.y[0] == 1.0F); - KMM_ASSERT(input.y[1] == 2.0F); - KMM_ASSERT(input.y[2] == 3.0F); - std::cout << "input is correct for range " << range << "!" << std::endl; -} - -int main() { - using namespace kmm::placeholders; - auto rt = kmm::make_runtime(); - auto y = rt.allocate({1.0F, 2.0F, 3.0F}); - auto structure = Example {.x = 123, .y = y}; - - rt.parallel_submit(kmm::TileDomain({1000}, {200}), kmm::Host(example), _x, structure); - rt.synchronize(); - - return EXIT_SUCCESS; -} diff --git a/examples/vector_add.cu b/examples/vector_add.cu deleted file mode 100644 index cc31ff19..00000000 --- a/examples/vector_add.cu +++ /dev/null @@ -1,96 +0,0 @@ -#include - -#include "spdlog/spdlog.h" - -#include "kmm/kmm.hpp" - -__global__ void initialize_range(kmm::Range range, kmm::GPUSubviewMut output) { - int64_t i = blockIdx.x * blockDim.x + threadIdx.x + range.begin; - if (i >= range.end) { - return; - } - - output[i] = float(i); -} - -__global__ void fill_range( - kmm::Range range, - float value, - kmm::GPUSubviewMut output -) { - int64_t i = blockIdx.x * blockDim.x + threadIdx.x + range.begin; - if (i >= range.end) { - return; - } - - output[i] = value; -} - -__global__ void vector_add( - kmm::Range range, - kmm::GPUSubviewMut output, - kmm::GPUSubview left, - kmm::GPUSubview right -) { - int64_t i = blockIdx.x * blockDim.x + threadIdx.x + range.begin; - - if (i >= range.end) { - return; - } - - output[i] = left[i] + right[i]; -} - -int main() { - using namespace kmm::placeholders; - - auto rt = kmm::make_runtime(); - spdlog::set_level(spdlog::level::trace); - long n = 200'000'000; - long chunk_size = n / 10; - dim3 block_size = 256; - - auto A = kmm::Array {n}; - auto B = kmm::Array {n}; - auto C = kmm::Array {n}; - auto domain = kmm::TileDomain(n, chunk_size); - - rt.parallel_submit( // - domain, - kmm::GPUKernel(initialize_range, block_size), - _x, - write(A[_x]) - ); - - rt.parallel_submit( - domain, - kmm::GPUKernel(fill_range, block_size), - _x, - float(1.0), - write(B[_x]) - ); - - rt.parallel_submit( - domain, - kmm::GPUKernel(vector_add, block_size), - _x, - write(C[_x]), - A[_x], - B[_x] - ); - - auto result = std::vector(n); - C.copy_to(result); - - // Correctness check - for (long i = 0; i < n; i++) { - if (result[i] != float(i) + 1.0F) { - std::cerr << "Wrong result at " << i << " : " << result[i] << " != " << float(i) + 1 - << std::endl; - return EXIT_FAILURE; - } - } - - std::cout << "Correctness check completed." << std::endl; - return EXIT_SUCCESS; -} diff --git a/external/unordered_dense b/external/unordered_dense new file mode 160000 index 00000000..e5b9441e --- /dev/null +++ b/external/unordered_dense @@ -0,0 +1 @@ +Subproject commit e5b9441ecf193f3e1e8f954527fc76edee20d7eb diff --git a/include/kmm/api/access.hpp b/include/kmm/api/access.hpp deleted file mode 100644 index aa1faeee..00000000 --- a/include/kmm/api/access.hpp +++ /dev/null @@ -1,130 +0,0 @@ -#pragma once - -#include "kmm/api/argument.hpp" -#include "kmm/api/mapper.hpp" - -namespace kmm { - -/** - * Encapsulates read access to an argument. - */ -template -struct Read { - Arg& argument; - M access_mapper = {}; -}; - -template -Read read(const Arg& argument, M access_mapper = {}) { - return {argument, access_mapper}; -} - -template -Read read(Read access) { - return access; -} - -/** - * Encapsulates write access to an argument. - */ -template -struct Write { - Arg& argument; - M access_mapper = {}; -}; - -template -Write write(Arg& argument, M access_mapper = {}) { - return {argument, access_mapper}; -} - -template -Write write(Read access) { - return {access.argument, access.access_mapper}; -} - -/** - * Encapsulates reduce access to an argument. - */ -template> -struct Reduce { - Arg& argument; - Reduction op; - M access_mapper = {}; - P private_mapper = {}; -}; - -template -struct Privatize { - M access_mapper; - - explicit Privatize(M access_mapper) : // - access_mapper(std::move(access_mapper)) {} -}; - -template -Privatize privatize(const M& mapper) { - return Privatize {mapper}; -} - -template -Privatize> privatize(const Is&... slices) { - return Privatize {bounds(slices...)}; -} - -template -Reduce reduce(Reduction op, Arg& argument, M access_mapper = {}) { - return {argument, op, access_mapper}; -} - -template -Reduce reduce( - Reduction op, - Privatize

private_mapper, - Arg& argument, - M access_mapper = {} -) { - return {argument, op, access_mapper, private_mapper.access_mapper}; -} - -template -Reduce reduce(Reduction op, Read access) { - return {access.argument, op, access.access_mapper}; -} - -template -Reduce reduce(Reduction op, Privatize

private_mapper, Read access) { - return {access.argument, op, access.access_mapper, private_mapper.access_mapper}; -} - -template typename Mode, size_t N, size_t I = 0> -struct MultiIndexAccess { - MultiIndexAccess(Arg& m_argument, MultiIndexMap m_mapper = {}) : - m_argument(m_argument), - m_mapper(m_mapper) {} - - template - auto operator[](const M& index) { - m_mapper.axes[I] = into_index_map(index); - - if constexpr (I + 1 == N) { - return Mode> {m_argument, {m_mapper}}; - } else { - return MultiIndexAccess(m_argument, m_mapper); - } - } - - private: - Arg& m_argument; - MultiIndexMap m_mapper = {}; -}; - -// Forward `Read` to `Read` only if `Arg` is not const -template -struct ArgumentHandler, std::enable_if_t>>: - ArgumentHandler> { - ArgumentHandler(Read access) : - ArgumentHandler>({access.argument, access.access_mapper}) {} -}; - -} // namespace kmm \ No newline at end of file diff --git a/include/kmm/api/argument.hpp b/include/kmm/api/argument.hpp deleted file mode 100644 index 18fd1d8a..00000000 --- a/include/kmm/api/argument.hpp +++ /dev/null @@ -1,92 +0,0 @@ -#pragma once - -#include "kmm/api/task_group.hpp" -#include "kmm/core/resource.hpp" - -namespace kmm { - -template -struct ArgumentHandler; - -template -struct ArgumentHandler: ArgumentHandler { - ArgumentHandler(const T& arg) : ArgumentHandler(arg) {} -}; - -template -struct ArgumentHandler: ArgumentHandler { - ArgumentHandler(T& arg) : ArgumentHandler(arg) {} -}; - -template -struct ArgumentHandler: ArgumentHandler { - ArgumentHandler(T&& arg) : ArgumentHandler(std::move(arg)) {} -}; - -template -using packed_argument_t = typename ArgumentHandler::type; - -template -packed_argument_t pack_argument(TaskInstance& task, T&& arg) { - return ArgumentHandler(std::forward(arg)).before_submit(task); -} - -template -struct ArgumentUnpack; - -template -auto unpack_argument(TaskContext& context, T&& arg) { - return ArgumentUnpack>::call(context, std::forward(arg)); -} - -template -struct Argument { - Argument(T value) : m_value(std::move(value)) {} - - static Argument pack(TaskInstance& builder, T value) { - return Argument {std::move(value)}; - } - - template - T unpack(TaskContext& context) { - return m_value; - } - - private: - T m_value; -}; - -template -struct ArgumentHandler { - using type = Argument; - - ArgumentHandler(T value) : m_value(std::move(value)) {} - - void initialize(const TaskGroupInit& init) { - // Nothing to do - } - - type before_submit(TaskInstance& builder) { - return Argument::pack(builder, m_value); - } - - void after_submit(const TaskSubmissionResult& result) { - // Nothing to do - } - - void commit(const TaskGroupCommit& commit) { - // Nothing to do - } - - private: - T m_value; -}; - -template -struct ArgumentUnpack> { - static auto call(TaskContext& context, Argument& data) { - return data.template unpack(context); - } -}; - -} // namespace kmm \ No newline at end of file diff --git a/include/kmm/api/array.hpp b/include/kmm/api/array.hpp index 4648ec70..4ac53a0a 100644 --- a/include/kmm/api/array.hpp +++ b/include/kmm/api/array.hpp @@ -1,495 +1,317 @@ #pragma once -#include -#include -#include +#include +#include +#include +#include +#include #include -#include "spdlog/spdlog.h" - -#include "kmm/api/access.hpp" -#include "kmm/api/argument.hpp" -#include "kmm/api/array_instance.hpp" -#include "kmm/api/view_argument.hpp" -#include "kmm/planner/read_planner.hpp" -#include "kmm/planner/reduction_planner.hpp" -#include "kmm/planner/write_planner.hpp" +#include "kmm/api/array_base.hpp" +#include "kmm/api/launch_arg.hpp" +#include "kmm/core/layout.hpp" +#include "kmm/core/view.hpp" +#include "kmm/runtime/buffer.hpp" +#include "kmm/runtime/device_event.hpp" +#include "kmm/runtime/identifiers.hpp" +#include "kmm/runtime/memops/fill.hpp" +#include "kmm/runtime/requisition.hpp" namespace kmm { -class ArrayBase { - public: - virtual ~ArrayBase() = default; - virtual const std::type_info& type_info() const = 0; - virtual size_t rank() const = 0; - virtual int64_t size(size_t axis) const = 0; - virtual const Runtime& runtime() const = 0; - virtual void synchronize() const = 0; - virtual void copy_bytes_to(void* output, size_t num_bytes) const = 0; -}; +template +class DomainArray: public ArrayBase { + // Grants every `DomainArray` instantiation access to every other's private constructor. + template + friend class DomainArray; + + // Grants `Reduce>::finalize()` access to the private raw + // constructor, so it can hand back a `DomainArray` without going through a public constructor + // that would (re-)trigger `Runtime::begin_reduction`. + friend class Reduce>; -template -class Array: public ArrayBase { public: - Array(Dim shape = {}) : m_shape(shape) {} + using self_type = DomainArray; + using element_type = T; + using domain_type = DomainT; + using policy_type = PolicyT; + using layout_type = Layout; + static constexpr size_t rank = layout_type::rank; + using mapping_type = typename layout_type::mapping_type; + using index_type = typename layout_type::index_type; + using ndindex_type = typename layout_type::ndindex_type; + using shape_type = typename layout_type::shape_type; + using range_type = typename layout_type::range_type; + using bounds_type = typename layout_type::bounds_type; + using stride_type = typename layout_type::stride_type; + using ndstrides_type = typename layout_type::ndstrides_type; - explicit Array(std::shared_ptr> b) : - m_instance(b), - m_shape(m_instance->distribution().array_size()) {} + /// The `DomainArray` sharing this array's element type but over a different `Layout`, + template + using rebind_layout = DomainArray< // + T, + typename NewLayoutT::domain_type, + typename NewLayoutT::policy_type>; - const std::type_info& type_info() const final { - return typeid(T); - } + using zero_origin_type = rebind_layout; + using move_origin_type = rebind_layout; + using reverse_axes_type = rebind_layout; - size_t rank() const final { - return N; - } + template + using drop_axis_type = rebind_layout>; - Dim shape() const { - return m_shape; - } + template + using insert_axis_type = rebind_layout>; - int64_t size(size_t axis) const final { - return m_shape.get_or_default(axis); - } + template + using slice_axis_type = + rebind_layout>; - int64_t size() const { - return m_shape.volume(); - } + template + using slice_type = rebind_layout>; - bool is_empty() const { - return m_shape.is_empty(); - } + DomainArray() = default; - bool has_instance() const { - return m_instance != nullptr; - } + DomainArray(layout_type layout) : m_layout(layout) {} - ArrayInstance& instance() const { - if (m_instance == nullptr) { - throw_uninitialized_array_exception(); - } + DomainArray(domain_type domain, policy_type policy = {}) : + DomainArray(make_layout(domain, policy).normalize_offset()) {} - return *m_instance; - } + DomainArray(Runtime runtime, layout_type layout, std::optional fill_value = std::nullopt) : + ArrayBase(Buffer( + runtime, + BufferLayout::for_type(static_cast(layout.offset_span().size())), + "array", + fill_value ? FillValue::from(*fill_value) : FillValue {} + )), + m_layout(layout.normalize_offset()) {} - const Distribution& distribution() const { - return instance().distribution(); - } + DomainArray( + Runtime runtime, + domain_type domain, + policy_type policy = {}, + std::optional fill_value = std::nullopt + ) : + DomainArray(runtime, make_layout(domain, policy), fill_value) {} - Dim chunk_size() const { - return distribution().chunk_size(); - } + template + DomainArray(const DomainArray& that) : + ArrayBase(that), + m_layout(that.layout()) {} - int64_t chunk_size(size_t axis) const { - return chunk_size().get_or_default(axis); + const layout_type& layout() const noexcept { + return m_layout; } - Runtime& runtime() const final { - return instance().runtime(); + const domain_type& domain() const noexcept { + return m_layout.domain(); } - void synchronize() const final { - if (m_instance) { - m_instance->synchronize(); - } + const mapping_type& mapping() const noexcept { + return m_layout.mapping(); } - void reset() { - m_instance = nullptr; + shape_type shape() const noexcept { + return m_layout.shape(); } - template - Read, M> access(M mapper = {}) { - return {*this, {std::move(mapper)}}; + /// The valid index range along the given axis. + range_type bounds(size_t axis) const noexcept { + return m_layout.bounds(axis); } - template - Read, M> access(M mapper = {}) const { - return {*this, {std::move(mapper)}}; + /// The bounds (begin/end per axis) covered by this array. + bounds_type bounds() const noexcept { + return m_layout.bounds(); } - template - auto operator[](M first_index) { - return MultiIndexAccess, Read, N>(*this)[first_index]; + /// The first valid index along the given axis. + index_type origin(size_t axis) const noexcept { + return m_layout.origin(axis); } - template - auto operator[](M first_index) const { - return MultiIndexAccess, Read, N>(*this)[first_index]; + /// The first valid index along each axis. + ndindex_type origin() const noexcept { + return m_layout.origin(); } - template - Read, MultiIndexMap> operator()(const Is&... index) { - return access(bounds(index...)); + stride_type stride(size_t axis) const noexcept { + return m_layout.stride(axis); } - template - Read, MultiIndexMap> operator()(const Is&... index) const { - return access(bounds(index...)); + /// The stride along each axis. + ndstrides_type strides() const noexcept { + return m_layout.strides(); } - void copy_bytes_to(void* output, size_t num_bytes) const { - KMM_ASSERT(num_bytes % sizeof(T) == 0); - KMM_ASSERT(checked_equals(num_bytes / sizeof(T), size())); - instance().copy_bytes_into(output); + index_type extent(size_t axis) const noexcept { + return m_layout.extent(axis); } - void copy_to(T* output) const { - instance().copy_bytes_into(output); + /// The start of the valid range along the given axis. + index_type begin(size_t axis) const noexcept { + return m_layout.begin(axis); } - template - void copy_to(T* output, I num_elements) const { - KMM_ASSERT(checked_equals(num_elements, size())); - instance().copy_bytes_into(output); + /// The end (exclusive) of the valid range along the given axis. + index_type end(size_t axis) const noexcept { + return m_layout.end(axis); } - void copy_to(std::vector& output) const { - output.resize(checked_cast(size())); - instance().copy_bytes_into(output.data()); + /// The start of the valid range along each axis. + ndindex_type begin() const noexcept { + return m_layout.begin(); } - std::vector copy_to_vector() const { - std::vector output(size()); - copy_to(output); - return output; + /// The end (exclusive) of the valid range along each axis. + ndindex_type end() const noexcept { + return m_layout.end(); } - void copy_bytes_from(const void* input, size_t num_bytes) const { - KMM_ASSERT(num_bytes % sizeof(T) == 0); - KMM_ASSERT(checked_equals(num_bytes / sizeof(T), size())); - instance().copy_bytes_from(input); + /// The total number of elements in this array. + index_type size() const noexcept { + return m_layout.size(); } - void copy_from(T* input) const { - instance().copy_bytes_from(input); + /// Whether this array covers zero elements. + bool is_empty() const noexcept { + return m_layout.is_empty(); } - template - void copy_from(T* input, I num_elements) const { - KMM_ASSERT(checked_equals(num_elements, size())); - instance().copy_bytes_from(input); + /// Returns this array rebased so its domain starts at the zero index. + zero_origin_type zero_origin() const noexcept { + return {buffer(), m_layout.zero_origin()}; } - void copy_from(const std::vector& input) const { - copy_from(input.data(), input.size()); + /// Returns this array shifted so it originates at the given index, keeping the same shape. + move_origin_type move_origin(ndindex_type new_origin) const noexcept { + return {buffer(), m_layout.move_origin(new_origin)}; } - private: - std::shared_ptr> m_instance; - Index m_offset; // Unused for now, always zero - Dim m_shape; -}; - -template -using Scalar = Array; - -template -struct ArgumentHandler>> { - using type = ViewArgument>; - - ArgumentHandler(Read> access) : - m_planner(access.argument.instance().shared_from_this()), - m_array_shape(access.argument.shape()) {} - - void initialize(const TaskGroupInit& init) {} - - type before_submit(TaskInstance& task) { - auto region = Bounds(m_array_shape); - size_t buffer_index = task.add_buffer_requirement( // - m_planner.prepare_access(task.graph, task.memory_id, region, task.dependencies) - ); - - auto domain = views::dynamic_domain {region.sizes()}; - return {buffer_index, domain}; - } - - void after_submit(const TaskSubmissionResult& result) { - m_planner.finalize_access(result.graph, result.event_id); + /// Returns this array restricted to the intersection of its bounds and the given bounds. + move_origin_type restrict_bounds(bounds_type new_bounds) const noexcept { + return {buffer(), m_layout.restrict_bounds(new_bounds)}; } - void commit(const TaskGroupCommit& commit) { - m_planner.commit(commit.graph); + /// Returns this array restricted along one axis to the intersection with [start, stop). + template + move_origin_type restrict_axis(index_type start, index_type stop) const noexcept { + return {buffer(), m_layout.template restrict_axis(start, stop)}; } - private: - ArrayReadPlanner m_planner; - Dim m_array_shape; -}; - -template -struct ArgumentHandler>: ArgumentHandler>> { - ArgumentHandler(Array array) : ArgumentHandler>>(read(array)) {} -}; - -template -struct ArgumentHandler, M>> { - using type = ViewArgument>; - - static_assert( - is_dimensionality_accepted_by_mapper, - "mapper of 'read' must return N-dimensional region" - ); - - ArgumentHandler(Read, M> access) : - m_planner(access.argument.instance().shared_from_this()), - m_array_shape(access.argument.shape()), - m_access_mapper(access.access_mapper) {} - - void initialize(const TaskGroupInit& init) {} - - type before_submit(TaskInstance& task) { - Bounds region = m_access_mapper(task.chunk, Bounds(m_array_shape)); - auto buffer_index = task.add_buffer_requirement( // - m_planner.prepare_access(task.graph, task.memory_id, region, task.dependencies) - ); - - auto domain = views::dynamic_subdomain {region.begin(), region.sizes()}; - return {buffer_index, domain}; + /// Returns this array with the given axis dropped, fixed at the given index. + template + drop_axis_type drop_axis(index_type index) const noexcept { + return {buffer(), m_layout.template drop_axis(index)}; } - void after_submit(const TaskSubmissionResult& result) { - m_planner.finalize_access(result.graph, result.event_id); + /// Returns this array with a new broadcast axis of the given extent inserted at the given + /// position. + template + insert_axis_type insert_axis(index_type extent = static_cast(1)) + const noexcept { + return {buffer(), m_layout.template insert_axis(extent)}; } - void commit(const TaskGroupCommit& commit) { - m_planner.commit(commit.graph); + /// Returns this array with the order of all axes reversed. + reverse_axes_type reverse_axes() const noexcept { + return {buffer(), m_layout.reverse_axes()}; } - private: - ArrayReadPlanner m_planner; - Dim m_array_shape; - M m_access_mapper; -}; - -template -struct ArgumentHandler>> { - using type = ViewArgument>; - - ArgumentHandler(Write> access) : m_array(access.argument) {} - - void initialize(const TaskGroupInit& init) { - if (!m_array.has_instance()) { - auto instance = ArrayInstance::create( // - init.runtime, - map_domain_to_distribution(m_array.shape(), init.domain, All()), - DataType::of() - ); - - m_array = Array(instance); - } - - m_planner = std::make_unique>(m_array.instance().shared_from_this()); + /// Returns this array with the given axis sliced according to the given slice token (e.g. + /// `all`, a `Range`, `new_axis`). + template + slice_axis_type slice_axis(const SliceT& slice) const noexcept { + return {buffer(), m_layout.template slice_axis(slice)}; } - type before_submit(TaskInstance& task) { - auto access_region = Bounds(m_array.shape()); - auto buffer_index = task.add_buffer_requirement( - m_planner->prepare_access(task.graph, task.memory_id, access_region, task.dependencies) - ); - - auto domain = views::dynamic_domain {access_region.sizes()}; - return {buffer_index, domain}; + /// Returns this array with the given axis narrowed to the range [start, end). + template + self_type slice_axis(index_type start, index_type end) const noexcept { + return self_type(buffer(), m_layout.template slice_axis(start, end)); } - void after_submit(const TaskSubmissionResult& result) { - m_planner->finalize_access(result.graph, result.event_id); + /// Returns this array sliced across all axes at once, one slice token per axis. + template + slice_type slice(const Slices&... slices) const noexcept { + return {buffer(), m_layout.slice(slices...)}; } - void commit(const TaskGroupCommit& commit) { - m_planner->commit(commit.graph); + template + slice_axis_type<0, SliceT> operator[](const SliceT& slice) const noexcept { + return {buffer(), m_layout.template slice_axis<0>(slice)}; } private: - Array& m_array; - std::unique_ptr> m_planner; + DomainArray(Buffer buffer, layout_type layout) noexcept : + ArrayBase(std::move(buffer)), + m_layout(layout) {} + + layout_type m_layout {}; }; -template -struct ArgumentHandler, M>> { - using type = ViewArgument>; +template +using Array = DomainArray, PolicyT>; - static_assert( - is_dimensionality_accepted_by_mapper, - "mapper of 'write' must return N-dimensional region" - ); +template +using SubArray = DomainArray, PolicyT>; - ArgumentHandler(Write, M> access) : - m_array(access.argument), - m_shape(m_array.shape()), - m_access_mapper(access.access_mapper) {} +namespace detail { - void initialize(const TaskGroupInit& init) { - if (!m_array.has_instance()) { - auto instance = ArrayInstance::create( // - init.runtime, - map_domain_to_distribution(m_array.shape(), init.domain, m_access_mapper), - DataType::of() - ); +/// Shared implementation of `LaunchArg` for `DomainArray`, parameterized on the access +/// mode granted to the resolved view: `AccessMode::Read` yields a `DomainView`, +/// `AccessMode::ReadWrite` a `DomainView`. +template +class LaunchArgArray { + public: + using view_element_type = std::conditional_t; + using resolve_type = DomainView>; - m_array = Array(instance); - } + explicit LaunchArgArray(const DomainArray& array) : m_array(array) {} - m_planner = std::make_unique>(m_array.instance().shared_from_this()); + void acquire(Runtime& runtime, Requisition& req) { + m_index = req.add(m_array.buffer().id(), Mode); } - type before_submit(TaskInstance& task) { - auto access_region = m_access_mapper(task.chunk, Bounds(m_shape)); - auto buffer_index = task.add_buffer_requirement( - m_planner->prepare_access(task.graph, task.memory_id, access_region, task.dependencies) - ); - - auto domain = views::dynamic_subdomain {access_region.begin(), access_region.sizes()}; - return {buffer_index, domain}; + resolve_type resolve(Runtime& runtime, Requisition& req) { + auto accessor = req.accessor(runtime, m_index); + auto* data = static_cast(accessor.address); + return resolve_type(data, m_array.layout()); } - void after_submit(const TaskSubmissionResult& result) { - m_planner->finalize_access(result.graph, result.event_id); - } - - void commit(const TaskGroupCommit& commit) { - m_planner->commit(commit.graph); - } + void release(Runtime& runtime, Requisition& req) {} private: - Array& m_array; - Dim m_shape; - M m_access_mapper; - std::unique_ptr> m_planner; + DomainArray m_array; + size_t m_index = 0; }; -template -struct ArgumentHandler>> { - using type = ViewArgument>; - - ArgumentHandler(Reduce> access) : - m_array(access.argument), - m_operation(access.op) {} - - void initialize(const TaskGroupInit& init) { - if (!m_array.has_instance()) { - auto instance = ArrayInstance::create( // - init.runtime, - map_domain_to_distribution( // - m_array.shape(), - init.domain, - All(), - true - ), - DataType::of() - ); - - m_array = Array(instance); - } - - m_planner = std::make_unique>( - m_array.instance().shared_from_this(), - m_operation - ); - } - - type before_submit(TaskInstance& task) { - auto access_region = Bounds(m_array.shape()); - - size_t buffer_index = task.add_buffer_requirement( - m_planner - ->prepare_access(task.graph, task.memory_id, access_region, 1, task.dependencies) - ); - - views::dynamic_domain domain = {access_region.sizes()}; - - return {buffer_index, domain}; - } - - void after_submit(const TaskSubmissionResult& result) { - m_planner->finalize_access(result.graph, result.event_id); - } +} // namespace detail - void commit(const TaskGroupCommit& commit) { - m_planner->commit(commit.graph); - } - - private: - Array& m_array; - Reduction m_operation; - std::unique_ptr> m_planner; +/// Read-only access to a `DomainArray` (the default when passed to `Context::scope` unwrapped). +template +class LaunchArg>: + public detail::LaunchArgArray { + public: + using detail::LaunchArgArray::LaunchArgArray; }; -template -struct ArgumentHandler, M, P>> { - static constexpr size_t K = mapper_dimensionality

; - using type = ViewArgument>; - - static_assert( - is_dimensionality_accepted_by_mapper, - "mapper of 'reduce' must return N-dimensional region" - ); - - static_assert( - is_dimensionality_accepted_by_mapper, - "private mapper of 'reduce' must return K-dimensional region" - ); - - ArgumentHandler(Reduce, M, P> access) : - m_array(access.argument), - m_operation(access.op), - m_access_mapper(access.access_mapper), - m_private_mapper(access.private_mapper) {} - - void initialize(const TaskGroupInit& init) { - if (!m_array.has_instance()) { - auto instance = ArrayInstance::create( // - init.runtime, - map_domain_to_distribution( // - m_array.shape(), - init.domain, - m_access_mapper, - true - ), - DataType::of() - ); - - m_array = Array(instance); - } - - m_planner = std::make_unique>( - m_array.instance().shared_from_this(), - m_operation - ); - } - - type before_submit(TaskInstance& task) { - auto access_region = m_access_mapper(task.chunk, Bounds(m_array.shape())); - auto private_region = m_private_mapper(task.chunk); - - auto rep = checked_cast(private_region.size()); - size_t buffer_index = task.add_buffer_requirement( - m_planner - ->prepare_access(task.graph, task.memory_id, access_region, rep, task.dependencies) - ); - - views::dynamic_subdomain domain = { - concat(private_region, access_region).begin(), - concat(private_region, access_region).sizes()}; - - return {buffer_index, domain}; - } - - void after_submit(const TaskSubmissionResult& result) { - m_planner->finalize_access(result.graph, result.event_id); - } - - void commit(const TaskGroupCommit& commit) { - m_planner->commit(commit.graph); - } +/// Read-only access to a `DomainArray` explicitly wrapped in `read(...)`. +template +class LaunchArg>>: + public detail::LaunchArgArray { + public: + explicit LaunchArg(Read> arg) : + detail::LaunchArgArray(arg.value) {} +}; - private: - Array& m_array; - Reduction m_operation; - std::unique_ptr> m_planner; - M m_access_mapper; - P m_private_mapper; +/// Read-write access to a `DomainArray` wrapped in `write(...)`. +template +class LaunchArg>>: + public detail::LaunchArgArray { + public: + explicit LaunchArg(Write> arg) : + detail::LaunchArgArray(arg.value) {} }; -} // namespace kmm \ No newline at end of file +} // namespace kmm diff --git a/include/kmm/api/array_base.hpp b/include/kmm/api/array_base.hpp new file mode 100644 index 00000000..66638aec --- /dev/null +++ b/include/kmm/api/array_base.hpp @@ -0,0 +1,36 @@ +#pragma once + +#include + +#include "kmm/api/buffer.hpp" + +namespace kmm { + +/// Base class shared by all `DomainArray` instantiations: wraps the `Buffer` backing the array's +/// data. Domain/layout information is specific to each `DomainArray` and lives there +/// instead; actual data access goes through `Context::scope`, not through this class. +class ArrayBase { + public: + ArrayBase() = default; + ArrayBase(const ArrayBase&) = default; + + const Buffer& buffer() const { + return m_buffer; + } + + Runtime runtime() const noexcept { + return m_buffer.runtime(); + } + + explicit operator bool() const { + return bool(m_buffer); + } + + protected: + explicit ArrayBase(Buffer buffer) noexcept : m_buffer(std::move(buffer)) {} + + private: + Buffer m_buffer; +}; + +} // namespace kmm diff --git a/include/kmm/api/array_instance.hpp b/include/kmm/api/array_instance.hpp deleted file mode 100644 index 0dd4091a..00000000 --- a/include/kmm/api/array_instance.hpp +++ /dev/null @@ -1,35 +0,0 @@ -#pragma once - -#include "kmm/planner/array_descriptor.hpp" - -namespace kmm { - -class Runtime; - -template -class ArrayInstance: - public ArrayDescriptor, - public std::enable_shared_from_this> { - KMM_NOT_COPYABLE_OR_MOVABLE(ArrayInstance) - - ArrayInstance(TaskGraph& stage, Runtime& rt, Distribution dist, DataType dtype); - - public: - static std::shared_ptr create(Runtime& rt, Distribution dist, DataType dtype); - ~ArrayInstance(); - - void copy_bytes_into(void* data); - void copy_bytes_from(const void* data); - void synchronize() const; - - Runtime& runtime() const { - return *m_rt; - } - - private: - std::shared_ptr m_rt; -}; - -[[noreturn]] void throw_uninitialized_array_exception(); - -} // namespace kmm \ No newline at end of file diff --git a/include/kmm/api/buffer.hpp b/include/kmm/api/buffer.hpp new file mode 100644 index 00000000..98b2f01e --- /dev/null +++ b/include/kmm/api/buffer.hpp @@ -0,0 +1,46 @@ +#pragma once + +#include +#include + +#include "kmm/runtime/buffer.hpp" +#include "kmm/runtime/identifiers.hpp" +#include "kmm/runtime/runtime.hpp" +#include "kmm/utils/refcnt_ptr.hpp" + +namespace kmm { + +class Runtime; + +class Buffer { + public: + struct Impl; + + Buffer() = default; + + /// Allocates a new buffer of `layout` within `context`. The buffer is released (via + /// `context`) once the last reference to it is dropped. If `fill_value` is non-empty, the + /// buffer's contents are set to repeated copies of it the first time it is materialized in + /// any memory. + Buffer(Runtime runtime, BufferLayout layout, std::string name, FillValue fill_value = {}); + + BufferId id() const; + BufferLayout layout() const; + Runtime runtime() const; + void prefetch(MemoryId memory_id, AccessMode mode = AccessMode::Read) const; + void poison(std::exception_ptr reason) const; + void invalidate() const; + void copy_to(void* dest, size_t nbytes, size_t offset = 0) const; + void copy_from(const void* dest, size_t nbytes, size_t offset = 0) const; + + explicit operator bool() const { + return bool(m_impl); + } + + private: + refcnt_ptr m_impl; +}; + +KMM_REFCNT_TRAITS_FWD(Buffer::Impl) + +} // namespace kmm diff --git a/include/kmm/api/context.hpp b/include/kmm/api/context.hpp new file mode 100644 index 00000000..64aa7517 --- /dev/null +++ b/include/kmm/api/context.hpp @@ -0,0 +1,104 @@ +#pragma once + +#include +#include +#include + +#include "kmm/api/array.hpp" +#include "kmm/api/buffer.hpp" +#include "kmm/api/reduce.hpp" +#include "kmm/api/scope.hpp" +#include "kmm/core/checked_compare.hpp" +#include "kmm/core/layout.hpp" +#include "kmm/runtime/identifiers.hpp" +#include "kmm/runtime/runtime.hpp" + +namespace kmm { + +/// The handle through which arrays are allocated and accessed. +class Context { + public: + /// Implicit conversion from `Runtime`, so `Runtime` can be passed where `Context` is expected. + Context(Runtime runtime) : m_runtime(std::move(runtime)) {} + + /// Allocates a new, unpopulated buffer of `layout`. The returned `Buffer` releases itself + /// (via this context) once its last reference is dropped. + Buffer allocate_buffer(BufferLayout layout, std::string name) { + return Buffer(runtime(), std::move(layout), std::move(name)); + } + + /// Allocates a new, uninitialized array of the given shape, without setting its content. + template + Array empty(Sizes... extents) { + return DomainArray, PolicyT>( + runtime(), + Shape {checked_cast(extents)...}, + PolicyT {} + ); + } + + /// Allocates a new array of the given shape fill with the given value. + template + Array fill(T value, Sizes... extents) { + return DomainArray, PolicyT>( + runtime(), + Shape {checked_cast(extents)...}, + PolicyT {}, + value + ); + } + + /// Allocates a new array of the given shape with zeros. + template + Array zeros(Sizes... extents) { + return fill(T {0}, extents...); + } + + /// Allocates a new array of the given shape filled with ones. + template + Array ones(Sizes... extents) { + return fill(T {1}, extents...); + } + + /// Allocates a new 1-D array and copies the contents of `values` into it. + template + Array from_vector(const std::vector& values) { + auto array = empty(values.size()); + array.buffer().copy_from(values.data(), values.size() * sizeof(T)); + return array; + } + + /// The `Runtime` backing this context. + Runtime& runtime() noexcept { + return m_runtime; + } + + /// Equivalent to `scope(MemoryId::host(), callback, args...)`. + template + decltype(auto) scope(F&& callback, Args&&... args) { + return scope(preferred_memory_id(), std::forward(callback), std::forward(args)...); + } + + template + decltype(auto) scope(MemoryId memory_id, F&& callback, Args&&... args) { + Scope scope { + runtime(), + memory_id, + m_root_transaction, + std::forward(args)...}; + + return std::apply(std::forward(callback), scope.resolve()); + } + + virtual MemoryId preferred_memory_id() const { + return MemoryId::host(); + } + + private: + MemoryTransaction m_root_transaction; + Runtime m_runtime; +}; + +} // namespace kmm + +#include "kmm/api/scope.hpp" diff --git a/include/kmm/api/device_context.hpp b/include/kmm/api/device_context.hpp new file mode 100644 index 00000000..15c0fa92 --- /dev/null +++ b/include/kmm/api/device_context.hpp @@ -0,0 +1,20 @@ +#pragma once + +#include "kmm/api/context.hpp" + +namespace kmm { + +class DeviceContext: public Context { + DeviceContext(Context base, DeviceId device_id) : + Context(std::move(base)), + m_device_id(device_id) {} + + MemoryId preferred_memory_id() const override { + return MemoryId::device(m_device_id); + } + + private: + DeviceId m_device_id; +}; + +} // namespace kmm \ No newline at end of file diff --git a/include/kmm/api/launch_arg.hpp b/include/kmm/api/launch_arg.hpp new file mode 100644 index 00000000..3d08e598 --- /dev/null +++ b/include/kmm/api/launch_arg.hpp @@ -0,0 +1,85 @@ +#pragma once + +#include +#include + +#include "kmm/runtime/device_event.hpp" +#include "kmm/runtime/memory_manager.hpp" + +namespace kmm { + +/// Tags `value` as needing read-write access when passed to `Context::scope` (the default for a +/// plain argument is read-only access). +template +struct Write { + T& value; +}; + +template +Write write(T& value) { + return Write {value}; +} + +/// Tags `value` as needing read-only access when passed to `Context::scope`. A plain argument is +/// already read-only by default; this is only useful to force a `Write`-tagged value back down to +/// read-only. +template +struct Read { + T& value; +}; + +template +Read read(T& value) { + return Read {value}; +} + +template +Read read(Write value) { + return Read {value.value}; +} + +template +class Reduce; + +class Runtime; +class Requisition; + +/// Default `LaunchArg` implementation, used for any `T` without a more specific specialization +/// (e.g. a plain scalar or POD struct). The caller's value is copied once into this object; there +/// is no backing buffer to acquire or release. `LaunchArg` and `LaunchArg` build on +/// top of this copy to expose it to the kernel by const or mutable reference instead of by value. +template +class LaunchArg { + public: + using resolve_type = T; + + explicit LaunchArg(T value) : m_value(std::move(value)) {} + + void acquire(Runtime& runtime, Requisition& req) {} + + resolve_type resolve(Runtime& runtime, Requisition& req) { + return m_value; + } + + void release(Runtime& runtime, Requisition& req) {} + + protected: + T m_value; +}; + +/// `LaunchArg` specialization used when `Context::scope` is given a const lvalue reference to +/// `T`. Just forwards to `LaunchArg`. +template +class LaunchArg: public LaunchArg { + public: + explicit LaunchArg(const T& value) : LaunchArg(value) {} +}; + +/// See `LaunchArg`. +template +class LaunchArg: public LaunchArg { + public: + explicit LaunchArg(T& value) : LaunchArg(value) {} +}; + +} // namespace kmm diff --git a/include/kmm/api/launcher.hpp b/include/kmm/api/launcher.hpp deleted file mode 100644 index 1779522b..00000000 --- a/include/kmm/api/launcher.hpp +++ /dev/null @@ -1,79 +0,0 @@ -#pragma once - -#include "kmm/core/domain.hpp" -#include "kmm/core/resource.hpp" - -namespace kmm { - -template -struct Host { - static constexpr ExecutionSpace execution_space = ExecutionSpace::Host; - - Host(F fun) : m_fun(fun) {} - - template - void operator()(Resource& resource, DomainChunk chunk, Args... args) { - m_fun(args...); - } - - private: - std::decay_t m_fun; -}; - -template -struct GPU { - static constexpr ExecutionSpace execution_space = ExecutionSpace::Device; - - GPU(F fun) : m_fun(fun) {} - - template - void operator()(Resource& resource, DomainChunk chunk, Args... args) { - m_fun(resource.cast(), args...); - } - - private: - std::decay_t m_fun; -}; - -template -struct GPUKernel { - static constexpr ExecutionSpace execution_space = ExecutionSpace::Device; - - GPUKernel(F kernel, dim3 block_size) : GPUKernel(kernel, block_size, block_size) {} - - GPUKernel(F kernel, dim3 block_size, dim3 elements_per_block, uint32_t shared_memory = 0) : - kernel(kernel), - block_size(block_size), - elements_per_block(elements_per_block), - shared_memory(shared_memory) {} - - template - void operator()(Resource& resource, DomainChunk chunk, Args... args) { - int64_t g[3] = { - chunk.size.get_or_default(0), - chunk.size.get_or_default(1), - chunk.size.get_or_default(2)}; - int64_t b[3] = {elements_per_block.x, elements_per_block.y, elements_per_block.z}; - dim3 grid_dim = { - checked_cast((g[0] / b[0]) + int64_t(g[0] % b[0] != 0)), - checked_cast((g[1] / b[1]) + int64_t(g[1] % b[1] != 0)), - checked_cast((g[2] / b[2]) + int64_t(g[2] % b[2] != 0)), - }; - - resource.cast().launch( // - grid_dim, - block_size, - shared_memory, - kernel, - args... - ); - } - - private: - std::decay_t kernel; - dim3 block_size; - dim3 elements_per_block; - uint32_t shared_memory; -}; - -} // namespace kmm diff --git a/include/kmm/api/mapper.hpp b/include/kmm/api/mapper.hpp deleted file mode 100644 index b7e5bace..00000000 --- a/include/kmm/api/mapper.hpp +++ /dev/null @@ -1,268 +0,0 @@ -#pragma once - -#include "spdlog/spdlog.h" - -#include "kmm/api/argument.hpp" -#include "kmm/core/domain.hpp" -#include "kmm/core/reduction.hpp" -#include "kmm/utils/geometry.hpp" -#include "kmm/utils/integer_fun.hpp" - -namespace kmm { - -struct All { - template - Bounds operator()(DomainChunk chunk, Bounds bounds) const { - return bounds; - } -}; - -struct Axis { - constexpr Axis() : m_axis(0) {} - explicit constexpr Axis(size_t axis) : m_axis(axis) {} - - Bounds<1> operator()(DomainChunk chunk) const { - return Bounds<1>::from_offset_size( - chunk.offset.get_or_default(m_axis), - chunk.size.get_or_default(m_axis) - ); - } - - Bounds<1> operator()(DomainChunk chunk, Bounds<1> bounds) const { - return (*this)(chunk).intersection(bounds); - } - - size_t get() const { - return m_axis; - } - - explicit operator size_t() const { - return get(); - } - - private: - size_t m_axis = 0; -}; - -struct IdentityMap { - template - Bounds operator()(DomainChunk chunk, Bounds bounds) const { - return Bounds::from_offset_size(Index::from(chunk.offset), Dim::from(chunk.size)); - } -}; - -// (scale * variable + offset + [0...length]) / divisor -struct IndexMap { - constexpr IndexMap(Axis variable = {}) : m_axis(variable) {} - - IndexMap( - Axis variable, - int64_t scale, - int64_t offset = 0, - int64_t length = 1, - int64_t divisor = 1 - ); - - static IndexMap range(IndexMap begin, IndexMap end); - IndexMap offset_by(int64_t offset) const; - IndexMap scale_by(int64_t factor) const; - IndexMap divide_by(int64_t divisor) const; - IndexMap negate() const; - Bounds<1> apply(DomainChunk chunk) const; - - Bounds<1> operator()(DomainChunk chunk) const { - return apply(chunk); - } - - Bounds<1> operator()(DomainChunk chunk, Bounds<1> bounds) const { - return apply(chunk).intersection(bounds); - } - - friend std::ostream& operator<<(std::ostream& f, const IndexMap& that); - - private: - Axis m_axis = {}; - int64_t m_scale = 1; - int64_t m_offset = 0; - int64_t m_length = 1; - int64_t m_divisor = 1; -}; - -inline IndexMap range(IndexMap begin, IndexMap end) { - return IndexMap::range(begin, end); -} - -inline IndexMap range(int64_t begin, int64_t end) { - return {Axis {}, 0, begin, end - begin}; -} - -inline IndexMap range(int64_t end) { - return range(0, end); -} - -inline IndexMap operator+(IndexMap a) { - return a; -} - -inline IndexMap operator+(IndexMap a, int64_t b) { - return a.offset_by(b); -} - -inline IndexMap operator+(int64_t a, IndexMap b) { - return b.offset_by(a); -} - -inline IndexMap operator-(IndexMap a) { - return a.negate(); -} - -inline IndexMap operator-(IndexMap a, int64_t b) { - return a + (-b); -} - -inline IndexMap operator-(int64_t a, IndexMap b) { - return a + (-b); -} - -inline IndexMap operator*(IndexMap a, int64_t b) { - return a.scale_by(b); -} - -inline IndexMap operator*(int64_t a, IndexMap b) { - return b.scale_by(a); -} - -inline IndexMap operator/(IndexMap a, int64_t b) { - return a.divide_by(b); -} - -template -struct MultiIndexMap { - Bounds operator()(DomainChunk chunk) const { - Bounds result; - - for (size_t i = 0; i < N; i++) { - result[i] = (this->axes[i])(chunk); - } - - return result; - } - - Bounds operator()(DomainChunk chunk, Bounds bounds) const { - Bounds result; - - for (size_t i = 0; i < N; i++) { - result[i] = (this->axes[i])(chunk, Bounds<1> {bounds[i]}); - } - - return result; - } - - IndexMap axes[N]; -}; - -template<> -struct MultiIndexMap<0> { - Bounds<0> operator()(DomainChunk chunk, Bounds<0> bounds = {}) const { - return {}; - } -}; - -inline IndexMap into_index_map(int64_t m) { - return {Axis {}, 0, m}; -} - -inline IndexMap into_index_map(Axis m) { - return m; -} - -inline IndexMap into_index_map(IndexMap m) { - return m; -} - -inline IndexMap into_index_map(All m) { - return {Axis(), 0, 0, std::numeric_limits::max()}; -} - -template -MultiIndexMap bounds(const Is&... slices) { - return {into_index_map(slices)...}; -} - -template -MultiIndexMap tile(const Is&... length) { - size_t variable = 0; - return {IndexMap( - Axis {variable++}, - checked_cast(length), - 0, - checked_cast(length) - )...}; -} - -namespace placeholders { -static constexpr All _; - -static constexpr Axis _x = Axis(0); -static constexpr Axis _y = Axis(1); -static constexpr Axis _z = Axis(2); - -static constexpr Axis _i = Axis(0); -static constexpr Axis _j = Axis(1); -static constexpr Axis _k = Axis(2); - -static constexpr Axis _0 = Axis(0); -static constexpr Axis _1 = Axis(1); -static constexpr Axis _2 = Axis(2); - -static constexpr MultiIndexMap<2> _xy = {_x, _y}; -static constexpr MultiIndexMap<3> _xyz = {_x, _y, _z}; - -static constexpr MultiIndexMap<2> _ij = {_x, _y}; -static constexpr MultiIndexMap<3> _ijk = {_x, _y, _z}; - -static constexpr IdentityMap one_to_one; -static constexpr All all; -} // namespace placeholders - -template<> -struct Argument: Argument> { - static Argument pack(TaskInstance& builder, IndexMap mapper) { - return {mapper(builder.chunk).get_or_default(0)}; - } -}; - -template<> -struct Argument: Argument> { - static Argument pack(TaskInstance& builder, Axis mapper) { - return {mapper(builder.chunk).get_or_default(0)}; - } -}; - -template -struct Argument>: Argument> { - static Argument pack(TaskInstance& builder, MultiIndexMap mapper) { - return {mapper(builder.chunk)}; - } -}; - -namespace detail { -template -struct RangeDim: std::integral_constant {}; - -template -struct RangeDim>: std::integral_constant {}; -} // namespace detail - -template -static constexpr size_t mapper_dimensionality = - detail::RangeDim>::value; - -template -static constexpr bool is_dimensionality_accepted_by_mapper = - detail::RangeDim>>::value == N; - -} // namespace kmm - -template<> -struct fmt::formatter: fmt::ostream_formatter {}; \ No newline at end of file diff --git a/include/kmm/api/parallel_submit.hpp b/include/kmm/api/parallel_submit.hpp deleted file mode 100644 index 6ad89eed..00000000 --- a/include/kmm/api/parallel_submit.hpp +++ /dev/null @@ -1,127 +0,0 @@ -#pragma once - -#include "kmm/api/argument.hpp" -#include "kmm/api/task_group.hpp" -#include "kmm/core/buffer.hpp" -#include "kmm/core/domain.hpp" -#include "kmm/core/identifiers.hpp" -#include "kmm/runtime/runtime.hpp" - -namespace kmm { - -class Runtime; -class TaskGraphState; - -namespace detail { - -template -class ComputeTaskImpl: public ComputeTask { - public: - ComputeTaskImpl(DomainChunk chunk, Launcher launcher, Args... args) : - m_chunk(chunk), - m_launcher(std::move(launcher)), - m_args(std::move(args)...) {} - - void execute(Resource& resource, TaskContext context) override { - execute_impl(std::index_sequence_for(), resource, context); - } - - template - void execute_impl(std::index_sequence, Resource& resource, TaskContext& context) { - static constexpr ExecutionSpace execution_space = Launcher::execution_space; - - m_launcher( - resource, - m_chunk, - ArgumentUnpack::call(context, std::get(m_args))... - ); - } - - private: - DomainChunk m_chunk; - Launcher m_launcher; - std::tuple m_args; -}; - -template -EventId parallel_submit_impl( - std::index_sequence, - Runtime& runtime, - const SystemInfo& system_info, - const Domain& domain, - Launcher launcher, - Args&&... args -) { - std::tuple...> handlers = {std::forward(args)...}; - - auto init = TaskGroupInit { - .runtime = runtime, // - .domain = domain}; - - (std::get(handlers).initialize(init), ...); - - return runtime.schedule([&](TaskGraph& graph) { - EventList events; - - for (const DomainChunk& chunk : domain.chunks) { - auto processor_id = chunk.owner_id; - - auto instance = TaskInstance { - .runtime = runtime, - .graph = graph, - .chunk = chunk, - .memory_id = system_info.affinity_memory(processor_id), - .buffers = {}, - .dependencies = {}}; - - auto task = std::make_unique...>>( - chunk, - launcher, - std::get(handlers).before_submit(instance)... - ); - - EventId event_id = graph.insert_compute_task( - processor_id, - std::move(task), - std::move(instance.buffers), - std::move(instance.dependencies) - ); - - events.push_back(event_id); - - auto result = TaskSubmissionResult { - .runtime = runtime, // - .graph = graph, - .event_id = event_id}; - - (std::get(handlers).after_submit(result), ...); - } - - auto commit = TaskGroupCommit {.runtime = runtime, .graph = graph}; - - (std::get(handlers).commit(commit), ...); - - return graph.join_events(events); - }); -} -} // namespace detail - -template -EventId parallel_submit( - Runtime& runtime, - const SystemInfo& system_info, - const Domain& partition, - Launcher launcher, - Args&&... args -) { - return detail::parallel_submit_impl( - std::index_sequence_for {}, - runtime, - system_info, - partition, - launcher, - std::forward(args)... - ); -} - -} // namespace kmm diff --git a/include/kmm/api/reduce.hpp b/include/kmm/api/reduce.hpp new file mode 100644 index 00000000..bc5e1dd4 --- /dev/null +++ b/include/kmm/api/reduce.hpp @@ -0,0 +1,140 @@ +#pragma once + +#include +#include + +#include "kmm/api/array.hpp" + +namespace kmm { + +/// A `DomainArray` that has been put into reduction mode -- constructed as +/// `Reduce(array)` or `Reduce(array, op)` (`op` defaults to `ReductionOp::Sum`). Only +/// reduce-mode access is possible on it (there is no `LaunchArg` specialization for plain +/// read/write access) -- attempting to read or write it is caught at compile time, not by an +/// implicit auto-finalize. Call `finalize()` to fold the accumulated values back into a regular +/// `DomainArray`, or just let it go out of scope: reduction mode is scoped to this object's +/// lifetime, so the destructor finalizes automatically if you don't. Move-only: exactly one +/// `Reduce` owns "when does this reduction end" at a time. +template +class Reduce> { + public: + using element_type = T; + using domain_type = DomainT; + using policy_type = PolicyT; + using layout_type = Layout; + + Reduce() = default; + + explicit Reduce( + const DomainArray& array, + ReductionOp op = ReductionOp::Sum + ) : + m_layout(array.layout()), + m_op(op), + m_buffer(array.buffer()) { + if (m_buffer) { + m_buffer.runtime().begin_reduction(m_buffer.id(), data_type_of(), op); + } + } + + Reduce(const Reduce&) = delete; + Reduce& operator=(const Reduce&) = delete; + + Reduce(Reduce&& that) noexcept : + m_layout(that.m_layout), + m_op(that.m_op), + m_buffer(std::move(that.m_buffer)) {} + + Reduce& operator=(Reduce&& that) noexcept { + if (this != &that) { + finish(); + m_buffer = std::move(that.m_buffer); + m_layout = that.m_layout; + m_op = that.m_op; + } + + return *this; + } + + ~Reduce() { + finish(); + } + + const layout_type& layout() const noexcept { + return m_layout; + } + + ReductionOp op() const noexcept { + return m_op; + } + + /// Finalizes the reduction. + DomainArray finalize() { + m_buffer.runtime().finalize_reduction(m_buffer.id()); + return DomainArray(std::move(m_buffer), m_layout); + } + + private: + void finish() noexcept { + if (!m_buffer) { + return; + } + + auto buffer = std::move(m_buffer); + bool unwinding = std::uncaught_exceptions() > 0; + + try { + if (!unwinding) { + buffer.runtime().finalize_reduction(buffer.id()); + } else { + buffer.runtime().rollback_reduction(buffer.id()); + } + } catch (...) { + buffer.runtime().rollback_reduction(buffer.id()); + } + } + + layout_type m_layout {}; + ReductionOp m_op = ReductionOp::Sum; + Buffer m_buffer; +}; + +/// Deduction guide: CTAD only ever considers the primary `Reduce` template (intentionally +/// left undefined -- see `launch_arg.hpp`), so this specialization must spell out its own guide +/// for `Reduce(array)` / `Reduce(array, op)` to deduce through it. +template +Reduce(const DomainArray&, ReductionOp = ReductionOp::Sum) + -> Reduce>; + +/// Reduce-mode access to a `Reduce`-wrapped array. There is no `LaunchArg` specialization for +/// plain `DomainArray<...>` reduce access -- a buffer must be put into reduction mode first (see +/// `Reduce`'s constructor), which is what makes reduce access to it representable at all. Holds +/// a reference rather than a copy since `Reduce` is move-only; safe because every use is confined +/// to `Context::scope`'s own call, which completes before any temporary argument is destroyed. +template +class LaunchArg>> { + public: + using resolve_type = DomainView>; + + explicit LaunchArg(const Reduce>& array) : m_array(array) {} + + void acquire(Runtime& runtime, Requisition& req) { + m_index = req.add_reduction(m_array.buffer().id()); + } + + // Assumes `req` has already been polled to `Poll::Ready` by `Context::scope` -- see the + // matching comment on `detail::LaunchArgArray::resolve`. + resolve_type resolve(Runtime& runtime, Requisition& req) { + auto accessor = req.accessor(runtime, m_index); + auto* data = static_cast(accessor.address); + return resolve_type(data, m_array.layout()); + } + + void release(Runtime& runtime, Requisition& req) {} + + private: + const Reduce>& m_array; + size_t m_index = 0; +}; + +} // namespace kmm diff --git a/include/kmm/api/runtime_handle.hpp b/include/kmm/api/runtime_handle.hpp deleted file mode 100644 index 1e5af9ce..00000000 --- a/include/kmm/api/runtime_handle.hpp +++ /dev/null @@ -1,261 +0,0 @@ -#pragma once - -#include -#include - -#include "kmm/api/array.hpp" -#include "kmm/api/launcher.hpp" -#include "kmm/api/parallel_submit.hpp" -#include "kmm/api/struct_argument.hpp" -#include "kmm/core/config.hpp" -#include "kmm/core/system_info.hpp" -#include "kmm/core/view.hpp" -#include "kmm/utils/checked_math.hpp" -#include "kmm/utils/panic.hpp" -#include "kmm/utils/range.hpp" - -namespace kmm { - -class Runtime; - -class RuntimeHandle { - struct Impl; - RuntimeHandle(std::shared_ptr impl); - - public: - RuntimeHandle(std::shared_ptr rt); - RuntimeHandle(Runtime& rt); - - /** - * Submit a single task to the runtime system. - * - * @param index_space The index space defining the task dimensions. - * @param target The target processor for the task. - * @param launcher The task launcher. - * @param args The arguments that are forwarded to the launcher. - * @return The event identifier for the submitted task. - */ - template - EventId submit(ResourceId target, L&& launcher, Args&&... args) const { - DomainChunk chunk = { - .owner_id = target, // - .offset = DomainIndex::zero(), - .size = DomainDim::one()}; - - return kmm::parallel_submit( - worker(), - info(), - Domain {{chunk}}, - std::forward(launcher), - std::forward(args)... - ); - } - - /** - * Submit a set of tasks to the runtime systems. - * - * @param dist The domain describing how the work is split. - * @param launcher The task launcher. - * @param args The arguments that are forwarded to the launcher. - * @return The event identifier for the submitted task. - */ - template - EventId parallel_submit(D&& domain, L&& launcher, Args&&... args) const { - return kmm::parallel_submit( - worker(), - info(), - IntoDomain>::call( - std::forward(domain), - info(), - std::decay_t::execution_space - ), - std::forward(launcher), - std::forward(args)... - ); - } - - /** - * Submit a set of tasks to the runtime systems. - * - * @param domain_size The index space defining the domain dimensions. - * @param partitioner The partitioner describing how the work is split. - * @param launcher The task launcher. - * @param args The arguments that are forwarded to the launcher. - * @return The event identifier for the submitted task. - */ - template - EventId parallel_submit(Dim domain_size, Dim chunk_size, L&& launcher, Args&&... args) - const { - return this->parallel_submit( - TileDomain(domain_size, chunk_size), - std::forward(launcher), - std::forward(args)... - ); - } - - /** - * Allocates an array in memory with the given shape and memory affinity. - * - * The pointer to the given buffer should contain `shape[0] * shape[1] * shape[2]...` - * elements. - * - * @param data Pointer to the array data. - * @param shape Shape of the array. - * @param memory_id Identifier of the memory region. - * @return The allocated Array object. - */ - template - Array allocate(const T* data, Dim shape, MemoryId memory_id) const { - auto handle = ArrayInstance::create( - worker(), - Distribution {shape, shape, {memory_id}}, - DataType::of() - ); - - handle->copy_bytes_from(data); - return Array {std::move(handle)}; - } - - /** - * Allocates an array in memory with the given shape. - * - * The pointer to the given buffer should contain `shape[0] * shape[1] * shape[2]...` - * elements. - * - * In which memory the data will be allocated is determined by `memory_affinity_for_address`. - * - * @param data Pointer to the array data. - * @param shape Shape of the array. - * @return The allocated Array object. - */ - template - Array allocate(const T* data, Dim shape) const { - return allocate(data, shape, memory_affinity_for_address(data)); - } - - /** - * Alias for `allocate(v.data(), v.sizes())` - */ - template - Array allocate(View v) const { - return allocate(v.data(), v.sizes()); - } - - /** - * Alias for `allocate(data, Dim{sizes...})` - */ - template - Array allocate(const T* data, const Is&... num_elements) const { - return allocate(data, Dim {checked_cast(num_elements)...}); - } - - /** - * Alias for `allocate(v.data(), v.size())` - */ - template - Array allocate(const std::vector& v) const { - return allocate(v.data(), v.size()); - } - - /** - * Alias for `allocate(v.begin(), v.size())` - */ - template - Array allocate(std::initializer_list v) const { - return allocate(v.begin(), v.size()); - } - - /** - * Returns the memory affinity for a given address. - */ - MemoryId memory_affinity_for_address(const void* address) const; - - /** - * Returns a new event that triggers when all the given events have triggered. - */ - EventId join(EventList events) const; - - /** - * Returns a new event that triggers when all the given events have triggered. Each argument - * must be convertible to an `EventId`. - */ - template - EventId join(Es... events) const { - return join(EventList(EventId(events)...)); - } - - /** - * Returns `true` if the event with the provided identifier has finished, or `false` otherwise. - */ - bool is_done(EventId) const; - - /** - * Block the current thread until the event with the provided identifier completes. - */ - void wait(EventId id) const; - - /** - * Block the current thread until the event with the provided id completes. Blocks until - * either the event completes or the deadline is exceeded, whatever comes first. - * - * @return `true` if the event with the provided id has finished, otherwise returns `false`. - */ - bool wait_until(EventId id, std::chrono::system_clock::time_point deadline) const; - - /** - * Block the current thread until the event with the provided id completes. Blocks until - * either the event completes or the duration is exceeded, whatever comes first. - * - * @return `true` if the event with the provided id has finished, otherwise returns `false`. - */ - bool wait_for(EventId id, std::chrono::system_clock::duration duration) const; - - /** - * Submit a barrier the runtime system. The barrier completes once all the tasks submitted - * to the runtime system so far have finished. - * - * @return The identifier of the barrier. - */ - EventId barrier() const; - - /** - * Blocks until all the tasks submitted to the runtime system have finished and the - * system has become idle. - */ - void synchronize() const; - - /** - * Return a new `RuntimeHandle` that is constrained to the given set of resources. In other - * words, it only can submit work onto those resources. - */ - RuntimeHandle constrain_to(std::vector resources) const; - - /** - * Return a new `RuntimeHandle` that is constrained to the given device. In other - * words, it only can submit work onto that device. - */ - RuntimeHandle constrain_to(DeviceId device) const; - - /** - * Return a new `RuntimeHandle` that is constrained to the given resource. In other - * words, it only can submit work onto that resource. - */ - RuntimeHandle constrain_to(ResourceId resource) const; - - /** - * Returns information about the current system. - */ - const SystemInfo& info() const; - - /** - * Returns the inner `Worker`. - */ - Runtime& worker() const; - - private: - std::shared_ptr m_data; -}; - -RuntimeHandle make_runtime(const RuntimeConfig& config = default_config_from_environment()); - -} // namespace kmm diff --git a/include/kmm/api/scope.hpp b/include/kmm/api/scope.hpp new file mode 100644 index 00000000..5c55cb71 --- /dev/null +++ b/include/kmm/api/scope.hpp @@ -0,0 +1,73 @@ +#pragma once + +#include +#include +#include +#include + +#include "kmm/api/launch_arg.hpp" +#include "kmm/core/macros.hpp" +#include "kmm/core/panic.hpp" +#include "kmm/runtime/device_event.hpp" +#include "kmm/runtime/identifiers.hpp" +#include "kmm/runtime/requisition.hpp" +#include "kmm/runtime/runtime.hpp" +#include "kmm/utils/poll.hpp" + +namespace kmm { + +template +class Scope { + KMM_NOT_COPYABLE_OR_MOVABLE(Scope) + + public: + explicit Scope(Runtime& runtime, MemoryId memory_id, MemoryTransaction parent, Args&&... args) : + m_runtime(runtime), + m_requisition(memory_id, std::move(parent)), + resources(LaunchArg(std::forward(args))...) { + acquire(std::index_sequence_for {}); + + m_runtime.submit(m_requisition); + } + + std::tuple::resolve_type...> resolve() { + if (m_requisition.stream().is_null()) { + m_runtime.synchronize(m_requisition.dependencies()); + } else { + m_requisition.stream().wait_on_events(m_requisition.dependencies()); + } + + return resolve(std::index_sequence_for {}); + } + + void release(DeviceEventSet deps) { + release_impl(std::index_sequence_for {}); + m_runtime.release(m_requisition, deps); + } + + ~Scope() { + release({}); + } + + private: + template + void acquire(std::index_sequence) { + (std::get(resources).acquire(m_runtime, m_requisition), ...); + } + + template + std::tuple::resolve_type...> resolve(std::index_sequence) { + return std::make_tuple(std::get(resources).resolve(m_runtime, m_requisition)...); + } + + template + void release_impl(std::index_sequence) { + (std::get(resources).release(m_runtime, m_requisition), ...); + } + + Runtime m_runtime; + Requisition m_requisition; + std::tuple...> resources; +}; + +} // namespace kmm diff --git a/include/kmm/api/struct_argument.hpp b/include/kmm/api/struct_argument.hpp deleted file mode 100644 index 5e3b2381..00000000 --- a/include/kmm/api/struct_argument.hpp +++ /dev/null @@ -1,94 +0,0 @@ -#pragma once - -#include "kmm/api/argument.hpp" - -namespace kmm { - -template -struct StructArgument { - static constexpr size_t num_fields = sizeof...(Fields); - std::tuple fields; -}; - -template -struct StructArgumentHandler { - using type = StructArgument::type...>; - - StructArgumentHandler(const Type&, Fields... fields) : m_handlers(fields...) {} - - void initialize(const TaskGroupInit& init) { - initialize_impl(init, std::make_index_sequence()); - } - - type before_submit(TaskInstance& task) { - return before_submit_impl(task, std::make_index_sequence()); - } - - void after_submit(const TaskSubmissionResult& result) { - after_submit_impl(result, std::make_index_sequence()); - } - - void commit(const TaskGroupCommit& commit) { - commit_impl(commit, std::make_index_sequence()); - } - - private: - template - void initialize_impl(const TaskGroupInit& init, std::index_sequence) { - (std::get(m_handlers).initialize(init), ...); - } - - template - type before_submit_impl(TaskInstance& task, std::index_sequence) { - return {.fields = {(std::get(m_handlers).before_submit(task))...}}; - } - - template - void after_submit_impl(const TaskSubmissionResult& result, std::index_sequence) { - (std::get(m_handlers).after_submit(result), ...); - } - - template - void commit_impl(const TaskGroupCommit& commit, std::index_sequence) { - (std::get(m_handlers).commit(commit), ...); - } - - std::tuple...> m_handlers; -}; - -template -struct StructArgumentUnpack { - static View call(TaskContext& context, StructArgument& data) { - return call_impl(context, data, std::index_sequence_for()); - } - - private: - template - static View - call_impl(TaskContext& context, StructArgument& data, std::index_sequence) { - return {ArgumentUnpack::call(context, std::get(data.fields))...}; - } -}; - -} // namespace kmm - -#define KMM_DEFINE_STRUCT_ARGUMENT_IMPL(UNIQUE_NAME, T, ...) \ - static auto UNIQUE_NAME(const T& it) { \ - return kmm::StructArgumentHandler(it, __VA_ARGS__); \ - } \ - template<> \ - struct kmm::ArgumentHandler: decltype(UNIQUE_NAME(std::declval())) { \ - ArgumentHandler(const T& it) : decltype(UNIQUE_NAME(it))(it, __VA_ARGS__) {} \ - }; - -#define KMM_DEFINE_STRUCT_ARGUMENT(T, ...) \ - KMM_DEFINE_STRUCT_ARGUMENT_IMPL( \ - KMM_CONCAT(__kmm_argument_handle_type_helper_, __LINE__), \ - T, \ - __VA_ARGS__ \ - ) - -#define KMM_DEFINE_STRUCT_VIEW(T, V) \ - template \ - struct kmm::ArgumentUnpack>: \ - kmm::StructArgumentUnpack {}; diff --git a/include/kmm/api/task_group.hpp b/include/kmm/api/task_group.hpp deleted file mode 100644 index 5bf7492c..00000000 --- a/include/kmm/api/task_group.hpp +++ /dev/null @@ -1,51 +0,0 @@ -#pragma once - -#include "kmm/core/domain.hpp" -#include "kmm/core/identifiers.hpp" - -namespace kmm { - -class Runtime; -class TaskGraph; - -struct TaskGroupInit { - KMM_NOT_COPYABLE_OR_MOVABLE(TaskGroupInit) - - public: - Runtime& runtime; - const Domain& domain; -}; - -struct TaskInstance { - KMM_NOT_COPYABLE_OR_MOVABLE(TaskInstance) - - public: - Runtime& runtime; - TaskGraph& graph; - DomainChunk chunk; - MemoryId memory_id; - std::vector buffers; - EventList dependencies; - - size_t add_buffer_requirement(BufferRequirement req) { - size_t index = buffers.size(); - buffers.push_back(std::move(req)); - return index; - } -}; - -struct TaskSubmissionResult { - KMM_NOT_COPYABLE_OR_MOVABLE(TaskSubmissionResult) - - public: - Runtime& runtime; - TaskGraph& graph; - EventId event_id; -}; - -struct TaskGroupCommit { - Runtime& runtime; - TaskGraph& graph; -}; - -} // namespace kmm \ No newline at end of file diff --git a/include/kmm/api/view_argument.hpp b/include/kmm/api/view_argument.hpp deleted file mode 100644 index ee1d1540..00000000 --- a/include/kmm/api/view_argument.hpp +++ /dev/null @@ -1,67 +0,0 @@ -#pragma once - -#include "kmm/api/argument.hpp" -#include "kmm/core/view.hpp" - -namespace kmm { - -template> -struct ViewArgument { - using value_type = T; - using domain_type = D; - using layout_type = L; - - ViewArgument(size_t buffer_index, D domain, L layout) : - buffer_index(buffer_index), - domain(domain), - layout(layout) {} - - ViewArgument(size_t buffer_index, D domain) : - ViewArgument(buffer_index, domain, L::from_domain(domain)) {} - - size_t buffer_index; - D domain; - L layout; -}; - -template -struct ArgumentUnpack> { - using type = AbstractView; - - static type call(const TaskContext& context, ViewArgument arg) { - T* data = static_cast(context.accessors.at(arg.buffer_index).address); - return {data, arg.domain, arg.layout}; - } -}; - -template -struct ArgumentUnpack> { - using type = AbstractView; - - static type call(const TaskContext& context, ViewArgument arg) { - const T* data = static_cast(context.accessors.at(arg.buffer_index).address); - return {data, arg.domain, arg.layout}; - } -}; - -template -struct ArgumentUnpack> { - using type = AbstractView; - - static type call(const TaskContext& context, ViewArgument arg) { - T* data = static_cast(context.accessors.at(arg.buffer_index).address); - return {data, arg.domain, arg.layout}; - } -}; - -template -struct ArgumentUnpack> { - using type = AbstractView; - - static type call(const TaskContext& context, ViewArgument arg) { - const T* data = static_cast(context.accessors.at(arg.buffer_index).address); - return {data, arg.domain, arg.layout}; - } -}; - -} // namespace kmm \ No newline at end of file diff --git a/include/kmm/core/backends.hpp b/include/kmm/core/backends.hpp deleted file mode 100644 index 834ac1c2..00000000 --- a/include/kmm/core/backends.hpp +++ /dev/null @@ -1,438 +0,0 @@ -#pragma once - -#include - -#include "kmm/utils/macros.hpp" - -#ifdef KMM_USE_CUDA - #include - #include - #include - #include - #include -#elif KMM_USE_HIP - #include - #include - #include - #include -#endif - -namespace kmm { - -#ifdef KMM_USE_CUDA -using half_type = __half; -using bfloat16_type = __nv_bfloat16; - - #define GPU_DEVICE_ATTRIBUTE_MAX CU_DEVICE_ATTRIBUTE_MAX - #define GPU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z CU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z - #define GPU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR \ - CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR - #define GPU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR \ - CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR - #define GPU_MEMHOSTALLOC_PORTABLE CU_MEMHOSTALLOC_PORTABLE - #define GPU_MEMHOSTALLOC_DEVICEMAP CU_MEMHOSTALLOC_DEVICEMAP - #define GPU_SUCCESS CUDA_SUCCESS - #define GPU_ERROR_OUT_OF_MEMORY CUDA_ERROR_OUT_OF_MEMORY - #define GPU_MEM_ALLOCATION_TYPE_PINNED CU_MEM_ALLOCATION_TYPE_PINNED - #define GPU_MEM_HANDLE_TYPE_NONE CU_MEM_HANDLE_TYPE_NONE - #define GPU_MEM_LOCATION_TYPE_DEVICE CU_MEM_LOCATION_TYPE_DEVICE - #define GPU_ERROR_UNKNOWN CUDA_ERROR_UNKNOWN - #define GPU_STREAM_NON_BLOCKING CU_STREAM_NON_BLOCKING - #define GPU_EVENT_WAIT_DEFAULT CU_EVENT_WAIT_DEFAULT - #define GPU_ERROR_NOT_READY CUDA_ERROR_NOT_READY - #define GPU_EVENT_DISABLE_TIMING CU_EVENT_DISABLE_TIMING - #define GPU_MEMORYTYPE_HOST CU_MEMORYTYPE_HOST - #define GPU_MEMORYTYPE_DEVICE CU_MEMORYTYPE_DEVICE - #define GPU_ERROR_NO_DEVICE CUDA_ERROR_NO_DEVICE - #define GPU_POINTER_ATTRIBUTE_MEMORY_TYPE CU_POINTER_ATTRIBUTE_MEMORY_TYPE - #define GPU_POINTER_ATTRIBUTE_DEVICE_ORDINAL CU_POINTER_ATTRIBUTE_DEVICE_ORDINAL - #define GPU_CTX_MAP_HOST CU_CTX_MAP_HOST - #define gpuCtxGetDevice cuCtxGetDevice - #define gpuDeviceGetName cuDeviceGetName - #define gpuDeviceGetAttribute cuDeviceGetAttribute - #define gpuMemGetInfo cuMemGetInfo - #define gpuMemcpyHtoDAsync cuMemcpyHtoDAsync - #define gpuMemcpyDtoHAsync cuMemcpyDtoHAsync - #define gpuStreamSynchronize cuStreamSynchronize - #define gpuMemsetD8Async cuMemsetD8Async - #define gpuMemsetD16Async cuMemsetD16Async - #define gpuMemsetD32Async cuMemsetD32Async - #define gpuMemcpyAsync cuMemcpyAsync - #define gpuMemHostAlloc cuMemHostAlloc - #define gpuMemFreeHost cuMemFreeHost - #define gpuMemAlloc cuMemAlloc - #define gpuMemFree cuMemFree - #define gpuMemPoolCreate cuMemPoolCreate - #define gpuMemPoolDestroy cuMemPoolDestroy - #define gpuMemAllocFromPoolAsync cuMemAllocFromPoolAsync - #define gpuMemFreeAsync cuMemFreeAsync - #define gpuGetStreamPriorityRange cuCtxGetStreamPriorityRange - #define gpuStreamCreateWithPriority cuStreamCreateWithPriority - #define gpuStreamQuery cuStreamQuery - #define gpuStreamDestroy cuStreamDestroy - #define gpuEventSynchronize cuEventSynchronize - #define gpuEventRecord cuEventRecord - #define gpuStreamWaitEvent cuStreamWaitEvent - #define gpuEventQuery cuEventQuery - #define gpuEventDestroy cuEventDestroy - #define gpuEventCreate cuEventCreate - #define gpuMemcpyHtoD cuMemcpyHtoD - #define gpuMemcpy2DAsync cuMemcpy2DAsync - #define gpuMemcpy2D cuMemcpy2D - #define gpuMemcpyDtoH cuMemcpyDtoH - #define gpuMemcpyDtoDAsync cuMemcpyDtoDAsync - #define gpuMemcpyDtoD cuMemcpyDtoD - #define gpuGetErrorName cuGetErrorName - #define gpuGetErrorString cuGetErrorString - #define GPUrtGetErrorName cudaGetErrorName - #define GPUrtGetErrorString cudaGetErrorString - #define gpuGetLastError cudaGetLastError - #define gpuInit cuInit - #define gpuDeviceGetCount cuDeviceGetCount - #define gpuDeviceGet cuDeviceGet - #define gpuCtxCreate cuCtxCreate - #define gpuCtxDestroy cuCtxDestroy - #define gpuDevicePrimaryCtxRetain cuDevicePrimaryCtxRetain - #define gpuDevicePrimaryCtxRelease cuDevicePrimaryCtxRelease - #define gpuCtxPushCurrent cuCtxPushCurrent - #define gpuCtxPopCurrent cuCtxPopCurrent - #define GPUrtLaunchKernel cudaLaunchKernel - #define gpuMemPoolTrimTo cuMemPoolTrimTo - #define gpuDeviceGetDefaultMemPool cuDeviceGetDefaultMemPool - #define gpuPointerGetAttribute cuPointerGetAttribute - -using GPUresult = CUresult; -using gpuError_t = cudaError_t; -using GPUdevice = CUdevice; -using GPUdevice_attribute = CUdevice_attribute; -using GPUcontext = CUcontext; -using GPUmemorytype = CUmemorytype; -using GPUstream_t = CUstream; -using GPUdeviceptr = CUdeviceptr; -using GPUmemoryPool = CUmemoryPool; -using GPUmemPoolProps = CUmemPoolProps; -using GPUmemAllocationType = CUmemAllocationType; -using GPUmemAllocationHandleType = CUmemAllocationHandleType; -using GPUmemLocationType = CUmemLocationType; -using GPUevent_t = CUevent; -using GPU_MEMCPY2D = CUDA_MEMCPY2D; - -GPUresult gpuMemcpyPeerAsync(GPUdeviceptr, GPUcontext, GPUdevice, GPUdeviceptr, GPUcontext, GPUdevice, size_t, GPUstream_t); - - // cuBLAS - #define blasCreate cublasCreate - #define blasSetStream cublasSetStream - #define blasDestroy cublasDestroy - #define blasGetStatusName cublasGetStatusName - #define blasGetStatusString cublasGetStatusString - -using blasStatus_t = cublasStatus_t; -using blasHandle_t = cublasHandle_t; - -#elif KMM_USE_HIP - -// HIP backend -// Experimental draft, not working -using half_type = __half; -using bfloat16_type = __hip_bfloat16; - - #define GPU_DEVICE_ATTRIBUTE_MAX 0 - #define GPU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK hipDeviceAttributeMaxThreadsPerBlock - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X hipDeviceAttributeMaxBlockDimX - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y hipDeviceAttributeMaxBlockDimY - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z hipDeviceAttributeMaxBlockDimZ - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X hipDeviceAttributeMaxGridDimX - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y hipDeviceAttributeMaxGridDimY - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z hipDeviceAttributeMaxGridDimZ - #define GPU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR hipDeviceAttributeComputeCapabilityMajor - #define GPU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR hipDeviceAttributeComputeCapabilityMinor - #define GPU_MEMHOSTALLOC_PORTABLE hipHostMallocPortable - #define GPU_MEMHOSTALLOC_DEVICEMAP hipHostMallocMapped - #define GPU_SUCCESS hipSuccess - #define GPU_ERROR_OUT_OF_MEMORY hipErrorOutOfMemory - #define GPU_MEM_ALLOCATION_TYPE_PINNED hipMemAllocationTypePinned - #define GPU_MEM_HANDLE_TYPE_NONE hipMemHandleTypeNone - #define GPU_MEM_LOCATION_TYPE_DEVICE hipMemLocationTypeDevice - #define GPU_ERROR_UNKNOWN hipErrorUnknown - #define GPU_STREAM_NON_BLOCKING hipStreamNonBlocking - #define GPU_EVENT_WAIT_DEFAULT 0 - #define GPU_ERROR_NOT_READY hipErrorNotReady - #define GPU_EVENT_DISABLE_TIMING hipEventDisableTiming - #define GPU_MEMORYTYPE_HOST hipMemoryTypeHost - #define GPU_MEMORYTYPE_DEVICE hipMemoryTypeDevice - #define GPU_ERROR_NO_DEVICE hipErrorNoDevice - #define GPU_POINTER_ATTRIBUTE_DEVICE_ORDINAL HIP_POINTER_ATTRIBUTE_DEVICE_ORDINAL - #define GPU_CTX_MAP_HOST 0 - #define gpuCtxGetDevice hipCtxGetDevice - #define gpuDeviceGetName hipDeviceGetName - #define gpuDeviceGetAttribute hipDeviceGetAttribute - #define gpuMemGetInfo hipMemGetInfo - #define gpuMemcpyDtoHAsync hipMemcpyDtoHAsync - #define gpuStreamSynchronize hipStreamSynchronize - #define gpuMemsetD8Async hipMemsetD8Async - #define gpuMemsetD16Async hipMemsetD16Async - #define gpuMemsetD32Async hipMemsetD32Async - #define gpuMemHostAlloc hipHostMalloc - #define gpuMemFreeHost hipHostFree - #define gpuMemAlloc hipMalloc - #define gpuMemFree hipFree - #define gpuMemPoolCreate hipMemPoolCreate - #define gpuMemPoolDestroy hipMemPoolDestroy - #define gpuMemAllocFromPoolAsync hipMallocFromPoolAsync - #define gpuMemFreeAsync hipFreeAsync - #define gpuGetStreamPriorityRange hipDeviceGetStreamPriorityRange - #define gpuStreamCreateWithPriority hipStreamCreateWithPriority - #define gpuStreamQuery hipStreamQuery - #define gpuStreamDestroy hipStreamDestroy - #define gpuEventSynchronize hipEventSynchronize - #define gpuEventRecord hipEventRecord - #define gpuStreamWaitEvent hipStreamWaitEvent - #define gpuEventQuery hipEventQuery - #define gpuEventDestroy hipEventDestroy - #define gpuEventCreate hipEventCreateWithFlags - #define gpuMemcpy2DAsync hipMemcpyParam2DAsync - #define gpuMemcpy2D hipMemcpyParam2D - #define gpuMemcpyDtoH hipMemcpyDtoH - #define gpuMemcpyDtoDAsync hipMemcpyDtoDAsync - #define gpuMemcpyDtoD hipMemcpyDtoD - #define gpuGetErrorName hipDrvGetErrorName - #define gpuGetErrorString hipDrvGetErrorString - #define GPUrtGetErrorName hipGetErrorName - #define GPUrtGetErrorString hipGetErrorString - #define gpuGetLastError hipGetLastError - #define gpuInit hipInit - #define gpuDeviceGetCount hipGetDeviceCount - #define gpuDeviceGet hipDeviceGet - #define gpuCtxCreate hipCtxCreate - #define gpuCtxDestroy hipCtxDestroy - #define gpuDevicePrimaryCtxRetain hipDevicePrimaryCtxRetain - #define gpuDevicePrimaryCtxRelease hipDevicePrimaryCtxRelease - #define gpuCtxPushCurrent hipCtxPushCurrent - #define gpuCtxPopCurrent hipCtxPopCurrent - #define GPUrtLaunchKernel hipLaunchKernel - #define gpuMemPoolTrimTo hipMemPoolTrimTo - #define gpuDeviceGetDefaultMemPool hipDeviceGetDefaultMemPool - #define gpuPointerGetAttribute hipPointerGetAttribute - #define GPU_POINTER_ATTRIBUTE_MEMORY_TYPE HIP_POINTER_ATTRIBUTE_MEMORY_TYPE - -using GPUresult = hipError_t; -using gpuError_t = hipError_t; -using GPUdevice = hipDevice_t; -using GPUdevice_attribute = hipDeviceAttribute_t; -using GPUcontext = hipCtx_t; -using GPUmemorytype = hipMemoryType; -using GPUstream_t = hipStream_t; -using GPUdeviceptr = hipDeviceptr_t; -using GPUmemoryPool = hipMemPool_t; -using GPUmemPoolProps = hipMemPoolProps; -using GPUmemAllocationType = hipMemAllocationType; -using GPUmemAllocationHandleType = hipMemAllocationHandleType; -using GPUmemLocationType = hipMemLocationType; -using GPUevent_t = hipEvent_t; -using GPU_MEMCPY2D = hip_Memcpy2D; - - // cuBLAS - #define blasCreate rocblas_create_handle - #define blasSetStream rocblas_set_stream - #define blasDestroy rocblas_destroy_handle - #define blasGetStatusString rocblas_status_to_string - -using blasStatus_t = rocblas_status; -using blasHandle_t = rocblas_handle; - -const char* blasGetStatusName(blasStatus_t); -GPUresult gpuMemcpyAsync(GPUdeviceptr, GPUdeviceptr, size_t, GPUstream_t); -GPUresult gpuMemcpyHtoDAsync(GPUdeviceptr, const void*, size_t, GPUstream_t); -GPUresult gpuMemcpyHtoD(GPUdeviceptr, const void*, size_t); -GPUresult gpuMemcpyPeerAsync(GPUdeviceptr, GPUcontext, GPUdevice, GPUdeviceptr, GPUcontext, GPUdevice, size_t, GPUstream_t); - -#else - #define GPU_DEVICE_ATTRIBUTE_MAX 1 - #define GPU_MEMHOSTALLOC_PORTABLE 0 - #define GPU_MEMHOSTALLOC_DEVICEMAP 0 - #define GPU_SUCCESS 0 - #define GPU_ERROR_OUT_OF_MEMORY 0 - #define GPU_ERROR_UNKNOWN 0 - #define GPU_ERROR_NOT_READY 0 - #define GPU_EVENT_DISABLE_TIMING 2 - #define GPU_ERROR_NO_DEVICE 100 - #define GPU_CTX_MAP_HOST 0x08 - -using half_type = unsigned char; -using bfloat16_type = char; -using size_t = std::size_t; - -using GPUdevice = int; -class dim3 { - public: - dim3(unsigned int x = 1, unsigned int y = 1, unsigned int z = 1) : x(x), y(y), z(z) {} - - unsigned int x; - unsigned int y; - unsigned int z; -}; -enum GPUmemAllocationType { GPU_MEM_ALLOCATION_TYPE_PINNED = 1 }; -enum GPUmemAllocationHandleType { GPU_MEM_HANDLE_TYPE_NONE = 0 }; -enum GPUmemLocationType { GPU_MEM_LOCATION_TYPE_DEVICE = 1 }; -struct GPUmemLocation { - int id; - GPUmemLocationType type; -}; -struct GPUmemPoolProps { - GPUmemAllocationType allocType; - GPUmemAllocationHandleType handleTypes; - GPUmemLocation location; - size_t maxSize; - unsigned char reserved[54]; - unsigned short usage; - void* win32SecurityAttributes; -}; -using GPUcontext = int*; -using GPUstream_t = int*; -using GPUdeviceptr = unsigned long long; -using GPUmemoryPool = int*; -using GPUevent_t = void*; -enum GPUmemorytype { GPU_MEMORYTYPE_HOST, GPU_MEMORYTYPE_DEVICE }; -struct GPU_MEMCPY2D { - size_t Height; - size_t WidthInBytes; - GPUdeviceptr dstDevice; - void* dstHost; - GPUmemorytype dstMemoryType; - size_t dstPitch; - size_t dstXInBytes; - size_t dstY; - GPUdeviceptr srcDevice; - const void* srcHost; - GPUmemorytype srcMemoryType; - size_t srcPitch; - size_t srcXInBytes; - size_t srcY; -}; -enum GPUresult {}; -enum gpuError_t {}; -using GPUdevice_attribute = int; -using GPUstream_flags = int; -using GPUevent_wait_flags = int; - - #define GPU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK GPUdevice_attribute(1) - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_X GPUdevice_attribute(2) - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Y GPUdevice_attribute(3) - #define GPU_DEVICE_ATTRIBUTE_MAX_BLOCK_DIM_Z GPUdevice_attribute(4) - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X GPUdevice_attribute(5) - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Y GPUdevice_attribute(6) - #define GPU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_Z GPUdevice_attribute(7) - #define GPU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR GPUdevice_attribute(75) - #define GPU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR GPUdevice_attribute(76) - #define GPU_STREAM_NON_BLOCKING GPUstream_flags(1) - #define GPU_EVENT_WAIT_DEFAULT GPUevent_wait_flags(0) - -enum GPUpointer_attribute {}; - #define GPU_POINTER_ATTRIBUTE_MEMORY_TYPE GPUpointer_attribute(2) - #define GPU_POINTER_ATTRIBUTE_DEVICE_ORDINAL GPUpointer_attribute(9) - -// Dummy functions -GPUresult gpuCtxGetDevice(GPUdevice*); -GPUresult gpuDeviceGetName(char*, int, GPUdevice); -GPUresult gpuDeviceGetAttribute(int*, GPUdevice_attribute, GPUdevice); -GPUresult gpuMemGetInfo(size_t*, size_t*); -GPUresult gpuMemcpyHtoDAsync(GPUdeviceptr, const void*, size_t, GPUstream_t); -GPUresult gpuMemcpyDtoHAsync(void*, GPUdeviceptr, size_t, GPUstream_t); -GPUresult gpuMemcpyPeerAsync( - GPUdeviceptr, - GPUcontext, - GPUdevice, - GPUdeviceptr, - GPUcontext, - GPUdevice, - size_t, - GPUstream_t -); -GPUresult gpuStreamSynchronize(GPUstream_t); -GPUresult gpuMemsetD8Async(GPUdeviceptr, unsigned char, size_t, GPUstream_t); -GPUresult gpuMemsetD16Async(GPUdeviceptr, unsigned short, size_t, GPUstream_t); -GPUresult gpuMemsetD32Async(GPUdeviceptr, unsigned int, size_t, GPUstream_t); -GPUresult gpuMemcpyAsync(GPUdeviceptr, GPUdeviceptr, size_t, GPUstream_t); -GPUresult gpuMemHostAlloc(void**, size_t, unsigned int); -GPUresult gpuMemFreeHost(void*); -GPUresult gpuMemAlloc(GPUdeviceptr*, size_t); -GPUresult gpuMemFree(GPUdeviceptr); -GPUresult gpuMemPoolCreate(GPUmemoryPool*, const GPUmemPoolProps*); -GPUresult gpuMemPoolDestroy(GPUmemoryPool); -GPUresult gpuMemAllocFromPoolAsync(GPUdeviceptr*, size_t, GPUmemoryPool, GPUstream_t); -GPUresult gpuMemFreeAsync(GPUdeviceptr, GPUstream_t); -GPUresult gpuGetStreamPriorityRange(int*, int*); -GPUresult gpuStreamCreateWithPriority(GPUstream_t*, unsigned int, int); -GPUresult gpuStreamQuery(GPUstream_t); -GPUresult gpuStreamDestroy(GPUstream_t); -GPUresult gpuEventSynchronize(GPUevent_t); -GPUresult gpuEventRecord(GPUevent_t, GPUstream_t); -GPUresult gpuStreamWaitEvent(GPUstream_t, GPUevent_t, unsigned int); -GPUresult gpuEventQuery(GPUevent_t); -GPUresult gpuEventDestroy(GPUevent_t); -GPUresult gpuEventCreate(GPUevent_t, unsigned int); -GPUresult gpuMemcpyHtoD(GPUdeviceptr, const void*, size_t); -GPUresult gpuMemcpy2DAsync(const GPU_MEMCPY2D*, GPUstream_t); -GPUresult gpuMemcpy2D(const GPU_MEMCPY2D*); -GPUresult gpuMemcpyDtoH(void*, GPUdeviceptr, size_t); -GPUresult gpuMemcpyDtoDAsync(GPUdeviceptr, GPUdeviceptr, size_t, GPUstream_t); -GPUresult gpuMemcpyDtoD(GPUdeviceptr, GPUdeviceptr, size_t); -GPUresult gpuGetErrorName(GPUresult, const char**); -GPUresult gpuGetErrorString(GPUresult, const char**); -const char* GPUrtGetErrorName(gpuError_t); -const char* GPUrtGetErrorString(gpuError_t); -gpuError_t gpuGetLastError(void); -GPUresult gpuInit(unsigned int); -GPUresult gpuDeviceGetCount(int*); -GPUresult gpuDeviceGet(GPUdevice*, int); -GPUresult gpuPointerGetAttribute(void*, GPUpointer_attribute, GPUdeviceptr); -GPUresult gpuCtxCreate(GPUcontext*, unsigned int, GPUdevice); -GPUresult gpuCtxDestroy(GPUcontext); -GPUresult gpuDevicePrimaryCtxRetain(GPUcontext*, GPUdevice); -GPUresult gpuDevicePrimaryCtxRelease(GPUdevice); -GPUresult gpuCtxPushCurrent(GPUcontext); -GPUresult gpuCtxPopCurrent(GPUcontext*); -gpuError_t GPUrtLaunchKernel(const void*, dim3, dim3, void**, size_t, GPUstream_t); -GPUresult gpuMemPoolTrimTo(GPUmemoryPool, size_t); -GPUresult gpuDeviceGetDefaultMemPool(GPUmemoryPool*, GPUdevice); - -// Atomics -template -T atomicAnd(T* input, T output) { - return 0; -}; -template -T atomicOr(T* input, T output) { - return 0; -}; -template -T atomicAdd(T* input, T output) { - return 0; -}; -template -T atomicMin(T* input, T output) { - return 0; -}; -template -T atomicMax(T* input, T output) { - return 0; -}; - -// Dummy BLAS -using blasHandle_t = void*; -enum blasStatus_t {}; -blasStatus_t blasCreate(blasHandle_t); -blasStatus_t blasSetStream(blasHandle_t, GPUstream_t); -blasStatus_t blasDestroy(blasHandle_t); -const char* blasGetStatusName(blasStatus_t); -const char* blasGetStatusString(blasStatus_t); - -#endif - -} // namespace kmm \ No newline at end of file diff --git a/include/kmm/core/bounds.hpp b/include/kmm/core/bounds.hpp new file mode 100644 index 00000000..3d3ae797 --- /dev/null +++ b/include/kmm/core/bounds.hpp @@ -0,0 +1,300 @@ +#pragma once + +#include "kmm/core/point.hpp" +#include "kmm/core/range.hpp" +#include "kmm/core/shape.hpp" +#include "kmm/core/type_utils.hpp" +#include "kmm/core/vec.hpp" + +namespace kmm { + +/// \addtogroup geometry +/// @{ + +/// An N-dimensional axis-aligned box, given by a `Range` per axis, backed by a +/// `Vec, N>`. +/// +/// A `Bounds` describes a sub-region of an N-dimensional domain: the half-open `[begin, end)` +/// range of valid indices along each axis. It can be constructed from a begin/end point pair, +/// from an offset and a `Shape`, or directly from a `Shape` (assuming a zero origin). Use +/// `contains`/`overlaps`/`intersection` to test and combine bounds. +template +class Bounds: public Vec, N> { + public: + using storage_type = Vec, N>; + + KMM_HOST_DEVICE + explicit constexpr Bounds(const storage_type& storage) : storage_type(storage) {} + + KMM_HOST_DEVICE + constexpr Bounds() : storage_type(fill(Range())) {} + + constexpr Bounds(const Bounds&) = default; + constexpr Bounds(Bounds&&) noexcept = default; + Bounds& operator=(const Bounds&) = default; + Bounds& operator=(Bounds&&) noexcept = default; + + template> + KMM_HOST_DEVICE Bounds(Range first, Ts&&... args) : storage_type {first, args...} {} + + template + KMM_HOST_DEVICE constexpr Bounds(const Bounds& that) { + if (!that.template is_convertible_to()) { + throw_overflow_exception(); + } + + *this = Bounds::from(that); + } + + /// Construct from begin and end point. + KMM_HOST_DEVICE static constexpr Bounds from_bounds( + const Point& begin, + const Point& end + ) { + storage_type result; + + for (size_t i = 0; is_less(i, N); i++) { + result[i] = {begin[i], end[i]}; + } + + return Bounds(result); + } + + /// Construct from offset and shape. + KMM_HOST_DEVICE static constexpr Bounds from_offset_size( + const Point& offset, + const Shape& shape + ) { + storage_type result; + + for (size_t i = 0; is_less(i, N); i++) { + result[i] = Range(shape[i]) + offset[i]; + } + + return Bounds(result); + } + + KMM_HOST_DEVICE + constexpr Bounds(const Shape& shape) : + Bounds(from_offset_size(Point::zero(), shape)) {} + + /// Returns an empty bounds (all ranges are 0...0). + KMM_HOST_DEVICE static constexpr Bounds empty() { + return Bounds(fill(Range())); + } + + /// Returns `Bounds` with one element (all ranges are 0...1). + KMM_HOST_DEVICE static constexpr Bounds one() { + return Bounds(fill(Range::one())); + } + + /// Returns `Bounds` constructed from another `Bounds`. + template + KMM_HOST_DEVICE static constexpr Bounds from(const Vec, M>& that) { + storage_type result; + + for (size_t i = 0; is_less(i, N); i++) { + result[i] = is_less(i, M) ? Range::from(that[i]) : Range(static_cast(1)); + } + + return Bounds(result); + } + + /// Returns `true` if this `Bounds` can also be represented as `Bounds` without + /// loss of information, `false` otherwise. + template + KMM_HOST_DEVICE bool is_convertible_to() const { + bool result = true; + + for (size_t i = 0; is_less(i, N); i++) { + if (i < M) { + result &= (*this)[i].template is_convertible_to(); + } else { + result &= (*this)[i] == Range::one(); + } + } + + return result; + } + + /// Returns the range along the `i`-th dimension and `default_value` if it is out of bounds. + KMM_HOST_DEVICE + Range get_or_default(size_t i, Range default_value = Range::one()) const { + if constexpr (N > 0) { + if (KMM_LIKELY(i < N)) { + return (*this)[i]; + } + } + + return default_value; + } + + /// Returns the begin value along the `i`-th dimension. + KMM_HOST_DEVICE + T begin(size_t axis) const { + return get_or_default(axis).start; + } + + /// Returns the end value along the `i`-th dimension (exclusive). + KMM_HOST_DEVICE + T end(size_t axis) const { + return get_or_default(axis).stop; + } + + /// Returns the number of elements along the `i`-th dimension. + KMM_HOST_DEVICE + T size(size_t axis) const { + return get_or_default(axis).size(); + } + + /// Returns N-d begin point. + KMM_HOST_DEVICE + Point begin() const { + Point result; + for (size_t axis = 0; is_less(axis, N); axis++) { + result[axis] = (*this)[axis].start; + } + return result; + } + + /// Returns N-d end point (exclusive) + KMM_HOST_DEVICE + Point end() const { + Point result; + for (size_t axis = 0; is_less(axis, N); axis++) { + result[axis] = (*this)[axis].stop; + } + return result; + } + + /// Returns the N-d shape (i.e., size along each dimension). + KMM_HOST_DEVICE + Shape shape() const { + Shape result; + for (size_t axis = 0; is_less(axis, N); axis++) { + result[axis] = (*this)[axis].size(); + } + return result; + } + + /// Returns `true` only if this bounds is empty. + KMM_HOST_DEVICE + bool is_empty() const { + bool result = false; + + for (size_t i = 0; is_less(i, N); i++) { + result |= this->begin(i) >= this->end(i); + } + + return result; + } + + /// Returns the product of `size(0) * size(1) * ...`. + KMM_HOST_DEVICE + T volume() const { + T result = 1; + + for (size_t i = 0; is_less(i, N); i++) { + result *= this->end(i) - this->begin(i); + } + + return this->is_empty() ? T {0} : result; + } + + KMM_HOST_DEVICE + Bounds intersection(const Bounds& that) const { + storage_type result; + + for (size_t i = 0; is_less(i, N); i++) { + result[i].start = this->begin(i) >= that.begin(i) ? this->begin(i) : that.begin(i); + result[i].stop = this->end(i) <= that.end(i) ? this->end(i) : that.end(i); + } + + return Bounds(result); + } + + /// Returns `true` if this bounds overlaps the given bounds. + KMM_HOST_DEVICE + bool overlaps(const Bounds& that) const { + bool result = true; + + for (size_t i = 0; is_less(i, N); i++) { + result &= this->begin(i) < this->end(i) && that.begin(i) < that.end(i) && // + this->begin(i) < that.end(i) && that.begin(i) < this->end(i); + } + + return result; + } + + /// Returns `true` if this bounds contains the given bounds. + KMM_HOST_DEVICE + bool contains(const Bounds& that) const { + bool is_contained = true; + bool that_is_empty = false; + + for (size_t i = 0; is_less(i, N); i++) { + is_contained &= that.begin(i) >= this->begin(i); + is_contained &= that.end(i) <= this->end(i); + that_is_empty |= that.begin(i) >= that.end(i); + } + + return is_contained || that_is_empty; + } + + /// Returns `true` if this bounds contains the given point. + KMM_HOST_DEVICE + bool contains(const Point& that) const { + bool result = true; + + for (size_t i = 0; is_less(i, N); i++) { + result &= (*this)[i].contains(that[i]); + } + + return result; + } + + /// Returns `true` if this bounds contains the given point `{first, rest, ...}`. + template> + KMM_HOST_DEVICE bool contains(const T& first, Ts&&... rest) const { + return contains(Point {first, rest...}); + } + + /// Returns `true` if this bounds overlaps the given shape. + KMM_HOST_DEVICE + bool overlaps(const Shape& that) const { + return overlaps(Bounds {that}); + } + + /// Returns `true` if this bounds contains the given shape. + KMM_HOST_DEVICE + bool contains(const Shape& that) const { + return contains(Bounds {that}); + } +}; + +template +Bounds(Ts&&...) -> Bounds; + +/// Constructs a `Bounds` from the given ranges. +template +KMM_HOST_DEVICE Bounds bounds(const Ts&... values) { + return Bounds {Vec, sizeof...(Ts)> {values...}}; +} + +/// @} + +} // namespace kmm + +#if !KMM_IS_RTC + #include + + #include "fmt/ostream.h" + + #include "kmm/utils/hash_utils.hpp" + +template +struct fmt::formatter>: fmt::ostream_formatter {}; + +template +struct std::hash>: std::hash, N>> {}; +#endif \ No newline at end of file diff --git a/include/kmm/core/buffer.hpp b/include/kmm/core/buffer.hpp deleted file mode 100644 index 77b8e1b2..00000000 --- a/include/kmm/core/buffer.hpp +++ /dev/null @@ -1,79 +0,0 @@ -#pragma once - -#include -#include -#include - -#include "kmm/core/data_type.hpp" -#include "kmm/core/identifiers.hpp" - -namespace kmm { - -/** - * Represents the layout of a buffer. For now, this is just its size and alignment. - */ -struct BufferLayout { - BufferLayout repeat(size_t n) { - size_t remainder = size_in_bytes % alignment; - size_t padding = remainder != 0 ? alignment - remainder : 0; - return {(size_in_bytes + padding) * n, alignment}; - } - - template - static BufferLayout for_type(size_t n = 1) { - return BufferLayout {sizeof(T), alignof(T)}.repeat(n); - } - - static BufferLayout for_type(DataType dtype, size_t n = 1) { - return BufferLayout {dtype.size_in_bytes(), dtype.alignment()}.repeat(n); - } - - size_t size_in_bytes = 0; - size_t alignment = 1; -}; - -/** - * This enum is used to specify how a buffer can be accessed: read-only, read-write, or exclusive. - */ -enum struct AccessMode { - Read, ///< Read-only access to the buffer. - ReadWrite, ///< Read and write access to the buffer. - Exclusive ///< Exclusive access, implying full control over the buffer. -}; - -/** - * Represents the requirements for accessing a buffer. - */ -struct BufferRequirement { - BufferId buffer_id; - MemoryId memory_id; - AccessMode access_mode; -}; - -/** - * Provides access to a buffer with specific properties. - */ -struct BufferAccessor { - MemoryId memory_id; - BufferLayout layout; - bool is_writable; - void* address; -}; - -inline std::ostream& operator<<(std::ostream& f, AccessMode mode) { - switch (mode) { - case AccessMode::Read: - return f << "Read"; - case AccessMode::ReadWrite: - return f << "ReadWrite"; - case AccessMode::Exclusive: - return f << "Exclusive"; - } - - return f; -} - -} // namespace kmm - -template<> -struct fmt::formatter: fmt::ostream_formatter {}; \ No newline at end of file diff --git a/include/kmm/core/checked_compare.hpp b/include/kmm/core/checked_compare.hpp new file mode 100644 index 00000000..4b845369 --- /dev/null +++ b/include/kmm/core/checked_compare.hpp @@ -0,0 +1,570 @@ +#pragma once + +#include "kmm/core/macros.hpp" +#include "kmm/core/panic.hpp" + +namespace kmm { + +namespace detail { + +enum class numeric_type_tag { // + signed_int, + unsigned_int, + floating_point, + other +}; + +template +struct numeric_type_traits { + static constexpr numeric_type_tag tag = numeric_type_tag::other; +}; + +template<> +struct numeric_type_traits { + static constexpr numeric_type_tag tag = numeric_type_tag::floating_point; + + KMM_HOST_DEVICE + static constexpr bool isnan(float f) { + return f != f; + } + + KMM_HOST_DEVICE + static float ceil(float f) { + return ::ceilf(f); + } + + KMM_HOST_DEVICE + static float floor(float f) { + return ::floorf(f); + } +}; + +template<> +struct numeric_type_traits { + static constexpr numeric_type_tag tag = numeric_type_tag::floating_point; + + KMM_HOST_DEVICE + static constexpr bool isnan(double f) { + return f != f; + } + + KMM_HOST_DEVICE + static double ceil(double f) { + return ::ceil(f); + } + + KMM_HOST_DEVICE + static double floor(double f) { + return ::floor(f); + } +}; + +template<> +struct numeric_type_traits { + static constexpr bool is_signed = false; + using unsigned_type = bool; + + static constexpr numeric_type_tag tag = numeric_type_tag::unsigned_int; + static constexpr bool min_inclusive = false; + static constexpr bool max_inclusive = true; + + static constexpr float min_inclusive_float = 0.0F; + static constexpr float max_exclusive_float = 2.0F; +}; + +#define KMM_DEFINE_INT_TRAITS(T) \ + template<> \ + struct numeric_type_traits { \ + static constexpr bool is_signed = true; \ + static constexpr numeric_type_tag tag = numeric_type_tag::signed_int; \ + using unsigned_type = unsigned T; \ + \ + static constexpr signed T max_inclusive = \ + (signed T)(unsigned_type(~unsigned_type(0)) >> 1); \ + static constexpr signed T min_inclusive = ~max_inclusive; \ + \ + static constexpr float min_inclusive_float = min_inclusive; \ + static constexpr float max_exclusive_float = unsigned_type(max_inclusive) + 1; \ + }; \ + \ + template<> \ + struct numeric_type_traits { \ + static constexpr bool is_signed = false; \ + static constexpr numeric_type_tag tag = numeric_type_tag::unsigned_int; \ + \ + static constexpr unsigned T min_inclusive = 0; \ + static constexpr unsigned T max_inclusive = ~static_cast(0); \ + \ + static constexpr float min_inclusive_float = min_inclusive; \ + static constexpr float max_exclusive_float = 2.0f * float(max_inclusive / 2 + 1); \ + }; + +KMM_DEFINE_INT_TRAITS(char) +KMM_DEFINE_INT_TRAITS(short) +KMM_DEFINE_INT_TRAITS(int) +KMM_DEFINE_INT_TRAITS(long) +KMM_DEFINE_INT_TRAITS(long long) + +template +struct checked_compare_base { + KMM_HOST_DEVICE + static constexpr bool is_equal(const L& left, const R& right) { + return left == right; + } + + KMM_HOST_DEVICE + static constexpr bool is_less(const L& left, const R& right) { + return left < right; + } + + KMM_HOST_DEVICE + static constexpr bool is_less_equal(const L& left, const R& right) { + return left <= right; + } +}; + +template< + typename L, + typename R, + numeric_type_tag = numeric_type_traits::tag, + numeric_type_tag = numeric_type_traits::tag> +struct checked_compare_impl; + +template +struct checked_compare_impl: + checked_compare_base {}; + +template +struct checked_compare_impl: + checked_compare_base {}; + +template +struct checked_compare_impl: + checked_compare_base {}; + +template +struct checked_compare_impl { + using UR = typename numeric_type_traits::unsigned_type; + + KMM_HOST_DEVICE + static constexpr bool is_equal(const L& left, const R& right) { + return right >= static_cast(0) && left == static_cast(right); + } + + KMM_HOST_DEVICE + static constexpr bool is_less(const L& left, const R& right) { + return right >= static_cast(0) && left < static_cast(right); + } + + KMM_HOST_DEVICE + static constexpr bool is_less_equal(const L& left, const R& right) { + return right >= static_cast(0) && left <= static_cast(right); + } +}; + +template +struct checked_compare_impl { + using UL = typename numeric_type_traits::unsigned_type; + + KMM_HOST_DEVICE + static constexpr bool is_equal(const L& left, const R& right) { + return left >= static_cast(0) && static_cast