diff --git a/projects/rocshmem/.github/CODEOWNERS b/projects/rocshmem/.github/CODEOWNERS new file mode 100755 index 0000000000..f4985a580e --- /dev/null +++ b/projects/rocshmem/.github/CODEOWNERS @@ -0,0 +1 @@ +* @avinashkethineedi @akolliasAMD @Yiltan @BKP @abouteiller @edgargabriel @gaoikawa @omor1 diff --git a/projects/rocshmem/.gitignore b/projects/rocshmem/.gitignore new file mode 100644 index 0000000000..567609b123 --- /dev/null +++ b/projects/rocshmem/.gitignore @@ -0,0 +1 @@ +build/ diff --git a/projects/rocshmem/.readthedocs.yaml b/projects/rocshmem/.readthedocs.yaml new file mode 100644 index 0000000000..3afbabbc82 --- /dev/null +++ b/projects/rocshmem/.readthedocs.yaml @@ -0,0 +1,18 @@ +# Read the Docs configuration file +# See https://docs.readthedocs.io/en/stable/config-file/v2.html for details + +version: 2 + +sphinx: + configuration: docs/conf.py + +formats: [] + +python: + install: + - requirements: docs/sphinx/requirements.txt + +build: + os: ubuntu-22.04 + tools: + python: "3.10" diff --git a/projects/rocshmem/AUTHORS.md b/projects/rocshmem/AUTHORS.md new file mode 100644 index 0000000000..79be998a65 --- /dev/null +++ b/projects/rocshmem/AUTHORS.md @@ -0,0 +1,16 @@ +## This is the list of rocSHMEM's significant contributors. + +#### This does not necessarily list everyone who has contributed code, especially since many employees of one corporation may be contributing. To see the full list of contributors, see the revision history in source control. + +- Khaled Hamidouche +- Brandon Potter +- Michael LeBeane +- Rohit Zambre +- Kishore Punniyamurthy +- Ruchi Shah +- Muhammad A. Awad +- Edgar Gabriel +- Avinash Kethineedi +- Yiltan Temucin +- Aurelien Bouteiller +- Omri Mor diff --git a/projects/rocshmem/CHANGELOG.md b/projects/rocshmem/CHANGELOG.md new file mode 100644 index 0000000000..fbacad5b7b --- /dev/null +++ b/projects/rocshmem/CHANGELOG.md @@ -0,0 +1,96 @@ +# Changelog for rocSHMEM +## Unreleased - rocSHMEM 3.x.x for ROCm 7.x.x +### Added +* Added new APIs: + * `rocshmem_TYPENAME_alltoall_wg` + +## Unreleased -- rocSHMEM 3.2.1 for ROCm x.x.x +### Added +### Changed +### Removed +### Resolved issues +### Known issues + +## rocSHMEM 3.2.0 for ROCm 7.2.0 +### Added +* Added the GDA conduit for AMD Pensando IONIC +### Changed +* Dependency libraries are loaded dynamically +* The following APIs now have an implementation for the GDA conduit + * `rocshmem_p` + * fetching atomics `rochsmem__fetch_` + * collective APIs +* The following APIs now have an implementation for the IPC conduit + * `rocshmem__atomic_{and,or,xor,swap}` + * `rocshmem__atomic_fetch_{and,or,xor,swap}` +### Known issues +* Only 64bit rocSHMEM atomic APIs are implemented for the GDA conduit + +## rocSHMEM 3.1.0 for ROCm 7.1.1 +### Added +* Allow for IPC, RO, GDA backends to be selected at runtime +* Added the GDA conduit for different NIC vendors + * Broadcom BNXT\_RE (Thor 2) + * Mellanox MLX5 (IB and RoCE ConnectX-7) +* Added new APIs: + * `rocshmem_get_device_ctx` + * `rocshmem_ctx_pe_quiet` + * `rocshmem_pe_quiet` + +### Changed +* The following APIs have been deprecated: + * `rocshmem_wg_init` + * `rocshmem_wg_finalize` + * `rocshmem_wg_init_thread` +* `rocshmem_ptr` can now return non-null pointer to + a shared memory region when the IPC transport is available to reach that region. + Previously, it would return a null pointer. +* `ROCSHMEM_RO_DISABLE_IPC` was renamed to `ROCSHMEM_DISABLE_MIXED_IPC`. + This enviroment variable was not documented for prior releases. + It is now documented to inform users who were using this undocumented feature. + +### Removed +* rocSHMEM no-longer requires rocPRIM and rocThrust as dependencies +* Removed MPI compile-time dependency + +### Known issues +* Only a subset of rocSHMEM APIs are implemented for the GDA conduit + +## rocSHMEM 3.0.0 for ROCm 7.0.0 +### Added + +* Added the Reverse Offload conduit +* Added new APIs: + * `rocshmem_ctx_barrier` + * `rocshmem_ctx_barrier_wave` + * `rocshmem_ctx_barrier_wg` + * `rocshmem_barrier_all` + * `rocshmem_barrier_all_wave` + * `rocshmem_barrier_all_wg` + * `rocshmem_ctx_sync` + * `rocshmem_ctx_sync_wave` + * `rocshmem_ctx_sync_wg` + * `rocshmem_sync_all` + * `rocshmem_sync_all_wave` + * `rocshmem_sync_all_wg` + * `rocshmem_init_attr` + * `rocshmem_get_uniqueid` + * `rocshmem_set_attr_uniqueid_args` +* Added dlmalloc based allocator +* Added XNACK support +* Added support for initialization with MPI communicators other than `MPI_COMM_WORLD` + +### Changed + +* Changed collective APIs to use `_wg` suffix rather than `_wg_` infix + +### Resolved Issues +* Resolved segfault in `rocshmem_wg_ctx_create`, now provides nullptr if ctx cannot be created + +## rocSHMEM 2.0.1 for ROCm 6.4.2 + +### Resolved Issues + +* Resolved incorrect output for `rocshmem_ctx_my_pe` and `rocshmem_ctx_n_pes` +* Resolved multi-team errors by providing team specific buffers in `rocshmem_ctx_wg_team_sync` +* Resolved missing implementation of `rocshmem_g` for IPC conduit diff --git a/projects/rocshmem/CMakeLists.txt b/projects/rocshmem/CMakeLists.txt new file mode 100644 index 0000000000..20becb009c --- /dev/null +++ b/projects/rocshmem/CMakeLists.txt @@ -0,0 +1,278 @@ +############################################################################### +# Copyright (c) Advanced Micro Devices, Inc. All rights reserved. +# +# SPDX-License-Identifier: MIT +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to +# deal in the Software without restriction, including without limitation the +# rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +# sell copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +# FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS +# IN THE SOFTWARE. +############################################################################### + +cmake_minimum_required(VERSION 3.16.3 FATAL_ERROR) + +############################################################################### +# AVOID IN SOURCE BUILD +############################################################################### +if(CMAKE_SOURCE_DIR STREQUAL CMAKE_BINARY_DIR AND + CMAKE_CURRENT_SOURCE_DIR STREQUAL CMAKE_SOURCE_DIR) + set(MSG "") + message(STATUS "Warning! Building from the source directory is not recommended") + message(STATUS "If unintended, please remove 'CMakeCache.txt' and 'CMakeFiles'") + message(STATUS "and build from a separate directory") + message(FATAL_ERROR "In-source build") +endif() + +############################################################################### +# CONFIGURATION OPTIONS +############################################################################### +option(DEBUG "Enable debug trace" OFF) +option(PROFILE "Enable statistics and timing support" OFF) +option(USE_RO "Enable RO conduit" ON) +option(USE_IPC "Enable IPC support (using HIP)" OFF) +option(USE_GDA "Enable GDA conduit" OFF) +option(USE_THREADS "Enable workgroup threads to share network queues" OFF) +option(USE_WF_COAL "Enable wavefront message coalescing" OFF) +option(USE_HEAP_DEVICE_FINEGRAIN "Heap uses GPU memory in finegrain mode" ON) +option(USE_HEAP_DEVICE_UNCACHED "Heap uses GPU memory in uncached mode" OFF) +option(USE_HEAP_DEVICE_COARSEGRAIN "Heap uses GPU memory in coarsegrain mode" OFF) +option(USE_HEAP_MANAGED "Heap uses managed memory" OFF) +option(USE_HEAP_HOST_HIP "Heap uses pinned host memory allocated with hip api" OFF) +option(USE_HEAP_HOST "Heap uses host memory allocated with malloc/free" OFF) +option(USE_ALLOC_DLMALLOC "Enable dlmalloc device memory allocator" ON) +option(USE_ALLOC_POW2BINS "Enable legacy Pow2Bins device memory allocator" OFF) +option(USE_FUNC_CALL "Force compiler to use function calls on library API" OFF) +option(USE_SHARED_CTX "Request support for shared ctx between WG" OFF) +option(USE_SINGLE_NODE "Enable single node support only." OFF) +option(USE_HDP_FLUSH "Force flush the HDP cache." OFF) +option(USE_HDP_FLUSH_HOST_SIDE "Use a polling thread to flush the HDP cache on the host." OFF) + +option(BUILD_FUNCTIONAL_TESTS "Build the functional tests (Requires MPI)" OFF) +option(BUILD_EXAMPLES "Build the examples" ON) +option(BUILD_UNIT_TESTS "Build the unit tests (Requires MPI)" OFF) +option(BUILD_TESTS_ONLY "Build only tests. Used to link agains rocSHMEM in a ROCm Release" OFF) +option(BUILD_TOOLS "Build binary tools (e.g., rocshmem_info)" ON) + +option(BUILD_LOCAL_GPU_TARGET_ONLY "Build only for GPUs detected on this machine" OFF) +option(BUILD_CODE_COVERAGE "Build with code coverage flags (gcc only)" OFF) + +option(GDA_IONIC "Build for AMD Pensando IONIC RDMA provider" OFF) +option(GDA_BNXT "Build for Broadcom RDMA provider" OFF) +option(GDA_MLX5 "Build for Mellanox MLX5 RDMA provider" OFF) + +set(USE_EXTERNAL_MPI AUTO CACHE STRING "Link with an external MPI (required if used MPI is ABI incompatible with Open MPI v5)") +set_property(CACHE USE_EXTERNAL_MPI PROPERTY STRINGS AUTO ON OFF) + +############################################################################### +# PROJECT +############################################################################### +include(${CMAKE_SOURCE_DIR}/cmake/setup_project.cmake) + +## Setup VERSION +file(READ include/rocshmem/rocshmem.hpp header_text) +if("${header_text}" MATCHES "constexpr char VERSION\\[\\] *= \"([0-9]+)\\.([0-9]+)\\.([0-9]+)\";") + set(VERSION_STRING ${CMAKE_MATCH_1}.${CMAKE_MATCH_2}.${CMAKE_MATCH_3}) +else() + message(FATAL_ERROR "Failed to parse Version") +endif() +message(STATUS "rocSHMEM Version: " "${VERSION_STRING}") + +project(rocshmem VERSION ${VERSION_STRING} LANGUAGES CXX) + +find_package(ROCmCMakeBuildTools PATHS /opt/rocm) +include(ROCMCreatePackage) +include(ROCMInstallTargets) +include(ROCMCheckTargetIds) + +rocm_setup_version(VERSION ${VERSION_STRING}) + + +############################################################################# +# SET GPU ARCHITECTURES +############################################################################# +include(cmake/rocm_local_targets.cmake) + +set(DEFAULT_GPUS + gfx90a:xnack-; + gfx90a:xnack+; + gfx1100; + gfx1201; + gfx942) + +if(${ROCM_MAJOR_VERSION} GREATER 6) + list(APPEND DEFAULT_GPUS gfx950) +endif() + +if($ENV{BUILD_LOCAL_GPU_TARGET_ONLY}) + set(BUILD_LOCAL_GPU_TARGET_ONLY ON) +endif() + +if (BUILD_LOCAL_GPU_TARGET_ONLY) + message(STATUS "Building only for local GPU target") + if (COMMAND rocm_local_targets) + rocm_local_targets(DEFAULT_GPUS) + else() + message(WARNING "Unable to determine local GPU targets. Falling back to default GPUs.") + endif() +endif() + +set(DEFAULT_GPU_TARGETS "${DEFAULT_GPUS}" CACHE STRING + "Target default GPUs if GPU_TARGETS is not defined.") + +if (COMMAND rocm_check_target_ids) + message(STATUS "Checking for ROCm support for GPU targets: " "${DEFAULT_GPU_TARGETS}") + rocm_check_target_ids(SUPPORTED_GPUS TARGETS ${DEFAULT_GPU_TARGETS}) +else() + message(WARNING "Unable to check for supported GPU targets.") + set(SUPPORTED_GPUS ${DEFAULT_GPU_TARGETS}) +endif() + +set(GPU_TARGETS "${SUPPORTED_GPUS}" CACHE STRING "GPU architectures to compile for") + +message(STATUS "Compiling for ${GPU_TARGETS}") + +############################################################################### +# CREATE ROCSHMEM LIBRARY +############################################################################### +if (NOT BUILD_TESTS_ONLY) + add_library(${PROJECT_NAME}) + add_library(roc::${PROJECT_NAME} ALIAS ${PROJECT_NAME}) + add_subdirectory(src) + + ############################################################################# + # PACKAGE DEPENDENCIES + ############################################################################# + if (NOT USE_EXTERNAL_MPI STREQUAL "OFF") + find_package(MPI) + else() + message ("-- External MPI detection disabled by user") + endif() + + if (MPI_FOUND) + set(HAVE_EXTERNAL_MPI ON) + else() + set(HAVE_EXTERNAL_MPI OFF) + set(BUILD_UNIT_TESTS OFF) + endif() + + if (USE_EXTERNAL_MPI STREQUAL "ON") + if(NOT HAVE_EXTERNAL_MPI) + message(FATAL_ERROR "External MPI support requested but MPI support not found. Build Aborted") + endif() + endif() + + find_package(hip REQUIRED PATHS /opt/rocm) + find_package(hsa-runtime64 REQUIRED) + + set(CMAKE_THREAD_PREFER_PTHREAD TRUE) + set(THREADS_PREFER_PTHREAD_FLAG TRUE) + find_package(Threads REQUIRED) + + configure_file(cmake/rocshmem_config.h.in include/rocshmem/rocshmem_config.h) + + ############################################################################# + # LINKING AND INCLUDE DIRECTORIES + ############################################################################# + target_compile_options( + ${PROJECT_NAME} + PUBLIC + -fgpu-rdc + ) + + target_include_directories( + ${PROJECT_NAME} + PUBLIC + $ + $ # rocshmem_config.h + $ # rocshmem_config.h from rocshmem.hpp + $ + $ + ) + + target_link_libraries( + ${PROJECT_NAME} + PUBLIC + $<$:MPI::MPI_CXX> + Threads::Threads + hip::device + hip::host + dl + hsa-runtime64::hsa-runtime64 + -fgpu-rdc + ) + + if(${ROCM_MAJOR_VERSION} LESS 7) + # ROCm 6.x requires us to explicitly enable warp sync builtins + target_compile_definitions(${PROJECT_NAME} PRIVATE HIP_ENABLE_WARP_SYNC_BUILTINS=1) + endif() + + ############################################################################# + # INSTALL + ############################################################################# + include(ROCMInstallTargets) + include(ROCMCreatePackage) + + rocm_install(TARGETS rocshmem) + + rocm_install( + DIRECTORY ${CMAKE_SOURCE_DIR}/include/ + DESTINATION ${CMAKE_INSTALL_INCLUDEDIR} + ) + + rocm_install( + FILES "${CMAKE_BINARY_DIR}/include/rocshmem/rocshmem_config.h" + DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/rocshmem + ) + + if (BUILD_TOOLS) + rocm_install( + PROGRAMS "${CMAKE_BINARY_DIR}/src/tools/rocshmem_info" + DESTINATION ${CMAKE_INSTALL_BINDIR} + ) + endif() + + rocm_package_add_dependencies( + DEPENDS + hsa-rocr + hip-runtime-amd + rocm-dev + ) + + rocm_export_targets( + TARGETS roc::rocshmem + NAMESPACE roc:: + ) + include(ROCMPackageConfigHelpers) + include(ROCMClients) + rocm_package_setup_component(clients) + rocm_package_setup_client_component(tests PACKAGE_NAME tests) + + rocm_create_package( + NAME "rocSHMEM" + DESCRIPTION "ROCm OpenSHMEM (rocSHMEM)" + MAINTAINER "rocSHMEM Maintainer " + ) +endif (NOT BUILD_TESTS_ONLY) + +############################################################################### +# TEST SUBDIRECTORIES +############################################################################### +add_subdirectory(tests) + +if (BUILD_EXAMPLES) + add_subdirectory(examples) +endif() + diff --git a/projects/rocshmem/CONTRIBUTING.md b/projects/rocshmem/CONTRIBUTING.md new file mode 100644 index 0000000000..b6a44d858e --- /dev/null +++ b/projects/rocshmem/CONTRIBUTING.md @@ -0,0 +1,59 @@ +## How to fork from us + +To keep our development fast and conflict free, we recommend you to [fork](https://github.com/ROCm/rocSHMEM/fork) our repository and start your work from our `develop` branch in your private repository. + +Afterwards, git clone your repository to your local machine. But that is not it! To keep track of the original develop repository, add it as another remote. + +``` +git remote add mainline https://github.com/ROCm/rocSHMEM.git +git checkout dev +``` + +As always in git, start a new branch with + +``` +git checkout -b topic- +``` + +and apply your changes there. + +## How to contribute to rocSHMEM + +### Did you find a bug? + +- Ensure the bug was not already reported by searching on GitHub under [Issues](https://github.com/ROCm/rocSHMEM/issues). + +- If you're unable to find an open issue addressing the problem, [open a new one](https://github.com/ROCm/rocSHMEM/issues/new). + +### Did you write a patch that fixes a bug? + +- Open a new GitHub [pull request](https://github.com/ROCm/rocSHMEM/compare) with the patch. + +- Ensure the PR description clearly describes the problem and solution. If there is an existing GitHub issue open describing this bug, please include it in the description so we can close it. + +- Ensure the PR is based on the `dev` branch of the rocSHMEM GitHub repository. + +- rocSHMEM requires new commits to include a "Signed-off-by" token in the commit message (typically enabled via the `git commit -s` option), indicating your agreement to the projects's [Developer's Certificate of Origin](https://developercertificate.org/) and compatability with the project [LICENSE](https://github.com/ROCm/rocSHMEM/blob/main/LICENSE): + + +> (a) The contribution was created in whole or in part by me and I +> have the right to submit it under the open source license +> indicated in the file; or +> +> (b) The contribution is based upon previous work that, to the best +> of my knowledge, is covered under an appropriate open source +> license and I have the right under that license to submit that +> work with modifications, whether created in whole or in part +> by me, under the same open source license (unless I am +> permitted to submit under a different license), as indicated +> in the file; or +> +> (c) The contribution was provided directly to me by some other +> person who certified (a), (b) or (c) and I have not modified +> it. +> +> (d) I understand and agree that this project and the contribution +> are public and that a record of the contribution (including all +> personal information I submit with it, including my sign-off) is +> maintained indefinitely and may be redistributed consistent with +> this project or the open source license(s) involved. diff --git a/projects/rocshmem/LICENSE.md b/projects/rocshmem/LICENSE.md new file mode 100644 index 0000000000..4ae4bb09d7 --- /dev/null +++ b/projects/rocshmem/LICENSE.md @@ -0,0 +1,23 @@ +MIT License + +Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + +SPDX-License-Identifier: MIT + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in all +copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +SOFTWARE. diff --git a/projects/rocshmem/README.md b/projects/rocshmem/README.md new file mode 100644 index 0000000000..56e1d39d12 --- /dev/null +++ b/projects/rocshmem/README.md @@ -0,0 +1,33 @@ +# ROCm OpenSHMEM (rocSHMEM) + +The ROCm OpenSHMEM (rocSHMEM) runtime is part of an AMD and AMD Research +initiative to provide GPU-centric networking through an OpenSHMEM-like interface. +This intra-kernel networking library simplifies application +code complexity and enables more fine-grained communication/computation +overlap than traditional host-driven networking. +rocSHMEM uses a single symmetric heap that is allocated on GPU memories. + +There are currently three backends for rocSHMEM; +IPC, Reverse Offload (RO), and GDA. +The backends primarily differ in their implementations of intra-kernel networking. + +The IPC backend implements communication primitives using load/store operations issued from the GPU. + +The Reverse Offload (RO) backend has the GPU runtime forward rocSHMEM networking operations +to the host-side runtime, which calls into a traditional MPI or OpenSHMEM +implementation. This forwarding of requests is transparent to the +programmer, who only sees the GPU-side interface. + +The GPU Direct Async (GDA) backend allows for rocSHMEM to issue communication operations to the NIC directly from the device-side code, without involving a CPU proxy. +within the GPU. +During initialization we prepare network resources for each NIC vendor using the vendor-appropriate +Direct Verbs APIs. +When calling the device-side rocSHMEM API, the device threads are used to construct Work Queue Entries (WQEs) and post the communication to the send queues of the NIC directly. +Completion Queues (CQs) are polled from the device-side code as well. + +The RO and GDA backend is provided as-is with limited support from AMD or AMD Research. + +## Installation and using rocSHMEM + +For information on how to install and use rocSHMEM, +[please see our documentation](https://rocm.docs.amd.com/projects/rocSHMEM/en/latest/). diff --git a/projects/rocshmem/cmake/FindPMIx.cmake b/projects/rocshmem/cmake/FindPMIx.cmake new file mode 100644 index 0000000000..2fcac1da96 --- /dev/null +++ b/projects/rocshmem/cmake/FindPMIx.cmake @@ -0,0 +1,68 @@ +############################################################################### +# Copyright (c) Advanced Micro Devices, Inc. All rights reserved. +# +# SPDX-License-Identifier: MIT +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to +# deal in the Software without restriction, including without limitation the +# rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +# sell copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +# FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS +# IN THE SOFTWARE. +############################################################################### + +# Find pmix installation. +# Different scenarios need to be covered: +# - pmix at user-provided location (i.e., in PMIX_ROOT) +# - pmix installed as part of Open MPI, i.e., in the MPI installation directories +# - pmix deployed with linux distros, Slurm, etc. + +find_package(PkgConfig QUIET) +if (PkgConfig_FOUND) + # Figure out and prepend the install dir for MPI + string(REGEX REPLACE "/include$" "" mpi_dir "${MPI_CXX_HEADER_DIR}") + foreach (mpiroot "${MPI_ROOT}" "$ENV{MPI_ROOT}" "${mpi_dir}") + if (mpiroot) + set(ENV{PKG_CONFIG_PATH} "${mpiroot}/lib/pkgconfig:$ENV{PKG_CONFIG_PATH}") + endif() + endforeach() + # prepend PMIX_ROOT + foreach (pmixroot "${PMIX_ROOT}" "$ENV{PMIX_ROOT}" "${PMIx_ROOT}" "$ENV{PMIx_ROOT}") + if (pmixroot) + set(ENV{PKG_CONFIG_PATH} "${pmixroot}/lib/pkgconfig:$ENV{PKG_CONFIG_PATH}") + endif() + endforeach() + pkg_check_modules(PC_PMIX QUIET pmix) +endif() + +find_path(PMIX_INCLUDE_DIR pmix.h + HINTS ${PC_PMIX_INCLUDE_DIRS} ${MPI_CXX_HEADER_DIR} ${MPI_ROOT} $ENV{MPI_ROOT} + PATH_SUFFIXES include) +if (PMIX_INCLUDE_DIR) + string(REGEX REPLACE "/include$" "" pmix_dir ${PMIX_INCLUDE_DIR}) + find_library(PMIX_LIBRARY pmix PATHS ${pmix_dir} PATH_SUFFIXES lib lib64 NO_DEFAULT_PATH) +endif() + +find_package_handle_standard_args(PMIx DEFAULT_MSG + PMIX_LIBRARY PMIX_INCLUDE_DIR) +mark_as_advanced(PMIX_LIBRARY PMIX_INCLUDE_DIR) + +if (PMIx_FOUND) +add_library(PMIx::pmix UNKNOWN IMPORTED) +set_target_properties(PMIx::pmix PROPERTIES + IMPORTED_LOCATION "${PMIX_LIBRARY}" + INTERFACE_COMPILE_OPTIONS "${PC_PMIX_CFLAGS_OTHER}" + INTERFACE_INCLUDE_DIRECTORIES "${PMIX_INCLUDE_DIR}" +) +endif() diff --git a/projects/rocshmem/cmake/rocm_local_targets.cmake b/projects/rocshmem/cmake/rocm_local_targets.cmake new file mode 100644 index 0000000000..ccad77351f --- /dev/null +++ b/projects/rocshmem/cmake/rocm_local_targets.cmake @@ -0,0 +1,52 @@ +############################################################################### +# Copyright (c) Advanced Micro Devices, Inc. All rights reserved. +# +# SPDX-License-Identifier: MIT +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to +# deal in the Software without restriction, including without limitation the +# rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +# sell copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +# FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS +# IN THE SOFTWARE. +############################################################################### + +# Find available local ROCM targets +# NOTE: This will eventually be part of ROCm-CMake and should be removed at that time +function(rocm_local_targets VARIABLE) + set(${VARIABLE} "NOTFOUND" PARENT_SCOPE) + find_program(_rocm_agent_enumerator rocm_agent_enumerator HINTS /opt/rocm/bin ENV ROCM_PATH) + if(NOT _rocm_agent_enumerator STREQUAL "_rocm_agent_enumerator-NOTFOUND") + execute_process( + COMMAND "${_rocm_agent_enumerator}" + RESULT_VARIABLE _found_agents + OUTPUT_VARIABLE _rocm_agents + ERROR_QUIET + ) + if (_found_agents EQUAL 0) + string(REPLACE "\n" ";" _rocm_agents "${_rocm_agents}") + unset(result) + foreach (agent IN LISTS _rocm_agents) + if (NOT agent STREQUAL "gfx000") + list(APPEND result "${agent}") + endif() + endforeach() + if(result) + list(REMOVE_DUPLICATES result) + set(${VARIABLE} "${result}" PARENT_SCOPE) + endif() + endif() + endif() +endfunction() + diff --git a/projects/rocshmem/cmake/rocshmem_config.h.in b/projects/rocshmem/cmake/rocshmem_config.h.in new file mode 100644 index 0000000000..b2e18664c3 --- /dev/null +++ b/projects/rocshmem/cmake/rocshmem_config.h.in @@ -0,0 +1,48 @@ +/****************************************************************************** + * Copyright (c) Advanced Micro Devices, Inc. All rights reserved. + * + * SPDX-License-Identifier: MIT + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to + * deal in the Software without restriction, including without limitation the + * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or + * sell copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS + * IN THE SOFTWARE. + *****************************************************************************/ + +#cmakedefine DEBUG +#cmakedefine PROFILE +#cmakedefine USE_RO +#cmakedefine USE_IPC +#cmakedefine USE_GDA +#cmakedefine USE_THREADS +#cmakedefine USE_SHARED_CTX +#cmakedefine USE_WF_COAL +#cmakedefine USE_HEAP_DEVICE_FINEGRAIN +#cmakedefine USE_HEAP_DEVICE_UNCACHED +#cmakedefine USE_HEAP_DEVICE_COARSEGRAIN +#cmakedefine USE_HEAP_MANAGED +#cmakedefine USE_HEAP_HOST_HIP +#cmakedefine USE_HEAP_HOST +#cmakedefine USE_ALLOC_DLMALLOC +#cmakedefine USE_ALLOC_POW2BINS +#cmakedefine USE_FUNC_CALL +#cmakedefine USE_SINGLE_NODE +#cmakedefine USE_HDP_FLUSH +#cmakedefine USE_HDP_FLUSH_HOST_SIDE +#cmakedefine GDA_IONIC +#cmakedefine GDA_BNXT +#cmakedefine GDA_MLX5 +#cmakedefine HAVE_EXTERNAL_MPI diff --git a/projects/rocshmem/cmake/setup_project.cmake b/projects/rocshmem/cmake/setup_project.cmake new file mode 100644 index 0000000000..4bf480b4ae --- /dev/null +++ b/projects/rocshmem/cmake/setup_project.cmake @@ -0,0 +1,82 @@ +############################################################################### +# Copyright (c) Advanced Micro Devices, Inc. All rights reserved. +# +# SPDX-License-Identifier: MIT +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to +# deal in the Software without restriction, including without limitation the +# rights to use, copy, modify, merge, publish, distribute, sublicense, and/or +# sell copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in +# all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +# FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS +# IN THE SOFTWARE. +############################################################################### + +############################################################################### +# DEFAULT BUILD TYPE +############################################################################### +set(CMAKE_BUILD_TYPE "Release" CACHE STRING + "build type: Release, Debug, RelWithDebInfo, MinSizeRel") + +############################################################################### +# DEPENDENCIES +############################################################################### + +# Try to establish ROCM_PATH (for find_package) +#================================================================================================== +if(NOT DEFINED ROCM_PATH) + # Guess default location + set(ROCM_PATH "/opt/rocm") + message(WARNING "Unable to find ROCM_PATH: Falling back to ${ROCM_PATH}") +else() + message(STATUS "ROCM_PATH found: ${ROCM_PATH}") +endif() +set(ENV{ROCM_PATH} ${ROCM_PATH}) + +## Check for ROCm version + +if(ROCM_PATH) + message(STATUS "Reading ROCM version from ${ROCM_PATH}/.info/version") + file(READ "${ROCM_PATH}/.info/version" rocm_version_string) +else() + message(FATAL_ERROR "Could not determine ROCM version (set EXPLICIT_ROCM_VERSION or set ROCM_PATH to a valid installation)") +endif() +string(REGEX MATCH "([0-9]+)\\.([0-9]+)\\.([0-9]+)" rocm_version_matches ${rocm_version_string}) +if (rocm_version_matches) + set(ROCM_MAJOR_VERSION ${CMAKE_MATCH_1}) + set(ROCM_MINOR_VERSION ${CMAKE_MATCH_2}) + set(ROCM_PATCH_VERSION ${CMAKE_MATCH_3}) + + message(STATUS "ROCm version: ${ROCM_MAJOR_VERSION}.${ROCM_MINOR_VERSION}.${ROCM_PATCH_VERSION}") +else() + message(WARNING "Failed to extract ROCm version.") +endif() + +foreach (root ${hip_ROOT} $ENV{hip_ROOT} ${ROCM_ROOT} $ENV{ROCM_ROOT} ${ROCM_PATH} $ENV{ROCM_PATH}) + if (IS_DIRECTORY ${root}) + list(PREPEND CMAKE_PREFIX_PATH ${root}) + endif() +endforeach() +if (NOT DEFINED CMAKE_CXX_COMPILER) + find_program(CMAKE_CXX_COMPILER hipcc PATHS /opt/rocm) +endif() + +############################################################################### +# GLOBAL COMPILE FLAGS +############################################################################### +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD 20) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CXX_FLAGS_DEBUG "-O0 -ggdb") + +list(APPEND CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/cmake) diff --git a/projects/rocshmem/docs/.gitignore b/projects/rocshmem/docs/.gitignore new file mode 100644 index 0000000000..7485356045 --- /dev/null +++ b/projects/rocshmem/docs/.gitignore @@ -0,0 +1,5 @@ +_build/ +_doxygen/ +doxygen/html/ +doxygen/xml/ +sphinx/_toc.yml diff --git a/projects/rocshmem/docs/README.md b/projects/rocshmem/docs/README.md new file mode 100644 index 0000000000..8aa5f29a2e --- /dev/null +++ b/projects/rocshmem/docs/README.md @@ -0,0 +1,21 @@ +# Building the rocSHMEM documentation + +## macOS + +To build html documentation locally: + +``` +brew install doxygen sphinx-doc +pip3.10 install -r ./sphinx/requirements.txt +python3.10 -m sphinx -T -E -b html -d _build/doctrees -D language=en . _build/html +open _build/html/index.html +``` + +To build pdf documentation we require a LaTeX installation on your machine. +Once LaTeX is installed, you may run the following: + +``` +pip3.10 install -r ./sphinx/requirements.txt +sphinx-build -M latexpdf . _build +open _build/latex/rocshmem.pdf +``` diff --git a/projects/rocshmem/docs/api/amo.rst b/projects/rocshmem/docs/api/amo.rst new file mode 100644 index 0000000000..00e6540c41 --- /dev/null +++ b/projects/rocshmem/docs/api/amo.rst @@ -0,0 +1,418 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-amo: + +--------------------------- +Atomic memory operations +--------------------------- + +You can call these functions from divergent control paths at the per-thread level. + +ROSHMEM_ATOMIC_FETCH +-------------------- +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_fetch(TYPE *source, int pe) +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch(rocshmem_ctx_t ctx, TYPE *source, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param pe: PE of the remote process. + + :returns: The value of ``dest``. + +**Description:** +This function atomically returns the value of ``dest`` to the calling PE. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in EXTENDED_AMO_TYPES_. + + +SHMEM_ATOMIC_SET +---------------- +.. cpp:function:: __device__ void rocshmem_TYPENAME_atomic_set(TYPE *dest, TYPE value, int pe); +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_atomic_set(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, int pe); + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically set. + :param pe: PE of the remote process. + + :returns: None. + +**Description:** +This function atomically sets the value ``value`` to ``dest`` on ``pe``. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in EXTENDED_AMO_TYPES_. + +SHMEM_ATOMIC_COMPARE_SWAP +------------------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_compare_swap(TYPE *dest, TYPE cond, TYPE value, TYPE pe); +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_compare_swap(rocshmem_ctx_t ctx, TYPE *dest, TYPE cond, TYPE value, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param cond: The value to be compare with. + :param value: The value to be atomically swapped. + :param pe: PE of the remote process. + + :return: The old value of ``dest``. + +**Description:** +This function atomically compares the value in ``dest`` with ``cond``. If they are equal, it stores ``value`` in ``dest``. +The operation returns the older value of ``dest`` to the calling PE. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in STANDARD_AMO_TYPES_. + +SHMEM_ATOMIC_SWAP +----------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_swap(TYPE *dest, TYPE value, TYPE pe); +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_swap(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically swapped. + :param pe: PE of the remote process. + + :return: The old value of ``dest``. + +**Description:** +This function atomically swaps the value ``val`` with ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in EXTENDED_AMO_TYPES_. + +SHMEM_ATOMIC_FETCH_INC +---------------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_fetch_inc(TYPE *dest, TYPE pe); +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_inc(rocshmem_ctx_t ctx, TYPE *dest, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param pe: PE of the remote process. + + :return: The old value of ``dest``. + +**Description:** +This function atomically adds ``1`` to ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in STANDARD_AMO_TYPES_. + +SHMEM_ATOMIC_INC +---------------- + +.. cpp:function:: __device__ void rocshmem_TYPENAME_atomic_inc(TYPE *dest, TYPE pe); +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_atomic_inc(rocshmem_ctx_t ctx, TYPE *dest, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param pe: PE of the remote process. + + :return: None. + +**Description:** +This function atomically adds ``1`` to ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in STANDARD_AMO_TYPES_. + +SHMEM_ATOMIC_FETCH_ADD +---------------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_fetch_add(TYPE *dest, TYPE value, TYPE pe); +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_add(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically added. + :param pe: PE of the remote process. + + :return: The old value of ``dest``. + +**Description:** +This function atomically adds ``value`` to ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in STANDARD_AMO_TYPES_. + +SHMEM_ATOMIC_ADD +---------------- + +.. cpp:function:: __device__ void rocshmem_TYPENAME_atomic_add(TYPE *dest, TYPE value, TYPE pe); +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_atomic_add(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically added. + :param pe: PE of the remote process. + + :return: None. + +**Description:** +This function atomically adds ``value`` to ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values can be seen in STANDARD_AMO_TYPES_. + +SHMEM_ATOMIC_FETCH_AND +---------------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_fetch_and(TYPE *dest, TYPE value, TYPE pe); +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_and(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically ``AND``. + :param pe: PE of the remote process. + + :return: The old value of ``dest``. + +**Description:** +This function atomically bitwise-and ``value`` to the value at ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in BITWISE_AMO_TYPES_. + +SHMEM_ATOMIC_AND +---------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_and(TYPE *dest, TYPE value, TYPE pe); +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_and(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically ``AND``. + :param pe: PE of the remote process. + + :return: None + +**Description:** +This function atomically bitwise-and ``value`` to the value at ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in BITWISE_AMO_TYPES_. + +SHMEM_ATOMIC_FETCH_OR +---------------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_fetch_or(TYPE *dest, TYPE value, TYPE pe) +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_or(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe) + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically ``OR``. + :param pe: PE of the remote process. + + :return: The old value of ``dest``. + +**Description:** +This function atomically bitwise-or ``value`` to the value at ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in BITWISE_AMO_TYPES_. + +SHMEM_ATOMIC_OR +--------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_or(TYPE *dest, TYPE value, TYPE pe) +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_or(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe) + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically ``OR``. + :param pe: PE of the remote process. + + :return: None. + +**Description:** +This function atomically bitwise-or ``value`` to the value at ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in BITWISE_AMO_TYPES_. + +SHMEM_ATOMIC_FETCH_XOR +---------------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_fetch_xor(TYPE *dest, TYPE value, TYPE pe); +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_fetch_xor(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe); + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically ``XOR``. + :param pe: PE of the remote process. + + :return: The old value of ``dest``. + +**Description:** +This function atomically bitwise-xor ``value`` to the value at ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in BITWISE_AMO_TYPES_. + +SHMEM_ATOMIC_XOR +---------------- + +.. cpp:function:: __device__ TYPE rocshmem_TYPENAME_atomic_xor(TYPE *dest, TYPE value, TYPE pe) +.. cpp:function:: __device__ TYPE rocshmem_ctx_TYPENAME_atomic_xor(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, TYPE pe) + + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: The value to be atomically ``XOR``. + :param pe: PE of the remote process. + + :return: None. + +**Description:** +This function atomically bitwise-xor ``value`` to the value at ``dest`` on ``pe``. +The operation is blocking. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in BITWISE_AMO_TYPES_. + +Supported AMO data types +------------------------ + +.. _STANDARD_AMO_TYPES: + +.. list-table:: Standard AMO Data Types + :widths: 10 20 20 + :header-rows: 1 + + * - TYPE + - TYPENAME + - Supported + * - int + - int + - Yes + * - long + - long + - Yes + * - long long + - longlong + - Yes + * - unsigned int + - uint + - Yes + * - unsigned long + - ulong + - Yes + * - unsigned long long + - ulonglong + - Yes + * - int32_t + - int32 + - Yes + * - int64_t + - int64 + - Yes + * - uint32_t + - uint32 + - Yes + * - uint64_t + - uint64 + - Yes + * - size_t + - size + - Yes + * - ptrdiff_t + - ptrdiff + - Yes + +.. _EXTENDED_AMO_TYPES: + +.. list-table:: Extended AMO Data Types + :widths: 10 20 20 + :header-rows: 1 + + * - TYPE + - TYPENAME + - Supported + * - float + - float + - Yes + * - double + - double + - Yes + * - int + - int + - Yes + * - long + - long + - Yes + * - long long + - longlong + - Yes + * - unsigned int + - uint + - Yes + * - unsigned long + - ulong + - Yes + * - unsigned long long + - ulonglong + - Yes + * - int32_t + - int32 + - Yes + * - int64_t + - int64 + - Yes + * - uint32_t + - uint32 + - Yes + * - uint64_t + - uint64 + - Yes + * - size_t + - size + - Yes + * - ptrdiff_t + - ptrdiff + - Yes + +.. _BITWISE_AMO_TYPES: + +.. list-table:: Bitwise AMO Data Types + :widths: 10 20 20 + :header-rows: 1 + + * - TYPE + - TYPENAME + - Supported + * - unsigned int + - uint + - Yes + * - unsigned long + - ulong + - Yes + * - unsigned long long + - ulonglong + - Yes + * - int32_t + - int32 + - Yes + * - int64_t + - int64 + - Yes + * - uint32_t + - uint32 + - Yes + * - uint64_t + - uint64 + - Yes + diff --git a/projects/rocshmem/docs/api/coll.rst b/projects/rocshmem/docs/api/coll.rst new file mode 100644 index 0000000000..63709b6a96 --- /dev/null +++ b/projects/rocshmem/docs/api/coll.rst @@ -0,0 +1,323 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-coll: + +--------------------------- +Collective routines +--------------------------- + +ROCSHMEM_BARRIER_ALL +-------------------- + +.. cpp:function:: __device__ void rocshmem_barrier_all() +.. cpp:function:: __device__ void rocshmem_barrier_all_wave() +.. cpp:function:: __device__ void rocshmem_barrier_all_wg() + + :returns: None. + +**Description:** +This routine performs a collective barrier across all PEs in the system. +The caller is blocked until the barrier is resolved and all updates local and remote are completed. +These APIs should be called from only one thread/wavefront/workgroup within the grid to avoid undefined behavior. + +ROCSHMEM_BARRIER_ALL_ON_STREAM +------------------------------- + +.. cpp:function:: __host__ void rocshmem_barrier_all_on_stream(hipStream_t stream) + + :param stream: HIP stream on which to enqueue the operation. + :returns: None. + +**Description:** +This routine enqueues a collective barrier operation on a HIP stream. The barrier is performed +across all PEs in the system. The operation is enqueued on the specified stream and will execute +asynchronously. The caller must synchronize the stream (e.g., using ``hipStreamSynchronize``) +to ensure completion. + +ROCSHMEM_BARRIER +---------------- + +.. cpp:function:: __device__ void rocshmem_ctx_barrier(rocshmem_ctx_t ctx, rocshmem_team_t team) +.. cpp:function:: __device__ void rocshmem_ctx_barrier_wave(rocshmem_ctx_t ctx, rocshmem_team_t team) +.. cpp:function:: __device__ void rocshmem_ctx_barrier_wg(rocshmem_ctx_t ctx, rocshmem_team_t team) + + :param ctx: Context with which to perform this operation. + :returns: None. + +**Description:** +This routine performs a collective barrier between all PEs in the system. +The caller is blocked until the barrier is resolved. + +ROCSHMEM_TEAM_SYNC +------------------ + +.. cpp:function:: __device__ void rocshmem_ctx_sync(rocshmem_ctx_t ctx, rocshmem_team_t team) +.. cpp:function:: __device__ void rocshmem_ctx_sync_wave(rocshmem_ctx_t ctx, rocshmem_team_t team) +.. cpp:function:: __device__ void rocshmem_ctx_sync_wg(rocshmem_ctx_t ctx, rocshmem_team_t team) + + :param ctx: Context with which to perform this operation. + :param team: Team with which to perform this operation. + :returns: None. + +**Description:** +This routine registers the arrival of a PE at a barrier. +The caller is blocked until the synchronization is resolved. + +Unlike the ``shmem_barrier_all`` routine, ``shmem_team_sync`` only ensures the +completion and visibility of previously issued memory stores, but does not +ensure the completion of remote memory updates issued via OpenSHMEM routines. + +ROCSHMEM_SYNC_ALL +----------------- + +.. cpp:function:: __device__ void rocshmem_sync_all() +.. cpp:function:: __device__ void rocshmem_sync_all_wave() +.. cpp:function:: __device__ void rocshmem_sync_all_wg() + + :returns: None. + +**Description:** +These routines behaves the same way as ``rocshmem_team_sync_*`` when called on the world team. +These APIs should be called from only one thread/wavefront/workgroup within the grid to avoid undefined behavior. + +ROSHMEM_ALLTOALL +---------------- + +.. cpp:function:: __device__ void rocshmem_TYPENAME_alltoall_wg(rocshmem_team_t team, TYPE *dest, const TYPE *source, int nelems) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_alltoall_wg(rocshmem_ctx_t ctx, rocshmem_team_t team, TYPE *dest, const TYPE *source, int nelems) + + :param team: The team participating in the collective. + :param dest: Destination address. Must be an address on the + symmetric heap. + :param source: Source address. Must be an address on the symmetric + heap. + :param nelems: Number of data blocks transferred per pair of PEs. + :returns: None. + +**Description:** +This routine exchanges a fixed amount of contiguous data blocks between all pairs +of PEs participating in the collective routine. +This function must be called as a work-group collective. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`RMA_TYPES`. + +ROCSHMEM_ALLTOALLMEM_ON_STREAM +------------------------------- + +.. cpp:function:: __host__ void rocshmem_alltoallmem_on_stream(rocshmem_team_t team, void *dest, const void *source, size_t size, hipStream_t stream) + + :param team: The team participating in the collective. + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param size: Number of bytes to transfer per pair of PEs. + :param stream: HIP stream on which to enqueue the operation. + :returns: None. + +**Description:** +This routine enqueues an alltoall collective operation on a HIP stream. The function +exchanges a fixed amount of contiguous data blocks between all pairs of PEs participating +in the collective routine. The operation is enqueued on the specified stream and will +execute asynchronously. The caller must synchronize the stream (e.g., using +``hipStreamSynchronize``) to ensure completion. + +This function creates a separate context for each workgroup to avoid contention on the +default context, allowing parallel execution across multiple streams. + +ROCSHMEM_BROADCAST +------------------ + +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_broadcast_wg(rocshmem_ctx_t ctx, rocshmem_team_t team, TYPE *dest, const TYPE *source, int nelems, int pe_root) + + :param ctx: Context with which to perform this collective. + :param team: The team participating in the collective. + :param dest: Destination address. Must be an address on the + symmetric heap. + :param source: Source address. Must be an address on the symmetric + heap. + :param nelems: Number of data blocks transferred per pair of PEs. + :returns: None. + +**Description:** +This routine performs a broadcast across PEs in the team. +The caller is blocked until the broadcast completes. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`RMA_TYPES`. + +ROCSHMEM_BROADCASTMEM_ON_STREAM +-------------------------------- + +.. cpp:function:: __host__ void rocshmem_broadcastmem_on_stream(rocshmem_team_t team, void *dest, const void *source, size_t nelems, int pe_root, hipStream_t stream) + + :param team: The team participating in the collective. + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: Number of bytes to broadcast. + :param pe_root: Root PE (relative to team) from which to broadcast. + :param stream: HIP stream on which to enqueue the operation. + :returns: None. + +**Description:** +This routine enqueues a broadcast collective operation on a HIP stream. The function broadcasts +data from the root PE to all other PEs participating in the collective routine. The operation +is enqueued on the specified stream and will execute asynchronously. The caller must synchronize +the stream (e.g., using ``hipStreamSynchronize``) to ensure completion. + +This function creates a separate context for each workgroup to avoid contention on the +default context, allowing parallel execution across multiple streams. + +ROCSHMEM_FCOLLECT +----------------- + +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_fcollect_wg(rocshmem_ctx_t ctx, rocshmem_team_t team, TYPE *dest, const TYPE *source, int nelems) + + :param ctx: Context with which to perform this collective. + :param team: The team participating in the collective. + :param dest: Destination address. Must be an address on the + symmetric heap. + :param source: Source address. Must be an address on the symmetric + heap. + :param nelems: Number of data blocks transferred per pair of PEs. + :returns: None. + +**Description:** +This routine concatenates blocks of data from multiple PEs to an array in every +PE participating in the collective routine. + +ROCSHMEM_REDUCTION +------------------ +.. cpp:function:: __device__ int rocshmem_ctx_TYPENAME_OPNAME_reduce_wg(rocshmem_ctx_t ctx, rocshmem_team_t team, TYPE *dest, const TYPE *source, int nreduce) + + :param ctx: Context with which to perform this collective. + :param team: The team participating in the collective. + :param dest: Destination address. Must be an address on the + symmetric heap. + :param source: Source address. Must be an address on the symmetric + heap. + :param nreduce: Number of data blocks transferred per pair of PEs. + :returns: Zero on successful local completion. Nonzero otherwise. + + +**Description:** +This routine performs an allreduce operation across PEs in the team. + +Valid ``TYPENAME``, ``TYPE``, and ``OPNAME`` values are listed in :ref:`REDUCE_TYPES`. + +Supported reduction types and operations +---------------------------------------- + +.. _REDUCE_TYPES: + +.. list-table:: Reduction Types, Names and Operations + :widths: 20 20 20 20 + :header-rows: 1 + + * - TYPE + - TYPENAME + - OPNAME + - Supported + * - char + - char + - max, min, sum, prod + - No + * - signed char + - schar + - max, min, sum, prod + - No + * - short + - short + - max, min, sum, prod + - Yes + * - int + - int + - max, min, sum, prod + - Yes + * - long + - long + - max, min, sum, prod + - Yes + * - long long + - longlong + - max, min, sum, prod + - Yes + * - ptrdiff_t + - ptrdiff + - max, min, sum, prod + - No + * - unsigned char + - uchar + - and, or, xor, max, min, sum, prod + - No + * - unsigned short + - ushort + - and, or, xor, max, min, sum, prod + - No + * - unsigned int + - uint + - and, or, xor, max, min, sum, prod + - No + * - unsigned long + - ulong + - and, or, xor, max, min, sum, prod + - No + * - unsigned long long + - ulonglong + - and, or, xor, max, min, sum, prod + - No + * - int8_t + - int8 + - and, or, xor, max, min, sum, prod + - No + * - int16_t + - int16 + - and, or, xor, max, min, sum, prod + - No + * - int32_t + - int32 + - and, or, xor, max, min, sum, prod + - No + * - int64_t + - int64 + - and, or, xor, max, min, sum, prod + - No + * - uint8_t + - uint8 + - and, or, xor, max, min, sum, prod + - No + * - uint16_t + - uint16 + - and, or, xor, max, min, sum, prod + - No + * - uint32_t + - uint32 + - and, or, xor, max, min, sum, prod + - No + * - uint64_t + - uint64 + - and, or, xor, max, min, sum, prod + - No + * - size_t + - size + - and, or, xor, max, min, sum, prod + - No + * - float + - float + - max, min, sum, prod + - Yes + * - double + - double + - max, min, sum, prod + - Yes + * - long double + - longdouble + - max, min, sum, prod + - No + * - double _Complex + - complexd + - sum, prod + - No + * - float _Complex + - complexf + - sum, prod + - No diff --git a/projects/rocshmem/docs/api/ctx.rst b/projects/rocshmem/docs/api/ctx.rst new file mode 100644 index 0000000000..a32c823482 --- /dev/null +++ b/projects/rocshmem/docs/api/ctx.rst @@ -0,0 +1,59 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-ctx: + +----------------------------------- +Context management routines +----------------------------------- + +ROCSHMEM_CTX_CREATE +------------------- + +.. cpp:function:: __device__ int rocshmem_wg_ctx_create(int64_t options, rocshmem_ctx_t *ctx) +.. cpp:function:: __device__ int rocshmem_wg_team_create_ctx(rocshmem_team_t team, long options, rocshmem_ctx_t *ctx) + + :param team: Team handle to derive the context from. + :param options: Options for context creation. Ignored in current design; use the value ``0``. + :param ctx: A handle to the newly created context. + + :returns: All threads returns ``0`` if the context was created successfully. + If any thread returns non-zero value, the operation fails, ctx is set to ``ROCSHMEM_CTX_INVALID`` and a + higher number of ``ROCSHMEM_MAX_NUM_CONTEXTS`` is required. + +**Description:** +This routine creates an rocSHMEM context. By design, the context is private to the calling work-group. +It must be called collectively by all threads in the work-group. If the context was created successfully, a value +of zero is returned and the context handle pointed to by ctx specifies a valid context; otherwise, a nonzero value +is returned and ctx is set to ``ROCSHMEM_CTX_INVALID``. An unsuccessful context creation call is not treated as an +error and the rocSHMEM library remains in a correct state. The creation call can be reattempted after additional +resources become available. + +ROCSHMEM_CTX_DESTROY +-------------------- + +.. cpp:function:: __device__ void rocshmem_wg_ctx_destroy(rocshmem_ctx_t *ctx) + + :param ctx: Context handle. + + :returns: None. + +**Description:** +This routine destroys an rocSHMEM context. It must be called collectively by all threads in the work-group. +If ctx has the value ``ROCSHMEM_CTX_INVALID``, no operation is performed. + +ROCSHMEM_GET_DEVICE_CTX +----------------------- + +.. cpp:function:: __host__ void * rocshmem_get_device_ctx() + + :param: None. + + :returns: Returns ``ROCSHMEM_CTX_DEFAULT`` device pointer that users. + can query from one instance of rocSHMEM host library and + use later for dynamic module initialization in + kernel bitcode device library in the same application. + +**Description:** +This routine queries rocSHMEM default device context from host API. diff --git a/projects/rocshmem/docs/api/env_variables.rst b/projects/rocshmem/docs/api/env_variables.rst new file mode 100644 index 0000000000..a827de1dc2 --- /dev/null +++ b/projects/rocshmem/docs/api/env_variables.rst @@ -0,0 +1,96 @@ +.. meta:: + :description: rocSHMEM environment variables reference + :keywords: rocSHMEM, ROCm, API, environment variables, environment, reference + +.. _rocshmem-api-env-variables: + +******************************************************************** +rocSHMEM environment variables +******************************************************************** + +This section describes the important environment variables used to +control the behavior of rocSHMEM. + +.. list-table:: + :header-rows: 1 + :widths: 35,14,51 + + * - **Environment variable** + - **Default value** + - **Value** + + * - | ``ROCSHMEM_HEAP_SIZE`` + | Defines the size of the rocSHMEM symmetric heap in bytes (per PE). + - ``1073741824`` (1 GB) + - | Size in bytes (per PE). + | Note: the heap is on GPU memory. + + * - | ``ROCSHMEM_MAX_NUM_CONTEXTS`` + | Defines the number of contexts an application can use. + - ``32`` + - Maximum number of contexts. + + * - | ``ROCSHMEM_MAX_NUM_TEAMS`` + | Defines the number of teams an application can use. + - ``40`` + - Maximum number of teams. + + * - | ``ROCSHMEM_BACKEND`` + | When rocSHMEM is compiled for all backends, this enviroment variable + | selects which backend to execute. The default value is an empty string and rocSHMEM auto-selects the most appropriate backend. + - `` `` + - | ``ipc``: IPC Backend + | ``ro``: Reverse Offload Backend + | ``gda``: GPU Direct Async Backend + + * - | ``ROCSHMEM_UNIQUEID_WITH_MPI`` + | Defines whether rocSHMEM is expected to use MPI when using the uniqueId based initialization. + - ``0`` + - | ``0``: Do not use MPI. + | ``1``: Use MPI. + + * - | ``ROCSHMEM_DISABLE_MIXED_IPC`` + | Defines whether to force using the network conduit even when IPC is available. + - ``0`` + - | ``0``: Use IPC when available. + | ``1``: Force network conduit. + + * - | ``ROCSHMEM_USE_IB_HCA`` + | Defines which NIC that this PE should be bound to. The default value is an empty string and rocSHMEM auto-detects the most appropriate NIC. + - `` `` + - | Example value: ``bnxt_re0`` + + * - | ``ROCSHMEM_BOOTSTRAP_SOCKET_IFNAME`` + | Chooses the interface to bootstrap rocSHMEM with. + | Only valid when not using MPI. + | The default value is an empty string and rocSHMEM auto-detects the most appropriate interface. + - `` `` + - | Example value: ``eno8303`` + + * - | ``ROCSHMEM_GDA_PROVIDER`` + | When rocSHMEM is compiled with support for multiple NIC vendors, + | the enviroment variable selects the desired provider. + | The default value is an empty string and rocSHMEM auto-detects the most appropriate NIC. + - `` `` + - | ``bnxt``: Broadcom Thor 2 + | ``pensando``: AMD Pensando Pollara + | ``ionic``: AMD Pensando Pollara (alias) + | ``mlx5``: Mellanox ConnectX-7 + + * - | ``ROCSHMEM_GDA_ALTERNATE_QP_PORTS`` + | Enables or disables alternating QP mappings across rocSHMEM contexts. + - ``1`` + - | ``0``: Disabled. + | ``1``: Enabled. This helps saturate bandwidth on multiport bonded interfaces. + + * - | ``ROCSHMEM_GDA_TRAFFIC_CLASS`` + | When using an NIC with an Ethernet link layer, this sets the traffic class for the QPs. + - ``0`` + - The traffic class number. + + * - | ``ROCSHMEM_GDA_PCIE_RELAXED_ORDERING`` + | Enables PCIe Relaxed Ordering when registering the symmetric heap with the RDMA NICs. + - ``0`` + - | ``0``: Disabled. + | ``1``: Enabled. + diff --git a/projects/rocshmem/docs/api/init.rst b/projects/rocshmem/docs/api/init.rst new file mode 100644 index 0000000000..ae358af2ad --- /dev/null +++ b/projects/rocshmem/docs/api/init.rst @@ -0,0 +1,161 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-init: + +--------------------------------------- +Library setup, exit, and query routines +--------------------------------------- + +ROCSHMEM_INIT +------------- + +.. cpp:function:: __host__ void rocshmem_init(void) + + :Parameters: None. + :returns: None. + +**Description:** +This routine initializes the rocSHMEM library and underlying transport layer. +Before ``rocshmem_init`` is called, +you must select the device that this PE is associated to by calling +`hipSetDevice +`_. + +.. WARNING:: + Routine `rocshmem_wg_init` has been deprecated. + +.. cpp:function:: [[deprecated]] __device__ void rocshmem_wg_init(void) + + :Parameters: None. + :returns: None. + +**Description:** +This routine has been deprecated, please do not use. +This routine initializes device-side rocSHMEM resources. +It must be called before any threads in this work-group invoke other rocSHMEM functions. +It must be called collectively by all threads in the work-group. + +ROCSHMEM_FINALIZE +----------------- +.. cpp:function:: __host__ void rocshmem_finalize(void) + + :Parameters: None. + :returns: None. + +**Description:** +This routine finalizes the rocSHMEM library. + +.. WARNING:: + Routine `rocshmem_wg_finalize` has been deprecated. + +.. cpp:function:: [[deprecated]] __device__ void rocshmem_wg_finalize(void) + + :Parameters: None. + :returns: None. + +**Description:** +This routine has been deprecated, please do not use. +This routine finalizes device-side rocSHMEM resources. +It must be called before work-group completion if the work-group also called ``rocshmem_wg_init``. +It must be called collectively by all threads in the work-group. + +ROCSHMEM_INIT_ATTR +------------------ +.. cpp:function:: __host__ int rocshmem_init_attr(unsigned int flags, rocshmem_init_attr_t *attr) + + :param flags: The initialization method to be used. + :param attr: Attribute structure specifying input characteristics. + + :returns int: Returns ``0`` on success; otherwise, returns a nonzero value. + +**Description:** +This routine initializes the rocSHMEM runtime and underlying transport layer using +the provided mode and attributes. +The parameter ``flags`` can be either +``ROCSHMEM_INIT_WITH_UNIQUEID`` or ``ROCSHMEM_INIT_WITH_MPI_COMM``. + +ROCSHMEM_GET_UNIQUEID +--------------------- +.. cpp:function:: __host__ int rocshmem_get_uniqueid(rocshmem_uniqueid_t *uid) + + :param uid: Pointer to a unique ID handle. + :returns: Returns ``0`` on success; otherwise, returns a nonzero value. + +**Description:** +This routine returns a unique ID. + +ROCSHMEM_SET_ATTR_UNIQUEID_ARGS +------------------------------- +.. cpp:function:: __host__ int rocshmem_set_attr_uniqueid_args(int rank, int nranks, rocshmem_uniqueid_t *uid, rocshmem_init_attr_t *attr) + + :param rank: Rank of the calling process. + :param nranks: Number of PEs. + :param uid: Unique ID used to identify the group processes. + :param attr: Attribute structure to be passed to ``rocshmem_init_attr_t``. + + :returns: Returns ``0`` on success; otherwise, returns a nonzero value. + +**Description:** +This routine initializes the ``rocshmem_init_attr_t`` struct. + +ROCSHMEM_N_PES +-------------- + +.. cpp:function:: __host__ int rocshmem_n_pes(void) + + :Parameters: None. + :returns: Total number of PEs. + +**Description:** +This routine queries the total number of PEs. +It can be called before ``rocshmem_init``. + +.. cpp:function:: __device__ int rocshmem_n_pes(void) +.. cpp:function:: __device__ int rocshmem_ctx_n_pes(rocshmem_ctx_t ctx) + + :param ctx: GPU side context handle. + :returns: Total number of PEs. + +**Description:** +This routine queries the total number of PEs for a given context. +It can be called per thread with no performance penalty. + +ROCSHMEM_MY_PE +-------------- + +.. cpp:function:: __host__ int rocshmem_my_pe(void) + + :Parameters: None. + :returns: PE ID of the caller. + +**Description:** +This routine queries the PE ID of the caller. +It can be called before ``rocshmem_init``. + +.. cpp:function:: __device__ int rocshmem_my_pe(void) +.. cpp:function:: __device__ int rocshmem_ctx_my_pe(rocshmem_ctx_t ctx) + + :param ctx: GPU side context handle. + :returns: PE ID of the caller. + +**Description:** +This routine queries the PE ID of the caller. +It can be called per thread with no performance penalty. + +ROCSHMEM_PTR +-------------- + +.. cpp:function:: __host__ void* rocshmem_ptr(const void *dest, int pe); +.. cpp:function:: __device__ void* rocshmem_ptr(const void *dest, int pe); + + :param dest: Local symmetric heap allocation pointer for current PE. + :param pe: Remote PE. + :returns: Returns remote symmetric heap device pointer from host-side API. + ``NULL`` is returned if a valid device pointer cannot be provided. + This pointer can be used to issue load/store from custom kernels + instead of using rocshmem device side get/put APIs for RMA operations. + +**Description:** +This routine queries rocSHMEM remote symmetric heap pointer. diff --git a/projects/rocshmem/docs/api/memory_management.rst b/projects/rocshmem/docs/api/memory_management.rst new file mode 100644 index 0000000000..a5c7b8b7b6 --- /dev/null +++ b/projects/rocshmem/docs/api/memory_management.rst @@ -0,0 +1,35 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-memory-management: + + +--------------------------- +Memory management routines +--------------------------- + +ROCSHMEM_MALLOC +--------------- + +.. cpp:function:: __host__ void *rocshmem_malloc(size_t size) + + :param size: Memory allocation size in bytes. + :returns: A pointer to the allocated memory on the symmetric heap. + If a valid allocation cannot be made, it returns ``NULL``. + +**Description:** +This routine allocates memory of ``size`` bytes from the symmetric heap. +This is a collective operation and must be called by all PEs. + +ROCSHMEM_FREE +------------- + +.. cpp:function:: __host__ void rocshmem_free(void *ptr) + + :param ptr: A pointer to previously allocated memory on the symmetric heap. + :returns: None. + +**Description:** +This routine frees a memory allocation from the symmetric heap. +It is a collective operation and must be called by all PEs. diff --git a/projects/rocshmem/docs/api/memory_ordering.rst b/projects/rocshmem/docs/api/memory_ordering.rst new file mode 100644 index 0000000000..00c12834f4 --- /dev/null +++ b/projects/rocshmem/docs/api/memory_ordering.rst @@ -0,0 +1,51 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-memory-ordering: + +--------------------------- +Memory ordering routines +--------------------------- + +ROCSHMEM_FENCE +-------------- + +.. cpp:function:: __device__ void rocshmem_fence() +.. cpp:function:: __device__ void rocshmem_fence(int pe) +.. cpp:function:: __device__ void rocshmem_ctx_fence(rocshmem_ctx_t ctx) +.. cpp:function:: __device__ void rocshmem_ctx_fence(rocshmem_ctx_t ctx, int pe) + + :param ctx: Context with which to perform this operation. + :param pe: Destination ``pe``. + :returns: None. + +**Description:** +This routine ensures order between messages in this context to follow OpenSHMEM semantics. + +ROCSHMEM_QUIET +-------------- + +.. cpp:function:: __device__ void rocshmem_ctx_quiet(rocshmem_ctx_t ctx) +.. cpp:function:: __device__ void rocshmem_quiet() + + :param ctx: Context with which to perform this operation. + :returns: None. + +**Description:** +This routine completes all previous operations posted to this context. + +ROCSHMEM_PE_QUIET +----------------- + +.. cpp:function:: __device__ void rocshmem_ctx_pe_quiet(shmem_ctx_t ctx, const int *target_pes, size_t npes) +.. cpp:function:: __device__ void rocshmem_pe_quiet(const int *target_pes, size_t npes) + + :param ctx: Context with which to perform this operation. + :param target_pes: Address of target PE array where the operations need to be completed + :param npes: The number of PEs in the target PE array + :returns: None. + +**Description:** +This routine completes all previous operations posted to this context +for the PEs in the `target_pes` array. diff --git a/projects/rocshmem/docs/api/pt2pt_sync.rst b/projects/rocshmem/docs/api/pt2pt_sync.rst new file mode 100644 index 0000000000..6985b61e02 --- /dev/null +++ b/projects/rocshmem/docs/api/pt2pt_sync.rst @@ -0,0 +1,142 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-pt2pt-sync: + +----------------------------------------- +Point-to-point synchronization routines +----------------------------------------- + +ROCSHMEM_WAIT_UNTIL +------------------- + +.. cpp:function:: __device__ void rocshmem_TYPENAME_wait_until(TYPE *ivars, int cmp, TYPE val) + + :param ivars: Pointer to memory on the symmetric heap to wait for. + :param cmp: Operation for the comparison. + :param val: Value to compare the memory at ``ivars`` to. + :returns: None. + +**Description:** +This routine blocks the caller until the condition ``(*ivars cmp val)`` is true. + +Valid ``cmp`` values are listed in :ref:`CMP_VALUES`. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`. + +ROCSHMEM_WAIT_UNTIL_ALL +----------------------- + +.. cpp:function:: __device__ void rocshmem_TYPENAME_wait_until_all(TYPE *ivars, size_t nelems, const int* status, int cmp, TYPE val) + + :param ivars: Pointer to memory on the symmetric heap to wait for. + :param nelems: Number of elements in the ``ivars`` array. + :param status: Array of length ``nelems`` to exclude elements from the wait. + :param cmp: Operation for the comparison. + :param val: Value to compare. + :returns: None. + +**Description:** +This routine blocks the caller until the condition ``(ivars[i] cmp val)`` is true for all ``ivars``. + +Valid ``cmp`` values are listed in :ref:`CMP_VALUES`. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`. + +ROCSHMEM_WAIT_UNTIL_ANY +----------------------- +.. cpp:function:: __device__ size_t rocshmem_TYPENAME_wait_until_any(TYPE *ivars, size_t nelems, const int* status, int cmp, TYPE val) + + :param ivars: Pointer to memory on the symmetric heap to wait for. + :param nelems: Number of elements in the ``ivars`` array. + :param status: Array of length ``nelems`` to exclude elements from the wait. + :param cmp: Operation for the comparison. + :param val: Value to compare. + :returns: The index of an element in the ``ivars`` array that satisfies the wait condition. If the wait set is empty, this routine returns ``SIZE_MAX``. + +**Description:** +This routine blocks the caller until any of the condition ``(ivars[i] cmp val)`` is true. + +Valid ``cmp`` values are listed in :ref:`CMP_VALUES`. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`. + +ROCSHMEM_WAIT_UNTIL_SOME +------------------------ + +.. cpp:function:: __device__ size_t rocshmem_TYPENAME_wait_until_some(TYPE *ivars, size_t nelems, size_t* indices, const int* status, int cmp, TYPE val) + + :param ivars: Pointer to memory on the symmetric heap to wait for. + :param nelems: Number of elements in the ``ivars`` array. + :param indices: List of indices with a length of at least ``nelems``. + :param status: Array of length ``nelems`` to exclude elements from the wait. + :param cmp: Operation for the comparison. + :param val: Value to compare. + :returns: The number of indices returned in the indices array. If the wait set is empty, this routine returns ``0``. + +**Description:** +This routine blocks the caller until any of the conditions ``(ivars[i] cmp val)`` is true. + +Valid ``cmp`` values are listed in :ref:`CMP_VALUES`. + +Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`STANDARD_AMO_TYPES`. + +ROCSHMEM_TEST +------------- + +.. cpp:function:: __device__ int rocshmem_TYPENAME_test(TYPE *ivars, int cmp, TYPE val) + + :param ivars: Pointer to memory on the symmetric heap to wait for. + :param cmp: Operation for the comparison. + :param val: Value to compare the memory at ``ivars`` to. + + :returns: ``1`` if the evaluation is true. ``0`` otherwise. + +**Description:** +This routine tests if the condition ``(*ivars cmp val)`` is true. + +ROCSHMEM_SIGNAL_WAIT_UNTIL_ON_STREAM +------------------------------------- + +.. cpp:function:: __host__ void rocshmem_signal_wait_until_on_stream(uint64_t *sig_addr, int cmp, uint64_t cmp_value, hipStream_t stream) + + :param sig_addr: Address of the signal variable on the symmetric heap. + :param cmp: Comparison operator (e.g., ROCSHMEM_CMP_EQ, ROCSHMEM_CMP_GE, etc.). + :param cmp_value: Value to compare against. + :param stream: HIP stream on which to enqueue the operation. + :returns: None. + +**Description:** +This routine enqueues a wait operation on a HIP stream. The function blocks the calling thread +until the signal variable at ``sig_addr`` satisfies the comparison condition ``(*sig_addr cmp cmp_value)``. +The wait operation is executed asynchronously on the specified stream. The caller must synchronize +the stream (e.g., using ``hipStreamSynchronize``) to ensure the wait condition has been satisfied. + +Valid ``cmp`` values are listed in :ref:`CMP_VALUES`. + +.. _CMP_VALUES: + +Supported comparisons +--------------------- + +The following table lists the point-to-point comparison constants: + +.. list-table:: Point-to-Point Comparison Constants + :widths: 20 20 + :header-rows: 1 + + * - Constant + - Description + * - ROCSHMEM_CMP_EQ + - Equal + * - ROCSHMEM_CMP_NE + - Not equal + * - ROCSHMEM_CMP_GT + - Greater than + * - ROCSHMEM_CMP_GE + - Greater than or equal to + * - ROCSHMEM_CMP_LT + - Less than + * - ROCSHMEM_CMP_LE + - Less than or equal to diff --git a/projects/rocshmem/docs/api/rma.rst b/projects/rocshmem/docs/api/rma.rst new file mode 100644 index 0000000000..d4bf4b09b5 --- /dev/null +++ b/projects/rocshmem/docs/api/rma.rst @@ -0,0 +1,278 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-rma: + +----------------------------------------- +Remote memory access routines +----------------------------------------- + +- Routines with the ``_wave`` and ``_wg`` suffixes require all threads in a wavefront and workgroup, respectively, + to call the routine with the same parameters. +- Routines with the ``_nbi`` substring will return as soon as the request is posted. +- Routines without the ``_nbi`` substring will block until the operation completes locally. +- Valid ``TYPENAME`` and ``TYPE`` values can be found in RMA_TYPES_. + +ROCSHMEM_PUT +------------ + +.. cpp:function:: __device__ void rocshmem_TYPENAME_put(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_wave(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_wg(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_nbi(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_nbi_wave(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_nbi_wg(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_nbi(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_nbi_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_nbi_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: The number of elements to transfer. + :param pe: PE of the remote process. + :returns: None. + +**Description:** +This routine writes contiguous data of ``nelems`` elements from source on the calling PE to ``dest`` at ``pe``. + +ROCSHMEM_PUTMEM +--------------- + +.. cpp:function:: __device__ void rocshmem_putmem(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_wave(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_wg(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_nbi(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_nbi(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_nbi_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_nbi_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: Size of the transfer in bytes. + :param pe: PE of the remote process. + + :returns: None. + +**Description:** +This routine writes contiguous data of ``nelems`` bytes from source on the calling PE to ``dest`` at ``pe``. + +ROCSHMEM_PUTMEM_ON_STREAM +-------------------------- + +.. cpp:function:: __host__ void rocshmem_putmem_on_stream(void *dest, const void *source, size_t nelems, int pe, hipStream_t stream) + + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: Size of the transfer in bytes. + :param pe: PE of the remote process. + :param stream: HIP stream on which to enqueue the operation. + + :returns: None. + +**Description:** +This routine enqueues a putmem RMA operation on a HIP stream. The function writes contiguous +data of ``nelems`` bytes from source on the calling PE to ``dest`` at ``pe``. The operation +is enqueued on the specified stream and will execute asynchronously. The caller must +synchronize the stream (e.g., using ``hipStreamSynchronize``) to ensure completion. + +ROCSHMEM_P +---------- + +.. cpp:function:: __device__ void rocshmem_TYPENAME_p(TYPE *dest, TYPE value, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_p(rocshmem_ctx_t ctx, TYPE *dest, TYPE value, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param value: Value to write to ``dest`` at ``pe``. + :param pe: PE of the remote process. + + :returns: None. + +**Description:** +This routine writes a single value to to ``dest`` at ``pe``. + +ROCSHMEM_GET +------------ + +.. cpp:function:: __device__ void rocshmem_TYPENAME_get(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_get_wave(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_get_wg(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_get_nbi(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_get_nbi_wave(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_get_nbi_wg(TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_get(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_get_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_get_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_get_nbi(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_get_nbi_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_get_nbi_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address; Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: The number of elements to transfer. + :param pe: PE of the remote process. + + :returns: None. + +**Description:** +This routine reads contiguous data of ``nelems`` elements from source on ``pe`` to ``dest`` on the calling PE. + +ROCSHMEM_GETMEM +--------------- + +.. cpp:function:: __device__ void rocshmem_getmem(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_getmem_wave(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_getmem_wg(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_getmem_nbi(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_getmem_nbi_wave(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_getmem_nbi_wg(void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_getmem(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_getmem_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_getmem_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_getmem_nbi(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_getmem_nbi_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_getmem_nbi_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: Size of the transfer in bytes. + :param pe: PE of the remote process. + + :returns: None. + +**Description:** +This routine reads contiguous data of ``nelems`` bytes from source on ``pe`` to ``dest`` on the calling PE. + +ROCSHMEM_GETMEM_ON_STREAM +-------------------------- + +.. cpp:function:: __host__ void rocshmem_getmem_on_stream(void *dest, const void *source, size_t nelems, int pe, hipStream_t stream) + + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: Size of the transfer in bytes. + :param pe: PE of the remote process. + :param stream: HIP stream on which to enqueue the operation. + + :returns: None. + +**Description:** +This routine enqueues a getmem RMA operation on a HIP stream. The function reads contiguous +data of ``nelems`` bytes from source on ``pe`` to ``dest`` on the calling PE. The operation +is enqueued on the specified stream and will execute asynchronously. The caller must +synchronize the stream (e.g., using ``hipStreamSynchronize``) to ensure completion. + +ROCSHMEM_G +---------- +.. cpp:function:: __device__ float rocshmem_ctx_float_g(rocshmem_ctx_t ctx, const float *source, int pe) +.. cpp:function:: __device__ float rocshmem_float_g(const float *source, int pe) + + :param ctx: Context with which to perform this operation. + :param source: Source address. Must be an address on the symmetric heap. + :param pe: PE of the remote process. + + :returns: The value read from source at ``pe``. + +**Description:** +This routine reads and returns single value from source at ``pe``. + +Supported RMA data types +------------------------ + +The following table lists the supported RMA data types: + +.. _RMA_TYPES: + +.. list-table:: RMA Data Types + :widths: 10 20 20 + :header-rows: 1 + + * - TYPE + - TYPENAME + - Supported + * - float + - float + - Yes + * - double + - double + - Yes + * - long double + - longdouble + - No + * - char + - char + - Yes + * - signed char + - schar + - Yes + * - short + - short + - Yes + * - int + - int + - Yes + * - long + - long + - Yes + * - long long + - longlong + - Yes + * - unsigned char + - uchar + - Yes + * - unsigned short + - ushort + - Yes + * - unsigned int + - uint + - Yes + * - unsigned long + - ulong + - Yes + * - unsigned long long + - ulonglong + - Yes + * - int8_t + - int8 + - No + * - int16_t + - int16 + - No + * - int32_t + - int32 + - No + * - int64_t + - int64 + - Yes + * - uint8_t + - uint8 + - No + * - uint16_t + - uint16 + - No + * - uint32_t + - uint32 + - No + * - uint64_t + - uint64 + - No + * - size_t + - size + - No + * - ptrdiff_t + - ptrdiff + - No + diff --git a/projects/rocshmem/docs/api/sigops.rst b/projects/rocshmem/docs/api/sigops.rst new file mode 100644 index 0000000000..578eb860bb --- /dev/null +++ b/projects/rocshmem/docs/api/sigops.rst @@ -0,0 +1,125 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-sigops: + +--------------------- +Signaling operations +--------------------- + +ROCSHMEM_PUTMEM_SIGNAL +---------------------- + +.. cpp:function:: __device__ void rocshmem_putmem_signal(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_signal_wave(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_signal_wg(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_signal_nbi(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_signal_nbi_wave(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_putmem_signal_nbi_wg(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_signal(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_signal_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_signal_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_signal_nbi(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_signal_nbi_wave(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_putmem_signal_nbi_wg(rocshmem_ctx_t ctx, void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: The number of bytes to transfer. + :param sig_addr: Signal address. Must be an address on the symmetric heap. + :param signal: Signal value. + :param sig_op: Atomic operation to apply the signal value. + :param pe: PE of the remote process. + :returns: None. + +**Description:** +This function writes contiguous data of ``nelems`` bytes from source on the calling PE to ``dest`` at ``pe``, +then applies ``sig_op`` at ``sig_addr`` with the signal value. +Valid ``sig_op values`` are listed in SIGNAL_OPERATORS_. + +ROCSHMEM_PUT_SIGNAL +------------------- + +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_signal(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_signal_wave(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_signal_wg(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_signal_nbi(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_signal_nbi_wave(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_TYPENAME_put_signal_nbi_wg(TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_signal(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_signal_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_signal_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_signal_nbi(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_signal_nbi_wave(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) +.. cpp:function:: __device__ void rocshmem_ctx_TYPENAME_put_signal_nbi_wg(rocshmem_ctx_t ctx, TYPE *dest, const TYPE *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe) + + :param ctx: Context with which to perform this operation. + :param dest: Destination address. Must be an address on the symmetric heap. + :param source: Source address. Must be an address on the symmetric heap. + :param nelems: The number of elements of size ``TYPE`` to transfer. + :param sig_addr: Signal address. Must be an address on the symmetric heap. + :param signal: Signal value. + :param sig_op: Atomic operation to apply the signal value. + :param pe: PE of the remote process. + :returns: None. + +**Description:** +This function writes contiguous data of ``nelems`` elements of ``TYPE`` from source on the calling PE to ``dest`` at ``pe``, +then applies ``sig_op`` at ``sig_addr`` with the signal value. +Valid ``sig_op values`` are listed in SIGNAL_OPERATORS_. +Valid ``TYPENAME`` and ``TYPE`` values are listed in :ref:`RMA_TYPES`. + +ROCSHMEM_PUTMEM_SIGNAL_ON_STREAM +--------------------------------- + +.. cpp:function:: __host__ void rocshmem_putmem_signal_on_stream(void *dest, const void *source, size_t nelems, uint64_t *sig_addr, uint64_t signal, int sig_op, int pe, hipStream_t stream) + + :param dest: Destination address on the remote PE. Must be an address on the symmetric heap. + :param source: Source address on the local PE. Must be an address on the symmetric heap. + :param nelems: Size of the transfer in bytes. + :param sig_addr: Address of signal variable on the remote PE. Must be an address on the symmetric heap. + :param signal: Signal value to be written. + :param sig_op: Signal operation (ROCSHMEM_SIGNAL_SET or ROCSHMEM_SIGNAL_ADD). + :param pe: PE number of the remote PE. + :param stream: HIP stream on which to enqueue the operation. + :returns: None. + +**Description:** +This routine enqueues a put-with-signal operation on a HIP stream. The function writes contiguous +data of ``nelems`` bytes from source on the calling PE to ``dest`` at ``pe``, then applies ``sig_op`` +at ``sig_addr`` with the signal value. The operation is enqueued on the specified stream and will +execute asynchronously. The caller must synchronize the stream (e.g., using ``hipStreamSynchronize``) +to ensure completion. + +Valid ``sig_op`` values are listed in SIGNAL_OPERATORS_. + +ROCSHMEM_SIGNAL_FETCH +--------------------- + +.. cpp:function:: __device__ uint64_t rocshmem_signal_fetch(const uint64_t *sig_addr) +.. cpp:function:: __device__ uint64_t rocshmem_signal_fetch_wg(const uint64_t *sig_addr) +.. cpp:function:: __device__ uint64_t rocshmem_signal_fetch_wave(const uint64_t *sig_addr) + + :param sig_addr: Signal address. Must be an address on the symmetric heap. + :returns: Value at ``sig_addr``. + +**Description:** +This function atomically fetches the value stored at ``sig_addr``. + +Signal operators +---------------- +.. _SIGNAL_OPERATORS: + +.. list-table:: Signal Operators + :widths: 20 40 + :header-rows: 1 + + * - Value + - Description + * - ROCSHMEM_SIGNAL_SET + - The signaling operation routines will atomically set the signal value at ``sig_addr``. + * - ROCSHMEM_SIGNAL_ADD + - The signaling operation routines will atomically add the signal value at ``sig_addr``. + diff --git a/projects/rocshmem/docs/api/teams.rst b/projects/rocshmem/docs/api/teams.rst new file mode 100644 index 0000000000..f775d90421 --- /dev/null +++ b/projects/rocshmem/docs/api/teams.rst @@ -0,0 +1,90 @@ +.. meta:: + :description: rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform. + :keywords: rocSHMEM, API, ROCm, documentation, HIP, Networking, Communication + +.. _rocshmem-api-teams: + +------------------------- +Team management routines +------------------------- + +ROCSHMEM_TEAM_MY_PE +------------------- + +.. cpp:function:: __host__ int rocshmem_team_my_pe(rocshmem_team_t team) + + :param team: The team to query. + :returns: PE ID of the caller in the provided team. + +**Description:** +This routine queries the PE ID of the caller in a team. + +ROCSHMEM_TEAM_N_PES +------------------- + +.. cpp:function:: __host__ int rocshmem_team_n_pes(rocshmem_team_t team) + + :param team: The team to query. + :returns: Number of PEs in the provided team. + +**Description:** +This routine queries the number of PEs in a team. + +ROCSHMEM_TEAM_TRANSLATE_PE +-------------------------- + +.. cpp:function:: __host__ int rocshmem_team_translate_pe(rocshmem_team_t src_team, int src_pe, rocshmem_team_t dest_team) + + :param src_team: Handle of the team from which to translate. + :param src_pe: PE-of-interest's index in ``src_team``. + :param dest_team: Handle of the team to which to translate. + :returns: PE of ``src_pe`` in ``dest_team``. + If any input is invalid or if ``src_pe`` is + not in both source and destination teams, a value of ``-1`` is returned. + +**Description:** +This routine translates the PE in ``src_team`` to that in ``dest_team``. + +ROCSHMEM_TEAM_SPLIT_STRIDED +--------------------------- + +.. cpp:function:: __host__ int rocshmem_team_split_strided(rocshmem_team_t parent_team, int start, int stride, int size, const rocshmem_team_config_t *config, long config_mask, rocshmem_team_t *new_team) + + :param parent_team: The team to split from. + :param start: The lowest PE number of the subset of the PEs + from the parent team that will form the new + team. + :param stride: The stride between team PE members in the + parent team that comprise the subset of PEs + that will form the new team. + :param size: The number of PEs in the new team. + :param config: Pointer to the config parameters for the new team. + :param config_mask: Bitwise mask representing parameters to use from config. + :param new_team: Pointer to the newly created team. + If an error occurs during team creation, or if the PE in + the parent team is not in the new team, the value will be + ``ROCSHMEM_TEAM_INVALID``. + + :returns: Zero upon successful team creation; non-zero if erroneous. + +**Description:** +This routine creates a new a team of PEs. It must be called by all PEs in the parent team. + +ROCSHMEM_TEAM_DESTROY +--------------------- + +.. cpp:function:: __host__ void rocshmem_team_destroy(rocshmem_team_t team) + + :param team: The team to destroy. The behavior is undefined if + the input team is ``ROCSHMEM_TEAM_WORLD`` or any other + invalid team. If the input is ``ROCSHMEM_TEAM_INVALID``, + this function will not perform any operation. + + :returns: None + +**Description:** +This routine destroys a team. It must be called by all PEs in the team. +You must destroy all private contexts created in the +team before destroying this team. Otherwise, the behavior +is undefined. This call will destroy only the shareable contexts +created from the referenced team. diff --git a/projects/rocshmem/docs/compile_and_run.rst b/projects/rocshmem/docs/compile_and_run.rst new file mode 100644 index 0000000000..77761e4ee8 --- /dev/null +++ b/projects/rocshmem/docs/compile_and_run.rst @@ -0,0 +1,60 @@ +.. meta:: + :description: Information on how to compile and run rocSHMEM applications. + :keywords: rocSHMEM, ROCm, library, API, compile, link, hipcc + +.. _running-applications: + +-------------------------------------------------- +Compiling and running rocSHMEM applications +-------------------------------------------------- + +This topic explains how to compile and run rocSHMEM applications. + +Compiling and linking with rocSHMEM +----------------------------------- + +rocSHMEM is a library that can be statically linked to your application during compilation with ``hipcc``. For more information, see :doc:`HIPCC `. + +When compiling your application with ``hipcc``, you must include the rocSHMEM header files and the rocSHMEM library. +Because rocSHMEM depends on MPI (Message Passing Interface), you must manually add the arguments for MPI linkage instead of using ``mpicc``. + +When using ``hipcc`` directly without a build system, it's recommended to perform the compilation and linking steps separately. + +Example compile and link commands are provided at the top of the example files in the ``examples`` directory: + +.. code-block:: bash + + # Compile + hipcc -c -fgpu-rdc -x hip rocshmem_allreduce_test.cc \ + -I/opt/rocm/include \ + -I$ROCSHMEM_INSTALL_DIR/include \ + -I$OPENMPI_UCX_INSTALL_DIR/include/ + + # Link + hipcc -fgpu-rdc --hip-link rocshmem_allreduce_test.o -o rocshmem_allreduce_test \ + $ROCSHMEM_INSTALL_DIR/lib/librocshmem.a \ + $OPENMPI_UCX_INSTALL_DIR/lib/libmpi.so \ + -L/opt/rocm/lib -lamdhip64 -lhsa-runtime64 + +If your project uses CMake, see +`Using CMake with AMD ROCm `_. + +Running a rocSHMEM application +------------------------------ + +Applications using rocSHMEM typically deploy multiple processes, usually one per GPU. +The MPI launcher, for example, ``mpiexec`` with Open MPI, is used to start the required number +of processes. For example, to launch two ``getmem`` example processes (available when compiled from source): + +.. code-block:: bash + + mpiexec --map-by numa --mca pml ucx --mca osc ucx -np 2 ./build/examples/rocshmem_getmem_test + +See the `Open MPI documentation `_ for more information about ``mpiexec`` command line parameters. + +.. note:: + + Some systems may have multiple MPI installations, some of which do not + have GPU support enabled. You must use the ``mpiexec`` from the expected + MPI library, especially when using the MPI built by yourself + as part of :ref:`install-dependencies`. diff --git a/projects/rocshmem/docs/conf.py b/projects/rocshmem/docs/conf.py new file mode 100644 index 0000000000..6deef59132 --- /dev/null +++ b/projects/rocshmem/docs/conf.py @@ -0,0 +1,36 @@ +# Configuration file for the Sphinx documentation builder. +# +# This file only contains a selection of the most common options. For a full +# list see the documentation: +# https://www.sphinx-doc.org/en/master/usage/configuration.html + +import re + +from rocm_docs import ROCmDocs + +with open('../include/rocshmem/rocshmem.hpp', encoding='utf-8') as f: + match = re.search(r'constexpr char VERSION\[\] = "([0-9.]+)[^0-9.]+', f.read()) + if not match: + raise ValueError("VERSION not found!") + version_number = match[1] +left_nav_title = f"rocSHMEM {version_number} documentation" + +# for PDF output on Read the Docs +project = "rocSHMEM" +author = "Advanced Micro Devices, Inc." +copyright = "Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved." +version = version_number +release = version_number + +external_toc_path = "./sphinx/_toc.yml" + +docs_core = ROCmDocs(left_nav_title) +docs_core.run_doxygen(doxygen_root="doxygen", doxygen_path="doxygen/xml") +docs_core.setup() + +external_projects_current_project = "rocshmem" +cpp_id_attributes = ["__host__", "__global__", "__device__"] +exclude_patterns = ["README.md"] + +for sphinx_var in ROCmDocs.SPHINX_VARS: + globals()[sphinx_var] = getattr(docs_core, sphinx_var) diff --git a/projects/rocshmem/docs/doxygen/Doxyfile b/projects/rocshmem/docs/doxygen/Doxyfile new file mode 100644 index 0000000000..47a33b785b --- /dev/null +++ b/projects/rocshmem/docs/doxygen/Doxyfile @@ -0,0 +1,2452 @@ +# Doxyfile 1.8.10 + +# This file describes the settings to be used by the documentation system +# doxygen (www.doxygen.org) for a project. +# +# All text after a double hash (##) is considered a comment and is placed in +# front of the TAG it is preceding. +# +# All text after a single hash (#) is considered a comment and will be ignored. +# The format is: +# TAG = value [value, ...] +# For lists, items can also be appended using: +# TAG += value [value, ...] +# Values that contain spaces should be placed between quotes (\" \"). + +#--------------------------------------------------------------------------- +# Project related configuration options +#--------------------------------------------------------------------------- + +# This tag specifies the encoding used for all characters in the config file +# that follow. The default is UTF-8 which is also the encoding used for all text +# before the first occurrence of this tag. Doxygen uses libiconv (or the iconv +# built into libc) for the transcoding. See http://www.gnu.org/software/libiconv +# for the list of possible encodings. +# The default value is: UTF-8. + +DOXYFILE_ENCODING = UTF-8 + +# The PROJECT_NAME tag is a single word (or a sequence of words surrounded by +# double-quotes, unless you are using Doxywizard) that should identify the +# project for which the documentation is generated. This name is used in the +# title of most generated pages and in a few other places. +# The default value is: My Project. + +PROJECT_NAME = "rocSHMEM" + +# 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 = 2.0.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 = "rocSHMEM intra-kernel networking runtime for AMD dGPUs on the ROCm platform" + +# 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 +# pixels and the maximum width should not exceed 200 pixels. Doxygen will copy +# the logo to the output directory. + +PROJECT_LOGO = + +# 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 = . + +# 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. +# The default value is: NO. + +CREATE_SUBDIRS = NO + +# 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 +# U+3044. +# The default value is: NO. + +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, 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 + +# 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. +# The default value is: YES. + +BRIEF_MEMBER_DESC = YES + +# If the REPEAT_BRIEF tag is set to YES, doxygen will prepend the brief +# description of a member or function before the detailed description +# +# Note: If both HIDE_UNDOC_MEMBERS and BRIEF_MEMBER_DESC are set to NO, the +# brief descriptions will be completely suppressed. +# The default value is: YES. + +REPEAT_BRIEF = YES + +# This tag implements a quasi-intelligent brief description abbreviator that is +# used to form the text in various listings. Each string in this list, if found +# as the leading text of the brief description, will be stripped from the text +# and the result, after processing the whole list, is used as the annotated +# text. Otherwise, the brief description is used as-is. If left blank, the +# following values are used ($name is automatically replaced with the name of +# the entity):The $name class, The $name widget, The $name file, is, provides, +# specifies, contains, represents, a, an and the. + +ABBREVIATE_BRIEF = "The $name class" \ + "The $name widget" \ + "The $name file" \ + is \ + provides \ + specifies \ + contains \ + represents \ + a \ + an \ + the + +# If the ALWAYS_DETAILED_SEC and REPEAT_BRIEF tags are both set to YES then +# doxygen will generate a detailed section even if there is only a brief +# description. +# The default value is: NO. + +ALWAYS_DETAILED_SEC = NO + +# If the INLINE_INHERITED_MEMB tag is set to YES, doxygen will show all +# inherited members of a class in the documentation of that class as if those +# members were ordinary class members. Constructors, destructors and assignment +# operators of the base classes will not be shown. +# The default value is: NO. + +INLINE_INHERITED_MEMB = NO + +# If the FULL_PATH_NAMES tag is set to YES, doxygen will prepend the full path +# before files name in the file list and in the header files. If set to NO the +# shortest path that makes the file name unique will be used +# The default value is: YES. + +FULL_PATH_NAMES = YES + +# The STRIP_FROM_PATH tag can be used to strip a user-defined part of the path. +# Stripping is only done if one of the specified strings matches the left-hand +# part of the path. The tag can be used to show relative paths in the file list. +# If left blank the directory from which doxygen is run is used as the path to +# strip. +# +# Note that you can specify absolute paths here, but also relative paths, which +# will be relative from the directory where doxygen is started. +# This tag requires that the tag FULL_PATH_NAMES is set to YES. + +STRIP_FROM_PATH = + +# The STRIP_FROM_INC_PATH tag can be used to strip a user-defined part of the +# path mentioned in the documentation of a class, which tells the reader which +# header file to include in order to use a class. If left blank only the name of +# the header file containing the class definition is used. Otherwise one should +# specify the list of include paths that are normally passed to the compiler +# using the -I flag. + +STRIP_FROM_INC_PATH = + +# If the SHORT_NAMES tag is set to YES, doxygen will generate much shorter (but +# less readable) file names. This can be useful is your file systems doesn't +# support long names like on DOS, Mac, or CD-ROM. +# The default value is: NO. + +SHORT_NAMES = NO + +# If the JAVADOC_AUTOBRIEF tag is set to YES then doxygen will interpret the +# first line (until the first dot) of a Javadoc-style comment as the brief +# description. If set to NO, the Javadoc-style will behave just like regular Qt- +# style comments (thus requiring an explicit @brief command for a brief +# description.) +# The default value is: NO. + +JAVADOC_AUTOBRIEF = NO + +# If the QT_AUTOBRIEF tag is set to YES then doxygen will interpret the first +# line (until the first dot) of a Qt-style comment as the brief description. If +# set to NO, the Qt-style will behave just like regular Qt-style comments (thus +# requiring an explicit \brief command for a brief description.) +# The default value is: NO. + +QT_AUTOBRIEF = NO + +# The MULTILINE_CPP_IS_BRIEF tag can be set to YES to make doxygen treat a +# multi-line C++ special comment block (i.e. a block of //! or /// comments) as +# a brief description. This used to be the default behavior. The new default is +# to treat a multi-line C++ comment block as a detailed description. Set this +# tag to YES if you prefer the old behavior instead. +# +# Note that setting this tag to YES also means that rational rose comments are +# not recognized any more. +# The default value is: NO. + +MULTILINE_CPP_IS_BRIEF = NO + +# If the INHERIT_DOCS tag is set to YES then an undocumented member inherits the +# documentation from any documented member that it re-implements. +# The default value is: YES. + +INHERIT_DOCS = YES + +# If the SEPARATE_MEMBER_PAGES tag is set to YES then doxygen will produce a new +# page for each member. If set to NO, the documentation of a member will be part +# of the file/class/namespace that contains it. +# The default value is: NO. + +SEPARATE_MEMBER_PAGES = NO + +# The TAB_SIZE tag can be used to set the number of spaces in a tab. Doxygen +# uses this value to replace tabs by spaces in code fragments. +# Minimum value: 1, maximum value: 16, default value: 4. + +TAB_SIZE = 4 + +# This tag can be used to specify a number of aliases that act as commands in +# the documentation. An alias has the form: +# name=value +# For example adding +# "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:". You can put \n's in the value part of an alias to insert +# newlines. + +ALIASES = + +# This tag can be used to specify a number of word-keyword mappings (TCL only). +# A mapping has the form "name=value". For example adding "class=itcl::class" +# will allow you to use the command class in the itcl::class meaning. + +TCL_SUBST = + +# Set the OPTIMIZE_OUTPUT_FOR_C tag to YES if your project consists of C sources +# only. Doxygen will then generate output that is more tailored for C. For +# instance, some of the names that are used will be different. The list of all +# members will be omitted, etc. +# The default value is: NO. + +OPTIMIZE_OUTPUT_FOR_C = NO + +# Set the OPTIMIZE_OUTPUT_JAVA tag to YES if your project consists of Java or +# Python sources only. Doxygen will then generate output that is more tailored +# for that language. For instance, namespaces will be presented as packages, +# qualified scopes will look different, etc. +# The default value is: NO. + +OPTIMIZE_OUTPUT_JAVA = NO + +# Set the OPTIMIZE_FOR_FORTRAN tag to YES if your project consists of Fortran +# sources. Doxygen will then generate output that is tailored for Fortran. +# The default value is: NO. + +OPTIMIZE_FOR_FORTRAN = NO + +# Set the OPTIMIZE_OUTPUT_VHDL tag to YES if your project consists of VHDL +# sources. Doxygen will then generate output that is tailored for VHDL. +# The default value is: NO. + +OPTIMIZE_OUTPUT_VHDL = NO + +# Doxygen selects the parser to use depending on the extension of the files it +# parses. With this tag you can assign which parser to use for a given +# 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, +# C#, C, C++, D, PHP, Objective-C, Python, 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), VHDL. For +# instance to make doxygen treat .inc files as Fortran files (default is PHP), +# and .f files as C (default is Fortran), use: inc=Fortran f=C. +# +# Note: For files without extension you can use no_extension as a placeholder. +# +# Note that for custom extensions you also need to set FILE_PATTERNS otherwise +# the files are not read by doxygen. + +EXTENSION_MAPPING = + +# If the MARKDOWN_SUPPORT tag is enabled then doxygen pre-processes all comments +# according to the Markdown format, which allows for more readable +# documentation. See http://daringfireball.net/projects/markdown/ for details. +# The output of markdown processing is further processed by doxygen, so you can +# mix doxygen, HTML, and XML commands with Markdown formatting. Disable only in +# case of backward compatibilities issues. +# The default value is: YES. + +MARKDOWN_SUPPORT = YES + +# 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 +# globally by setting AUTOLINK_SUPPORT to NO. +# The default value is: YES. + +AUTOLINK_SUPPORT = YES + +# If you use STL classes (i.e. std::string, std::vector, etc.) but do not want +# to include (a tag file for) the STL sources as input, then you should set this +# tag to YES in order to let doxygen match functions declarations and +# definitions whose arguments contain STL classes (e.g. func(std::string); +# versus func(std::string) {}). This also make the inheritance and collaboration +# diagrams that involve STL classes more complete and accurate. +# The default value is: NO. + +BUILTIN_STL_SUPPORT = YES + +# If you use Microsoft's C++/CLI language, you should set this option to YES to +# enable parsing support. +# The default value is: NO. + +CPP_CLI_SUPPORT = NO + +# Set the SIP_SUPPORT tag to YES if your project consists of sip (see: +# http://www.riverbankcomputing.co.uk/software/sip/intro) sources only. Doxygen +# will parse them like normal C++ but will assume all classes use public instead +# of private inheritance when no explicit protection keyword is present. +# The default value is: NO. + +SIP_SUPPORT = NO + +# For Microsoft's IDL there are propget and propput attributes to indicate +# getter and setter methods for a property. Setting this option to YES will make +# doxygen to replace the get and set methods by a property in the documentation. +# This will only work if the methods are indeed getting or setting a simple +# type. If this is not the case, or you want to show the methods anyway, you +# should set this option to NO. +# The default value is: YES. + +IDL_PROPERTY_SUPPORT = YES + +# If member grouping is used in the documentation and the DISTRIBUTE_GROUP_DOC +# tag is set to YES then doxygen will reuse the documentation of the first +# member in the group (if any) for the other members of the group. By default +# all members of a group must be documented explicitly. +# The default value is: NO. + +DISTRIBUTE_GROUP_DOC = YES + +# If one adds a struct or class to a group and this option is enabled, then also +# any nested class or struct is added to the same group. By default this option +# is disabled and one has to add nested compounds explicitly via \ingroup. +# The default value is: NO. + +GROUP_NESTED_COMPOUNDS = NO + +# Set the SUBGROUPING tag to YES to allow class member groups of the same type +# (for instance a group of public functions) to be put as a subgroup of that +# type (e.g. under the Public Functions section). Set it to NO to prevent +# subgrouping. Alternatively, this can be done per class using the +# \nosubgrouping command. +# The default value is: YES. + +SUBGROUPING = YES + +# When the INLINE_GROUPED_CLASSES tag is set to YES, classes, structs and unions +# are shown inside the group in which they are included (e.g. using \ingroup) +# instead of on a separate page (for HTML and Man pages) or section (for LaTeX +# and RTF). +# +# Note that this feature does not work in combination with +# SEPARATE_MEMBER_PAGES. +# The default value is: NO. + +INLINE_GROUPED_CLASSES = NO + +# When the INLINE_SIMPLE_STRUCTS tag is set to YES, structs, classes, and unions +# with only public data fields or simple typedef fields will be shown inline in +# the documentation of the scope in which they are defined (i.e. file, +# namespace, or group documentation), provided this scope is documented. If set +# to NO, structs, classes, and unions are shown on a separate page (for HTML and +# Man pages) or section (for LaTeX and RTF). +# The default value is: NO. + +INLINE_SIMPLE_STRUCTS = NO + +# When TYPEDEF_HIDES_STRUCT tag is enabled, a typedef of a struct, union, or +# enum is documented as struct, union, or enum with the name of the typedef. So +# typedef struct TypeS {} TypeT, will appear in the documentation as a struct +# with name TypeT. When disabled the typedef will appear as a member of a file, +# namespace, or class. And the struct will be named TypeS. This can typically be +# useful for C code in case the coding convention dictates that all compound +# types are typedef'ed and only the typedef is referenced, never the tag name. +# The default value is: NO. + +TYPEDEF_HIDES_STRUCT = YES + +# The size of the symbol lookup cache can be set using LOOKUP_CACHE_SIZE. This +# cache is used to resolve symbols given their name and scope. Since this can be +# an expensive process and often the same symbol appears multiple times in the +# code, doxygen keeps a cache of pre-resolved symbols. If the cache is too small +# doxygen will become slower. If the cache is too large, memory is wasted. The +# cache size is given by this formula: 2^(16+LOOKUP_CACHE_SIZE). The valid range +# is 0..9, the default is 0, corresponding to a cache size of 2^16=65536 +# symbols. At the end of a run doxygen will report the cache usage and suggest +# the optimal cache size from a speed point of view. +# Minimum value: 0, maximum value: 9, default value: 0. + +LOOKUP_CACHE_SIZE = 0 + +#--------------------------------------------------------------------------- +# Build related configuration options +#--------------------------------------------------------------------------- + +# If the EXTRACT_ALL tag is set to YES, doxygen will assume all entities in +# documentation are documented, even if no documentation was available. Private +# class members and static file members will be hidden unless the +# EXTRACT_PRIVATE respectively EXTRACT_STATIC tags are set to YES. +# Note: This will also disable the warnings about undocumented members that are +# normally produced when WARNINGS is set to YES. +# The default value is: NO. + +EXTRACT_ALL = YES + +# If the EXTRACT_PRIVATE tag is set to YES, all private members of a class will +# be included in the documentation. +# The default value is: NO. + +EXTRACT_PRIVATE = NO + +# If the EXTRACT_PACKAGE tag is set to YES, all members with package or internal +# scope will be included in the documentation. +# The default value is: NO. + +EXTRACT_PACKAGE = NO + +# If the EXTRACT_STATIC tag is set to YES, all static members of a file will be +# included in the documentation. +# The default value is: NO. + +EXTRACT_STATIC = NO + +# If the EXTRACT_LOCAL_CLASSES tag is set to YES, classes (and structs) defined +# locally in source files will be included in the documentation. If set to NO, +# only classes defined in header files are included. Does not have any effect +# for Java sources. +# The default value is: YES. + +EXTRACT_LOCAL_CLASSES = YES + +# This flag is only useful for Objective-C code. If set to YES, local methods, +# which are defined in the implementation section but not in the interface are +# included in the documentation. If set to NO, only methods in the interface are +# included. +# The default value is: NO. + +EXTRACT_LOCAL_METHODS = NO + +# If this flag is set to YES, the members of anonymous namespaces will be +# extracted and appear in the documentation as a namespace called +# 'anonymous_namespace{file}', where file will be replaced with the base name of +# the file that contains the anonymous namespace. By default anonymous namespace +# are hidden. +# The default value is: NO. + +EXTRACT_ANON_NSPACES = NO + +# If the HIDE_UNDOC_MEMBERS tag is set to YES, doxygen will hide all +# undocumented members inside documented classes or files. If set to NO these +# members will be included in the various overviews, but no documentation +# section is generated. This option has no effect if EXTRACT_ALL is enabled. +# The default value is: NO. + +HIDE_UNDOC_MEMBERS = NO + +# 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 +# has no effect if EXTRACT_ALL is enabled. +# The default value is: NO. + +HIDE_UNDOC_CLASSES = NO + +# If the HIDE_FRIEND_COMPOUNDS tag is set to YES, doxygen will hide all friend +# (class|struct|union) declarations. If set to NO, these declarations will be +# included in the documentation. +# The default value is: NO. + +HIDE_FRIEND_COMPOUNDS = NO + +# If the HIDE_IN_BODY_DOCS tag is set to YES, doxygen will hide any +# documentation blocks found inside the body of a function. If set to NO, these +# blocks will be appended to the function's detailed documentation block. +# The default value is: NO. + +HIDE_IN_BODY_DOCS = NO + +# The INTERNAL_DOCS tag determines if documentation that is typed after a +# \internal command is included. If the tag is set to NO then the documentation +# will be excluded. Set it to YES to include the internal documentation. +# The default value is: NO. + +INTERNAL_DOCS = NO + +# If the CASE_SENSE_NAMES tag is set to NO then doxygen will only generate file +# names in lower-case letters. If set to YES, upper-case letters are also +# allowed. This is useful if you have classes or files whose names only differ +# in case and if your file system supports case sensitive file names. Windows +# and Mac users are advised to set this option to NO. +# The default value is: system dependent. + +CASE_SENSE_NAMES = NO + +# 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 +# scope will be hidden. +# The default value is: NO. + +HIDE_SCOPE_NAMES = NO + +# If the HIDE_COMPOUND_REFERENCE tag is set to NO (default) then doxygen will +# append additional text to a page's title, such as Class Reference. If set to +# YES the compound reference will be hidden. +# The default value is: NO. + +HIDE_COMPOUND_REFERENCE= NO + +# 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. + +SHOW_INCLUDE_FILES = YES + +# If the SHOW_GROUPED_MEMB_INC tag is set to YES then Doxygen will add for each +# grouped member an include statement to the documentation, telling the reader +# which file to include in order to use the member. +# The default value is: NO. + +SHOW_GROUPED_MEMB_INC = NO + +# If the FORCE_LOCAL_INCLUDES tag is set to YES then doxygen will list include +# files with double quotes in the documentation rather than with sharp brackets. +# The default value is: NO. + +FORCE_LOCAL_INCLUDES = NO + +# If the INLINE_INFO tag is set to YES then a tag [inline] is inserted in the +# documentation for inline members. +# The default value is: YES. + +INLINE_INFO = YES + +# If the SORT_MEMBER_DOCS tag is set to YES then doxygen will sort the +# (detailed) documentation of file and class members alphabetically by member +# name. If set to NO, the members will appear in declaration order. +# The default value is: YES. + +SORT_MEMBER_DOCS = YES + +# If the SORT_BRIEF_DOCS tag is set to YES then doxygen will sort the brief +# descriptions of file, namespace and class members alphabetically by member +# name. If set to NO, the members will appear in declaration order. Note that +# this will also influence the order of the classes in the class list. +# The default value is: NO. + +SORT_BRIEF_DOCS = NO + +# If the SORT_MEMBERS_CTORS_1ST tag is set to YES then doxygen will sort the +# (brief and detailed) documentation of class members so that constructors and +# destructors are listed first. If set to NO the constructors will appear in the +# respective orders defined by SORT_BRIEF_DOCS and SORT_MEMBER_DOCS. +# Note: If SORT_BRIEF_DOCS is set to NO this option is ignored for sorting brief +# member documentation. +# Note: If SORT_MEMBER_DOCS is set to NO this option is ignored for sorting +# detailed member documentation. +# The default value is: NO. + +SORT_MEMBERS_CTORS_1ST = NO + +# If the SORT_GROUP_NAMES tag is set to YES then doxygen will sort the hierarchy +# of group names into alphabetical order. If set to NO the group names will +# appear in their defined order. +# The default value is: NO. + +SORT_GROUP_NAMES = NO + +# If the SORT_BY_SCOPE_NAME tag is set to YES, the class list will be sorted by +# fully-qualified names, including namespaces. If set to NO, the class list will +# be sorted only by class name, not including the namespace part. +# Note: This option is not very useful if HIDE_SCOPE_NAMES is set to YES. +# Note: This option applies only to the class list, not to the alphabetical +# list. +# The default value is: NO. + +SORT_BY_SCOPE_NAME = NO + +# If the STRICT_PROTO_MATCHING option is enabled and doxygen fails to do proper +# type resolution of all parameters of a function it will reject a match between +# the prototype and the implementation of a member function even if there is +# only one candidate or it is obvious which candidate to choose by doing a +# simple string match. By disabling STRICT_PROTO_MATCHING doxygen will still +# accept a match between prototype and implementation in such cases. +# The default value is: NO. + +STRICT_PROTO_MATCHING = NO + +# The GENERATE_TODOLIST tag can be used to enable (YES) or disable (NO) the todo +# list. This list is created by putting \todo commands in the documentation. +# The default value is: YES. + +GENERATE_TODOLIST = YES + +# The GENERATE_TESTLIST tag can be used to enable (YES) or disable (NO) the test +# list. This list is created by putting \test commands in the documentation. +# The default value is: YES. + +GENERATE_TESTLIST = YES + +# The GENERATE_BUGLIST tag can be used to enable (YES) or disable (NO) the bug +# list. This list is created by putting \bug commands in the documentation. +# The default value is: YES. + +GENERATE_BUGLIST = YES + +# The GENERATE_DEPRECATEDLIST tag can be used to enable (YES) or disable (NO) +# the deprecated list. This list is created by putting \deprecated commands in +# the documentation. +# The default value is: YES. + +GENERATE_DEPRECATEDLIST= YES + +# The ENABLED_SECTIONS tag can be used to enable conditional documentation +# sections, marked by \if ... \endif and \cond +# ... \endcond blocks. + +ENABLED_SECTIONS = + +# The MAX_INITIALIZER_LINES tag determines the maximum number of lines that the +# initial value of a variable or macro / define can have for it to appear in the +# documentation. If the initializer consists of more lines than specified here +# it will be hidden. Use a value of 0 to hide initializers completely. The +# appearance of the value of individual variables and macros / defines can be +# controlled using \showinitializer or \hideinitializer command in the +# documentation regardless of this setting. +# Minimum value: 0, maximum value: 10000, default value: 30. + +MAX_INITIALIZER_LINES = 30 + +# Set the SHOW_USED_FILES tag to NO to disable the list of files generated at +# the bottom of the documentation of classes and structs. If set to YES, the +# list will mention the files that were used to generate the documentation. +# The default value is: YES. + +SHOW_USED_FILES = YES + +# Set the SHOW_FILES tag to NO to disable the generation of the Files page. This +# will remove the Files entry from the Quick Index and from the Folder Tree View +# (if specified). +# The default value is: YES. + +SHOW_FILES = YES + +# Set the SHOW_NAMESPACES tag to NO to disable the generation of the Namespaces +# page. This will remove the Namespaces entry from the Quick Index and from the +# Folder Tree View (if specified). +# The default value is: YES. + +SHOW_NAMESPACES = YES + +# The FILE_VERSION_FILTER tag can be used to specify a program or script that +# doxygen should invoke to get the current version for each file (typically from +# the version control system). Doxygen will invoke the program by executing (via +# popen()) the command command input-file, where command is the value of the +# FILE_VERSION_FILTER tag, and input-file is the name of an input file provided +# by doxygen. Whatever the program writes to standard output is used as the file +# version. For an example see the documentation. + +FILE_VERSION_FILTER = + +# The LAYOUT_FILE tag can be used to specify a layout file which will be parsed +# by doxygen. The layout file controls the global structure of the generated +# 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. +# +# 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 +# tag is left empty. + +LAYOUT_FILE = + +# The CITE_BIB_FILES tag can be used to specify one or more bib files containing +# the reference definitions. This must be a list of .bib files. The .bib +# extension is automatically appended if omitted. This requires the bibtex tool +# to be installed. See also http://en.wikipedia.org/wiki/BibTeX for more info. +# For LaTeX the style of the bibliography can be controlled using +# LATEX_BIB_STYLE. To use this feature you need bibtex and perl available in the +# search path. See also \cite for info how to create references. + +CITE_BIB_FILES = + +#--------------------------------------------------------------------------- +# Configuration options related to warning and progress messages +#--------------------------------------------------------------------------- + +# The QUIET tag can be used to turn on/off the messages that are generated to +# standard output by doxygen. If QUIET is set to YES this implies that the +# messages are off. +# The default value is: NO. + +QUIET = NO + +# The WARNINGS tag can be used to turn on/off the warning messages that are +# generated to standard error (stderr) by doxygen. If WARNINGS is set to YES +# this implies that the warnings are on. +# +# Tip: Turn warnings on while writing the documentation. +# The default value is: YES. + +WARNINGS = YES + +# If the WARN_IF_UNDOCUMENTED tag is set to YES then doxygen will generate +# warnings for undocumented members. If EXTRACT_ALL is set to YES then this flag +# will automatically be disabled. +# The default value is: 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 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 + +# 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 or incomplete +# parameter documentation, but not about the absence of documentation. +# The default value is: NO. + +WARN_NO_PARAMDOC = NO + +# The WARN_FORMAT tag determines the format of the warning messages that doxygen +# can produce. The string should contain the $file, $line, and $text tags, which +# will be replaced by the file and line number from which the warning originated +# 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) +# The default value is: $file:$line: $text. + +WARN_FORMAT = "$file:$line: $text" + +# 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). + +WARN_LOGFILE = + +#--------------------------------------------------------------------------- +# Configuration options related to the input files +#--------------------------------------------------------------------------- + +# The INPUT tag is used to specify the files and/or directories that contain +# documented source files. You may enter file names like myfile.cpp or +# directories like /usr/src/myproject. Separate the files or directories with +# spaces. See also FILE_PATTERNS and EXTENSION_MAPPING +# Note: If this tag is empty the current directory is searched. + +INPUT = ../../include + +# 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: http://www.gnu.org/software/libiconv) for the list of +# possible encodings. +# The default value is: UTF-8. + +INPUT_ENCODING = UTF-8 + +# 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. +# +# 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 +# read by doxygen. +# +# 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, *.py, *.f90, *.f, *.for, *.tcl, *.vhd, +# *.vhdl, *.ucf, *.qsf, *.as and *.js. + +FILE_PATTERNS = *.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 \ + *.py \ + *.tcl \ + *.vhd \ + *.vhdl \ + *.ucf \ + *.qsf \ + *.as \ + *.js + +# The RECURSIVE tag can be used to specify whether or not subdirectories should +# be searched for input files as well. +# The default value is: NO. + +RECURSIVE = NO + +# The EXCLUDE tag can be used to specify files and/or directories that should be +# excluded from the INPUT source files. This way you can easily exclude a +# subdirectory from a directory tree whose root is specified with the INPUT tag. +# +# Note that relative paths are relative to the directory from which doxygen is +# run. + +EXCLUDE = + +# The EXCLUDE_SYMLINKS tag can be used to select whether or not files or +# directories that are symbolic links (a Unix file system feature) are excluded +# from the input. +# The default value is: NO. + +EXCLUDE_SYMLINKS = NO + +# If the value of the INPUT tag contains directories, you can use the +# EXCLUDE_PATTERNS tag to specify one or more wildcard patterns to exclude +# certain files from those directories. +# +# Note that the wildcards are matched against the file with absolute path, so to +# exclude all test directories for example use the pattern */test/* + +EXCLUDE_PATTERNS = + +# The EXCLUDE_SYMBOLS tag can be used to specify one or more symbol names +# (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, +# 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 = + +# The EXAMPLE_PATH tag can be used to specify one or more files or directories +# that contain example code fragments that are included (see the \include +# command). + +EXAMPLE_PATH = + +# If the value of the EXAMPLE_PATH tag contains directories, you can use the +# EXAMPLE_PATTERNS tag to specify one or more wildcard pattern (like *.cpp and +# *.h) to filter out the source-files in the directories. If left blank all +# files are included. + +EXAMPLE_PATTERNS = * + +# If the EXAMPLE_RECURSIVE tag is set to YES then subdirectories will be +# searched for input files to be used with the \include or \dontinclude commands +# irrespective of the value of the RECURSIVE tag. +# The default value is: NO. + +EXAMPLE_RECURSIVE = NO + +# The IMAGE_PATH tag can be used to specify one or more files or directories +# that contain images that are to be included in the documentation (see the +# \image command). + +IMAGE_PATH = + +# The INPUT_FILTER tag can be used to specify a program that doxygen should +# invoke to filter for each input file. Doxygen will invoke the filter program +# by executing (via popen()) the command: +# +# +# +# where is the value of the INPUT_FILTER tag, and is the +# name of an input file. Doxygen will then use the output that the filter +# program writes to standard output. If FILTER_PATTERNS is specified, this tag +# will be ignored. +# +# Note that the filter must not add or remove lines; it is applied before the +# code is scanned, but not when the output code is generated. If lines are added +# or removed, the anchors will not be placed correctly. + +INPUT_FILTER = + +# The FILTER_PATTERNS tag can be used to specify filters on a per file pattern +# basis. Doxygen will compare the file name with each pattern and apply the +# filter if there is a match. The filters are a list of the form: pattern=filter +# (like *.cpp=my_cpp_filter). See INPUT_FILTER for further information on how +# filters are used. If the FILTER_PATTERNS tag is empty or if none of the +# patterns match the file name, INPUT_FILTER is applied. + +FILTER_PATTERNS = + +# If the FILTER_SOURCE_FILES tag is set to YES, the input filter (if set using +# INPUT_FILTER) will also be used to filter the input files that are used for +# producing the source files to browse (i.e. when SOURCE_BROWSER is set to YES). +# The default value is: NO. + +FILTER_SOURCE_FILES = NO + +# The FILTER_SOURCE_PATTERNS tag can be used to specify source filters per file +# pattern. A pattern will override the setting for FILTER_PATTERN (if any) and +# it is also possible to disable source filtering for a specific pattern using +# *.ext= (so without naming a filter). +# This tag requires that the tag FILTER_SOURCE_FILES is set to YES. + +FILTER_SOURCE_PATTERNS = + +# If the USE_MDFILE_AS_MAINPAGE tag refers to the name of a markdown file that +# is part of the input, its contents will be placed on the main page +# (index.html). This can be useful if you have a project on for instance GitHub +# and want to reuse the introduction page also for the doxygen output. + +USE_MDFILE_AS_MAINPAGE = ../../README.md + +#--------------------------------------------------------------------------- +# Configuration options related to source browsing +#--------------------------------------------------------------------------- + +# If the SOURCE_BROWSER tag is set to YES then a list of source files will be +# generated. Documented entities will be cross-referenced with these sources. +# +# Note: To get rid of all source code in the generated output, make sure that +# also VERBATIM_HEADERS is set to NO. +# The default value is: NO. + +SOURCE_BROWSER = NO + +# Setting the INLINE_SOURCES tag to YES will include the body of functions, +# classes and enums directly into the documentation. +# The default value is: NO. + +INLINE_SOURCES = NO + +# Setting the STRIP_CODE_COMMENTS tag to YES will instruct doxygen to hide any +# special comment blocks from generated source code fragments. Normal C, C++ and +# Fortran comments will always remain visible. +# The default value is: YES. + +STRIP_CODE_COMMENTS = YES + +# If the REFERENCED_BY_RELATION tag is set to YES then for each documented +# function all documented functions referencing it will be listed. +# The default value is: NO. + +REFERENCED_BY_RELATION = NO + +# If the REFERENCES_RELATION tag is set to YES then for each documented function +# all documented entities called/used by that function will be listed. +# The default value is: NO. + +REFERENCES_RELATION = NO + +# If the REFERENCES_LINK_SOURCE tag is set to YES and SOURCE_BROWSER tag is set +# to YES then the hyperlinks from functions in REFERENCES_RELATION and +# REFERENCED_BY_RELATION lists will link to the source code. Otherwise they will +# link to the documentation. +# The default value is: YES. + +REFERENCES_LINK_SOURCE = YES + +# If SOURCE_TOOLTIPS is enabled (the default) then hovering a hyperlink in the +# source code will show a tooltip with additional information such as prototype, +# brief description and links to the definition and documentation. Since this +# will make the HTML file larger and loading of large files a bit slower, you +# can opt to disable this feature. +# The default value is: YES. +# This tag requires that the tag SOURCE_BROWSER is set to YES. + +SOURCE_TOOLTIPS = YES + +# If the USE_HTAGS tag is set to YES then the references to source code will +# point to the HTML generated by the htags(1) tool instead of doxygen built-in +# source browser. The htags tool is part of GNU's global source tagging system +# (see http://www.gnu.org/software/global/global.html). You will need version +# 4.8.6 or higher. +# +# To use it do the following: +# - Install the latest version of global +# - Enable SOURCE_BROWSER and USE_HTAGS in the config file +# - Make sure the INPUT points to the root of the source tree +# - Run doxygen as normal +# +# Doxygen will invoke htags (and that will in turn invoke gtags), so these +# tools must be available from the command line (i.e. in the search path). +# +# The result: instead of the source browser generated by doxygen, the links to +# source code will now point to the output of htags. +# The default value is: NO. +# This tag requires that the tag SOURCE_BROWSER is set to YES. + +USE_HTAGS = NO + +# If the VERBATIM_HEADERS tag is set the YES then doxygen will generate a +# verbatim copy of the header file for each class for which an include is +# specified. Set to NO to disable this. +# See also: Section \class. +# The default value is: YES. + +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 +# compiled with the --with-libclang option. +# The default value is: NO. + +CLANG_ASSISTED_PARSING = NO + +# 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 = + +#--------------------------------------------------------------------------- +# Configuration options related to the alphabetical class index +#--------------------------------------------------------------------------- + +# If the ALPHABETICAL_INDEX tag is set to YES, an alphabetical index of all +# compounds will be generated. Enable this if the project contains a lot of +# classes, structs, unions or interfaces. +# The default value is: YES. + +ALPHABETICAL_INDEX = YES + +# The COLS_IN_ALPHA_INDEX tag can be used to specify the number of columns in +# which the alphabetical index list will be split. +# Minimum value: 1, maximum value: 20, default value: 5. +# This tag requires that the tag ALPHABETICAL_INDEX is set to YES. + +COLS_IN_ALPHA_INDEX = 5 + +# 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 = + +#--------------------------------------------------------------------------- +# Configuration options related to the HTML output +#--------------------------------------------------------------------------- + +# If the GENERATE_HTML tag is set to YES, doxygen will generate HTML output +# The default value is: YES. + +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 +# it. +# The default directory is: html. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_OUTPUT = html + +# The HTML_FILE_EXTENSION tag can be used to specify the file extension for each +# generated HTML page (for example: .htm, .php, .asp). +# The default value is: .html. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_FILE_EXTENSION = .html + +# The HTML_HEADER tag can be used to specify a user-defined HTML header file for +# each generated HTML page. If the tag is left blank doxygen will generate a +# standard header. +# +# To get valid HTML the header file that includes any scripts and style sheets +# that doxygen needs, which is dependent on the configuration options used (e.g. +# the setting GENERATE_TREEVIEW). It is highly recommended to start with a +# default header using +# doxygen -w html new_header.html new_footer.html new_stylesheet.css +# YourConfigFile +# and then modify the file new_header.html. See also section "Doxygen usage" +# for information on how to generate the default header that doxygen normally +# uses. +# Note: The header is subject to change so you typically have to regenerate the +# default header when upgrading to a newer version of doxygen. For a description +# of the possible markers and block names see the documentation. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_HEADER = + +# The HTML_FOOTER tag can be used to specify a user-defined HTML footer for each +# generated HTML page. If the tag is left blank doxygen will generate a standard +# footer. See HTML_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. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_FOOTER = + +# The HTML_STYLESHEET tag can be used to specify a user-defined cascading style +# sheet that is used by each HTML page. It can be used to fine-tune the look of +# the HTML output. If left blank doxygen will generate a default style sheet. +# See also section "Doxygen usage" for information on how to generate the style +# sheet that doxygen normally uses. +# Note: It is recommended to use HTML_EXTRA_STYLESHEET instead of this tag, as +# it is more robust and this tag (HTML_STYLESHEET) will in the future become +# obsolete. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_STYLESHEET = + +# The HTML_EXTRA_STYLESHEET tag can be used to specify additional user-defined +# cascading style sheets that are included after the standard style sheets +# created by doxygen. Using this option one can overrule certain style aspects. +# This is preferred over using HTML_STYLESHEET since it does not replace the +# standard style sheet and is therefore more robust against future updates. +# 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). For an example see the documentation. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_EXTRA_STYLESHEET = + +# 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 +# that these files will be copied to the base HTML output directory. Use the +# $relpath^ marker in the HTML_HEADER and/or HTML_FOOTER files to load these +# files. In the HTML_STYLESHEET file, use the file name only. Also note that the +# files will be copied as-is; there are no commands or markers available. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_EXTRA_FILES = + +# 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 colorwheel, see +# http://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. +# Minimum value: 0, maximum value: 359, default value: 220. +# This tag requires that the tag GENERATE_HTML is set to YES. + +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 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. + +HTML_COLORSTYLE_SAT = 100 + +# The HTML_COLORSTYLE_GAMMA tag controls the gamma correction applied to the +# luminance component of the colors in the HTML output. Values below 100 +# gradually make the output lighter, whereas values above 100 make the output +# darker. The value divided by 100 is the actual gamma applied, so 80 represents +# a gamma of 0.8, The value 220 represents a gamma of 2.2, and 100 does not +# change the gamma. +# Minimum value: 40, maximum value: 240, default value: 80. +# This tag requires that the tag GENERATE_HTML is set to YES. + +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_SECTIONS tag is set to YES then the generated HTML +# documentation will contain sections that can be hidden and shown after the +# page has loaded. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_DYNAMIC_SECTIONS = NO + +# 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 +# such a level that at most the specified number of entries are visible (unless +# a fully collapsed tree already exceeds this amount). So setting the number of +# entries 1 will produce a full collapsed tree by default. 0 is a special value +# representing an infinite number of entries and will result in a full expanded +# tree by default. +# Minimum value: 0, maximum value: 9999, default value: 100. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_INDEX_NUM_ENTRIES = 100 + +# If the GENERATE_DOCSET tag is set to YES, additional index files will be +# generated that can be used as input for Apple's Xcode 3 integrated development +# environment (see: http://developer.apple.com/tools/xcode/), introduced with +# OSX 10.5 (Leopard). To create a documentation set, doxygen will generate a +# Makefile in the HTML output directory. Running make will produce the docset in +# that directory and running make install will install the docset in +# ~/Library/Developer/Shared/Documentation/DocSets so that Xcode will find it at +# startup. See http://developer.apple.com/tools/creatingdocsetswithdoxygen.html +# for more information. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_DOCSET = NO + +# This tag determines the name 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. +# The default value is: Doxygen generated docs. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +DOCSET_FEEDNAME = "Doxygen generated docs" + +# 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. +# The default value is: org.doxygen.Project. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +DOCSET_BUNDLE_ID = org.doxygen.Project + +# The DOCSET_PUBLISHER_ID tag specifies a string that should uniquely identify +# the documentation publisher. This should be a reverse domain-name style +# string, e.g. com.mycompany.MyDocSet.documentation. +# The default value is: org.doxygen.Publisher. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +DOCSET_PUBLISHER_ID = org.doxygen.Publisher + +# The DOCSET_PUBLISHER_NAME tag identifies the documentation publisher. +# The default value is: Publisher. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +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 +# (see: http://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 +# files are now used as the Windows 98 help format, and will replace the old +# Windows help format (.hlp) on all Windows platforms in the future. Compressed +# HTML files also contain an index, a table of contents, and you can search for +# words in the documentation. The HTML workshop also contains a viewer for +# compressed HTML files. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_HTMLHELP = NO + +# The CHM_FILE tag can be used to specify the file name of the resulting .chm +# file. You can add a path in front of the file if the result should not be +# written to the html output directory. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +CHM_FILE = + +# The HHC_LOCATION tag can be used to specify the location (absolute path +# including file name) of the HTML help compiler (hhc.exe). If non-empty, +# doxygen will try to run the HTML help compiler on the generated index.hhp. +# The file has to be specified with full path. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +HHC_LOCATION = + +# The GENERATE_CHI flag controls if a separate .chi index file is generated +# (YES) or that it should be included in the master .chm file (NO). +# The default value is: NO. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +GENERATE_CHI = NO + +# The CHM_INDEX_ENCODING is used to encode HtmlHelp index (hhk), content (hhc) +# and project file content. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +CHM_INDEX_ENCODING = + +# The BINARY_TOC flag controls whether a binary table of contents is generated +# (YES) or a normal table of contents (NO) in the .chm file. Furthermore it +# enables the Previous and Next buttons. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +BINARY_TOC = NO + +# The TOC_EXPAND flag can be set to YES to add extra items for group members to +# the table of contents of the HTML help documentation and to the tree view. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +TOC_EXPAND = NO + +# 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 +# (.qch) of the generated HTML documentation. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_QHP = NO + +# If the QHG_LOCATION tag is specified, the QCH_FILE tag can be used to specify +# the file name of the resulting .qch file. The path specified is relative to +# the HTML output folder. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QCH_FILE = + +# The QHP_NAMESPACE tag specifies the namespace to use when generating Qt Help +# Project output. For more information please see Qt Help Project / Namespace +# (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#namespace). +# The default value is: org.doxygen.Project. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_NAMESPACE = org.doxygen.Project + +# The QHP_VIRTUAL_FOLDER tag specifies the namespace to use when generating Qt +# Help Project output. For more information please see Qt Help Project / Virtual +# Folders (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#virtual- +# folders). +# The default value is: doc. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_VIRTUAL_FOLDER = doc + +# If the QHP_CUST_FILTER_NAME tag is set, it specifies the name of a custom +# filter to add. For more information please see Qt Help Project / Custom +# Filters (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#custom- +# filters). +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_CUST_FILTER_NAME = + +# The QHP_CUST_FILTER_ATTRS tag specifies the list of the attributes of the +# custom filter to add. For more information please see Qt Help Project / Custom +# Filters (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#custom- +# filters). +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_CUST_FILTER_ATTRS = + +# The QHP_SECT_FILTER_ATTRS tag specifies the list of the attributes this +# project's filter section matches. Qt Help Project / Filter Attributes (see: +# http://qt-project.org/doc/qt-4.8/qthelpproject.html#filter-attributes). +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_SECT_FILTER_ATTRS = + +# The QHG_LOCATION tag can be used to specify the location of Qt's +# qhelpgenerator. If non-empty doxygen will try to run qhelpgenerator on the +# generated .qhp file. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHG_LOCATION = + +# If the GENERATE_ECLIPSEHELP tag is set to YES, additional index files will be +# generated, together with the HTML files, they form an Eclipse help plugin. To +# install this plugin and make it available under the help contents menu in +# Eclipse, the contents of the directory containing the HTML and XML files needs +# to be copied into the plugins directory of eclipse. The name of the directory +# within the plugins directory should be the same as the ECLIPSE_DOC_ID value. +# After copying Eclipse needs to be restarted before the help appears. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_ECLIPSEHELP = NO + +# A unique identifier for the Eclipse help plugin. When installing the plugin +# the directory name containing the HTML and XML files should also have this +# name. Each documentation set should have its own identifier. +# The default value is: org.doxygen.Project. +# This tag requires that the tag GENERATE_ECLIPSEHELP is set to YES. + +ECLIPSE_DOC_ID = org.doxygen.Project + +# If you want full control over the layout of the generated HTML pages it might +# be necessary to disable the index and replace it with your own. The +# DISABLE_INDEX tag can be used to turn on/off the condensed index (tabs) at top +# of each HTML page. A value of NO enables the index and the value YES disables +# it. Since the tabs in the index contain the same information as the navigation +# tree, you can set this option to YES if you also set GENERATE_TREEVIEW to YES. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +DISABLE_INDEX = NO + +# The GENERATE_TREEVIEW tag is used to specify whether a tree-like index +# structure should be generated to display hierarchical information. If the tag +# value is set to YES, a side panel will be generated containing a tree-like +# index structure (just like the one that is generated for HTML Help). For this +# 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. 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 + +# 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. +# +# Note that a value of 0 will completely suppress the enum values from appearing +# in the overview section. +# Minimum value: 0, maximum value: 20, default value: 4. +# This tag requires that the tag GENERATE_HTML is set to YES. + +ENUM_VALUES_PER_LINE = 1 + +# If the treeview is enabled (see GENERATE_TREEVIEW) then this tag can be used +# to set the initial width (in pixels) of the frame in which the tree is shown. +# Minimum value: 0, maximum value: 1500, default value: 250. +# This tag requires that the tag GENERATE_HTML is set to YES. + +TREEVIEW_WIDTH = 250 + +# If the EXT_LINKS_IN_WINDOW option is set to YES, doxygen will open links to +# external symbols imported via tag files in a separate window. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +EXT_LINKS_IN_WINDOW = NO + +# Use this tag to change the font size of LaTeX formulas included as images in +# the HTML documentation. When you change the font size after a successful +# doxygen run you need to manually remove any form_*.png images from the HTML +# output directory to force them to be regenerated. +# Minimum value: 8, maximum value: 50, default value: 10. +# This tag requires that the tag GENERATE_HTML is set to YES. + +FORMULA_FONTSIZE = 10 + +# Use the FORMULA_TRANPARENT 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 + +# Enable the USE_MATHJAX option to render LaTeX formulas using MathJax (see +# http://www.mathjax.org) which uses client side Javascript for the rendering +# instead of using pre-rendered bitmaps. Use this if you do not have LaTeX +# installed or if you want to formulas look prettier in the HTML output. When +# enabled you may also need to install MathJax separately and configure the path +# to it using the MATHJAX_RELPATH option. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +USE_MATHJAX = YES + +# When MathJax is enabled you can set the default output format to be used for +# the MathJax output. See the MathJax site (see: +# http://docs.mathjax.org/en/latest/output.html) for more details. +# Possible values are: HTML-CSS (which is slower, but has the best +# 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. + +MATHJAX_FORMAT = HTML-CSS + +# When MathJax is enabled you need to specify the location relative to the HTML +# output directory using the MATHJAX_RELPATH option. The destination directory +# should contain the MathJax.js script. For instance, if the mathjax directory +# is located at the same level as the HTML output directory, then +# 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 http://www.mathjax.org before deployment. +# The default value is: http://cdn.mathjax.org/mathjax/latest. +# This tag requires that the tag USE_MATHJAX is set to YES. + +MATHJAX_RELPATH = http://cdn.mathjax.org/mathjax/latest + +# The MATHJAX_EXTENSIONS tag can be used to specify one or more MathJax +# extension names that should be enabled during MathJax rendering. For example +# MATHJAX_EXTENSIONS = TeX/AMSmath TeX/AMSsymbols +# This tag requires that the tag USE_MATHJAX is set to YES. + +MATHJAX_EXTENSIONS = + +# The MATHJAX_CODEFILE tag can be used to specify a file with javascript pieces +# of code that will be used on startup of the MathJax code. See the MathJax site +# (see: http://docs.mathjax.org/en/latest/output.html) for more details. For an +# example see the documentation. +# This tag requires that the tag USE_MATHJAX is set to YES. + +MATHJAX_CODEFILE = + +# When the SEARCHENGINE tag is enabled doxygen will generate a search box for +# the HTML output. The underlying search engine uses javascript and DHTML and +# should work on any modern browser. Note that when using HTML help +# (GENERATE_HTMLHELP), Qt help (GENERATE_QHP), or docsets (GENERATE_DOCSET) +# there is already a search function so this one should typically be disabled. +# For large projects the javascript based search engine can be slow, then +# enabling SERVER_BASED_SEARCH may provide a better solution. It is possible to +# search using the keyboard; to jump to the search box use + S +# (what the is depends on the OS and browser, but it is typically +# , /