diff --git a/CMakeLists.txt b/CMakeLists.txt index 65e63ec..8982717 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -34,6 +34,7 @@ option(WITH_NVIDIA "Enable CUDA backend" OFF) option(WITH_ILUVATAR "Enable Iluvatar GPU backend" OFF) option(WITH_HYGON "Enable Hygon GPU backend" OFF) option(WITH_METAX "Enable MetaX backend" OFF) +option(WITH_MARS "Enable Mars backend" OFF) option(WITH_CAMBRICON "Enable Cambricon backend" OFF) option(WITH_MOORE "Enable Moore backend" OFF) option(WITH_ASCEND "Enable Ascend backend" OFF) @@ -76,9 +77,48 @@ if(AUTO_DETECT_DEVICES) endif() endif() - if(DEFINED ENV{MACA_PATH}) + set(_infinirt_hpcc_env_root "") + if(DEFINED ENV{HPCC_PATH} AND + EXISTS "$ENV{HPCC_PATH}/include/hcr/hc_runtime.h") + set(_infinirt_hpcc_env_root "$ENV{HPCC_PATH}") + endif() + + set(_infinirt_maca_env_root "") + if(DEFINED ENV{MACA_PATH} AND + EXISTS "$ENV{MACA_PATH}/include/mcr/mc_runtime.h") + set(_infinirt_maca_env_root "$ENV{MACA_PATH}") + endif() + + set(_infinirt_default_hpcc_available FALSE) + if(EXISTS "/opt/hpcc/include/hcr/hc_runtime.h") + set(_infinirt_default_hpcc_available TRUE) + endif() + + set(_infinirt_default_maca_available FALSE) + if(EXISTS "/opt/maca/include/mcr/mc_runtime.h") + set(_infinirt_default_maca_available TRUE) + endif() + + if(WITH_MARS OR WITH_METAX) + message(STATUS "Skipping MetaX-family auto-detection because a backend was explicitly enabled") + elseif(_infinirt_hpcc_env_root AND _infinirt_maca_env_root) + message(FATAL_ERROR + "Both HPCC_PATH and MACA_PATH identify valid SDKs. Select Mars with -DWITH_MARS=ON or MetaX with -DWITH_METAX=ON.") + elseif(_infinirt_hpcc_env_root) + set(WITH_MARS ON) + message(STATUS "Auto-detected Mars HPCC environment from HPCC_PATH") + elseif(_infinirt_maca_env_root) set(WITH_METAX ON) - message(STATUS "Auto-detected MetaX environment from MACA_PATH") + message(STATUS "Auto-detected MetaX MACA environment from MACA_PATH") + elseif(_infinirt_default_hpcc_available AND _infinirt_default_maca_available) + message(FATAL_ERROR + "Both /opt/hpcc and /opt/maca contain valid SDKs. Select Mars with -DWITH_MARS=ON or MetaX with -DWITH_METAX=ON.") + elseif(_infinirt_default_hpcc_available) + set(WITH_MARS ON) + message(STATUS "Auto-detected Mars HPCC environment from /opt/hpcc") + elseif(_infinirt_default_maca_available) + set(WITH_METAX ON) + message(STATUS "Auto-detected MetaX MACA environment from /opt/maca") else() execute_process( COMMAND sh -c "grep -h 9999 /sys/bus/pci/devices/*/vendor 2>/dev/null" @@ -89,8 +129,7 @@ if(AUTO_DETECT_DEVICES) string(FIND "${_pci_vendor_output}" "9999" _found_pos) if(_found_pos GREATER -1) - set(WITH_METAX ON) - message(STATUS "Detected MetaX GPU from PCI vendor ID 0x9999") + message(STATUS "Detected a MetaX-family GPU from PCI vendor ID 0x9999, but no MACA or HPCC SDK was found") else() set(WITH_METAX OFF) message(STATUS "No MetaX GPU detected") @@ -118,7 +157,7 @@ if(AUTO_DETECT_DEVICES) if(WITH_NVIDIA) set(_non_nvidia_gpu_detected FALSE) - foreach(_gpu_backend WITH_ILUVATAR WITH_HYGON WITH_METAX WITH_CAMBRICON WITH_MOORE WITH_ASCEND) + foreach(_gpu_backend WITH_ILUVATAR WITH_HYGON WITH_METAX WITH_MARS WITH_CAMBRICON WITH_MOORE WITH_ASCEND) if(${_gpu_backend}) set(_non_nvidia_gpu_detected TRUE) endif() @@ -135,14 +174,14 @@ include_directories(${CMAKE_CURRENT_SOURCE_DIR}/src) # Only one CUDA-like GPU backend can be enabled at a time. set(_gpu_backend_count 0) -foreach(_gpu_backend WITH_NVIDIA WITH_ILUVATAR WITH_HYGON WITH_METAX WITH_MOORE WITH_ASCEND) +foreach(_gpu_backend WITH_NVIDIA WITH_ILUVATAR WITH_HYGON WITH_METAX WITH_MARS WITH_MOORE WITH_ASCEND) if(${_gpu_backend}) math(EXPR _gpu_backend_count "${_gpu_backend_count} + 1") endif() endforeach() if(_gpu_backend_count GREATER 1) - message(FATAL_ERROR "`WITH_NVIDIA`, `WITH_ILUVATAR`, `WITH_HYGON`, `WITH_METAX`, `WITH_MOORE`, and `WITH_ASCEND` are mutually exclusive. Build one GPU backend at a time.") + message(FATAL_ERROR "`WITH_NVIDIA`, `WITH_ILUVATAR`, `WITH_HYGON`, `WITH_METAX`, `WITH_MARS`, `WITH_MOORE`, and `WITH_ASCEND` are mutually exclusive. Build one GPU backend at a time.") endif() if(WITH_NVIDIA) @@ -220,18 +259,38 @@ endif() if(WITH_METAX) add_compile_definitions(WITH_METAX=1) - # Normally can be found at: `/opt/maca/`. - set(MACA_PATH $ENV{MACA_PATH}) - set(CMAKE_C_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) + set(MACA_PATH "$ENV{MACA_PATH}" CACHE PATH "Path to the MetaX MACA SDK") + if(NOT MACA_PATH) + set(MACA_PATH "/opt/maca") + endif() + if(NOT EXISTS "${MACA_PATH}/include/mcr/mc_runtime.h") + message(FATAL_ERROR "`WITH_METAX` requires a MACA SDK. Set MACA_PATH to a root containing include/mcr/mc_runtime.h.") + endif() set(CMAKE_CXX_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) + message(STATUS "MetaX: using MACA SDK at ${MACA_PATH}") include_directories("${MACA_PATH}/include") link_directories("${MACA_PATH}/lib") - - # Libraries: mcruntime / mcdnn / mcblas. find_library(MACA_RUNTIME_LIB NAMES mcruntime HINTS "${MACA_PATH}/lib" REQUIRED) endif() +if(WITH_MARS) + add_compile_definitions(WITH_MARS=1) + + set(HPCC_PATH "$ENV{HPCC_PATH}" CACHE PATH "Path to the Mars HPCC SDK") + if(NOT HPCC_PATH) + set(HPCC_PATH "/opt/hpcc") + endif() + if(NOT EXISTS "${HPCC_PATH}/include/hcr/hc_runtime.h") + message(FATAL_ERROR "`WITH_MARS` requires an HPCC SDK. Set HPCC_PATH to a root containing include/hcr/hc_runtime.h.") + endif() + message(STATUS "Mars: using HPCC SDK at ${HPCC_PATH}") + + include_directories("${HPCC_PATH}/include") + link_directories("${HPCC_PATH}/lib") + find_library(HPCC_RUNTIME_LIB NAMES hcruntime HINTS "${HPCC_PATH}/lib" REQUIRED) +endif() + if(WITH_MOORE) add_compile_definitions(WITH_MOORE=1) @@ -294,7 +353,7 @@ if(WITH_ASCEND) endif() # If all other platforms are not enabled, CPU is enabled by default. -if(NOT WITH_NVIDIA AND NOT WITH_ILUVATAR AND NOT WITH_HYGON AND NOT WITH_METAX AND NOT WITH_MOORE AND NOT WITH_CAMBRICON AND NOT WITH_ASCEND) +if(NOT WITH_NVIDIA AND NOT WITH_ILUVATAR AND NOT WITH_HYGON AND NOT WITH_METAX AND NOT WITH_MARS AND NOT WITH_MOORE AND NOT WITH_CAMBRICON AND NOT WITH_ASCEND) set(WITH_CPU ON) add_compile_definitions(WITH_CPU=1) endif() @@ -315,6 +374,9 @@ endif() if(WITH_METAX) list(APPEND INFINI_RT_PUBLIC_HEADER_DEVICES metax) endif() +if(WITH_MARS) + list(APPEND INFINI_RT_PUBLIC_HEADER_DEVICES mars) +endif() if(WITH_MOORE) list(APPEND INFINI_RT_PUBLIC_HEADER_DEVICES moore) endif() @@ -347,7 +409,7 @@ if(WITH_NVIDIA OR WITH_ILUVATAR OR WITH_HYGON) endif() set(INFINI_RT_PACKAGE_VENDOR_BACKEND "") -foreach(_backend METAX MOORE CAMBRICON ASCEND) +foreach(_backend METAX MARS MOORE CAMBRICON ASCEND) if(WITH_${_backend}) set(INFINI_RT_PACKAGE_VENDOR_BACKEND "${_backend}") endif() diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index b387366..1d51ca8 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -117,6 +117,7 @@ Common backend options are: -DWITH_ILUVATAR=ON -DWITH_HYGON=ON -DWITH_METAX=ON +-DWITH_MARS=ON -DWITH_MOORE=ON -DWITH_CAMBRICON=ON -DWITH_ASCEND=ON @@ -200,7 +201,8 @@ ruff check . 1. **CMake cannot find a backend SDK**: Check the backend-specific environment variables documented in `docs/backends.md`, such as `DTK_ROOT`, - `MACA_PATH`, `MUSA_ROOT`, `NEUWARE_HOME`, or `ASCEND_HOME_PATH`. + `MACA_PATH`, `HPCC_PATH`, `MUSA_ROOT`, `NEUWARE_HOME`, or + `ASCEND_HOME_PATH`. 2. **Switching between backends gives stale build errors**: Use a separate build directory per platform or delete the stale build directory. 3. **Generated headers are stale**: Re-run CMake configure. Public headers are diff --git a/cmake/InfiniRTConfig.cmake.in b/cmake/InfiniRTConfig.cmake.in index 09bce5a..137cdba 100644 --- a/cmake/InfiniRTConfig.cmake.in +++ b/cmake/InfiniRTConfig.cmake.in @@ -10,6 +10,7 @@ set(InfiniRT_WITH_NVIDIA @WITH_NVIDIA@) set(InfiniRT_WITH_ILUVATAR @WITH_ILUVATAR@) set(InfiniRT_WITH_HYGON @WITH_HYGON@) set(InfiniRT_WITH_METAX @WITH_METAX@) +set(InfiniRT_WITH_MARS @WITH_MARS@) set(InfiniRT_WITH_MOORE @WITH_MOORE@) set(InfiniRT_WITH_CAMBRICON @WITH_CAMBRICON@) set(InfiniRT_WITH_ASCEND @WITH_ASCEND@) @@ -34,6 +35,7 @@ if(_InfiniRT_VENDOR_BACKEND STREQUAL "METAX") if(DEFINED ENV{MACA_PATH}) list(APPEND _InfiniRT_ROOT_HINTS "$ENV{MACA_PATH}") endif() + list(APPEND _InfiniRT_ROOT_HINTS /opt/maca) find_path(_InfiniRT_BACKEND_INCLUDE_DIR NAMES mcr/mc_runtime.h HINTS ${_InfiniRT_ROOT_HINTS} @@ -52,6 +54,30 @@ if(_InfiniRT_VENDOR_BACKEND STREQUAL "METAX") _infinirt_define_backend_runtime( "${_InfiniRT_BACKEND_INCLUDE_DIR}" "${_InfiniRT_BACKEND_RUNTIME_LIBRARY}") +elseif(_InfiniRT_VENDOR_BACKEND STREQUAL "MARS") + set(_InfiniRT_ROOT_HINTS "${InfiniRT_MARS_ROOT}") + if(DEFINED ENV{HPCC_PATH}) + list(APPEND _InfiniRT_ROOT_HINTS "$ENV{HPCC_PATH}") + endif() + list(APPEND _InfiniRT_ROOT_HINTS /opt/hpcc) + find_path(_InfiniRT_BACKEND_INCLUDE_DIR + NAMES hcr/hc_runtime.h + HINTS ${_InfiniRT_ROOT_HINTS} + PATH_SUFFIXES include) + find_library(_InfiniRT_BACKEND_RUNTIME_LIBRARY + NAMES hcruntime + HINTS ${_InfiniRT_ROOT_HINTS} + PATH_SUFFIXES lib lib64) + if(NOT _InfiniRT_BACKEND_INCLUDE_DIR OR + NOT _InfiniRT_BACKEND_RUNTIME_LIBRARY) + set(InfiniRT_FOUND FALSE) + set(InfiniRT_NOT_FOUND_MESSAGE + "Mars HPCC SDK not found. Set InfiniRT_MARS_ROOT or HPCC_PATH.") + return() + endif() + _infinirt_define_backend_runtime( + "${_InfiniRT_BACKEND_INCLUDE_DIR}" + "${_InfiniRT_BACKEND_RUNTIME_LIBRARY}") elseif(_InfiniRT_VENDOR_BACKEND STREQUAL "MOORE") set(_InfiniRT_ROOT_HINTS "${InfiniRT_MOORE_ROOT}") foreach(_InfiniRT_ENV MUSA_ROOT MUSA_HOME MUSA_PATH) diff --git a/docs/api/core-types.md b/docs/api/core-types.md index 3517a9b..61fd307 100644 --- a/docs/api/core-types.md +++ b/docs/api/core-types.md @@ -22,6 +22,7 @@ Known device types include: - `Device::Type::kNvidia` - `Device::Type::kIluvatar` - `Device::Type::kMetax` +- `Device::Type::kMars` - `Device::Type::kMoore` - `Device::Type::kHygon` - `Device::Type::kCambricon` diff --git a/docs/backends.md b/docs/backends.md index 26ad1ff..a55566c 100644 --- a/docs/backends.md +++ b/docs/backends.md @@ -12,6 +12,7 @@ backend. | NVIDIA | `WITH_NVIDIA` | Requires CUDA toolkit. | | Iluvatar | `WITH_ILUVATAR` | CUDA-compatible backend using clang with ivcore flags. | | MetaX | `WITH_METAX` | Requires `MACA_PATH`. | +| Mars | `WITH_MARS` | Requires the HPCC SDK from `HPCC_PATH` or `/opt/hpcc`. | | Moore | `WITH_MOORE` | Requires `MUSA_ROOT`, `MUSA_HOME`, or `MUSA_PATH`. | | Hygon | `WITH_HYGON` | Requires DTK and a DTK CUDA toolkit. | | Cambricon | `WITH_CAMBRICON` | Requires `NEUWARE_HOME`. | @@ -20,6 +21,9 @@ backend. Only one accelerator backend can be enabled at a time. CPU can be enabled together with the selected accelerator backend. +When both HPCC and MACA SDKs are installed, select `WITH_MARS` or `WITH_METAX` +explicitly. Automatic detection rejects an ambiguous MetaX-family environment. + ## Runtime API Support All enabled runtime backends provide the core device and memory entry points @@ -33,7 +37,8 @@ Current test expectations are: | CPU | No | Yes | No | Yes | No | Yes | No | | NVIDIA | Yes | Yes | Yes | Yes | Yes | Yes | Yes | | Iluvatar | Yes | Yes | Yes | Yes | Yes | Yes | Yes | -| MetaX | Yes | Yes | Yes | Yes | Yes | Yes | No | +| MetaX | Yes | Yes | Yes | Yes | Yes | Yes | Yes | +| Mars | Yes | Yes | Yes | Yes | Yes | Yes | Yes | | Moore | Yes | Yes | No | Yes | Yes | Yes | No | | Hygon | Yes | Yes | Yes | Yes | Yes | Yes | No | | Cambricon | Yes | No | No | No | No | No | No | diff --git a/docs/build.md b/docs/build.md index b04e6df..b329859 100644 --- a/docs/build.md +++ b/docs/build.md @@ -9,6 +9,7 @@ InfiniRT uses CMake and C++17. -DWITH_NVIDIA=ON -DWITH_ILUVATAR=ON -DWITH_METAX=ON +-DWITH_MARS=ON -DWITH_MOORE=ON -DWITH_HYGON=ON -DWITH_CAMBRICON=ON diff --git a/scripts/generate_public_headers.py b/scripts/generate_public_headers.py index 3363741..854e71a 100644 --- a/scripts/generate_public_headers.py +++ b/scripts/generate_public_headers.py @@ -2,6 +2,7 @@ import dataclasses import pathlib import re +import shutil _DETAIL_PREFIX = "infini/rt/detail" @@ -28,9 +29,14 @@ ("hygon", "runtime_.h", "native/cuda/hygon/runtime_.h"), ), "metax": ( - ("metax", "data_type_.h", "native/cuda/metax/data_type_.h"), - ("metax", "device_.h", "native/cuda/metax/device_.h"), - ("metax", "runtime_.h", "native/cuda/metax/runtime_.h"), + ("metax", "data_type_.h", "native/maca/metax/data_type_.h"), + ("metax", "device_.h", "native/maca/metax/device_.h"), + ("metax", "runtime_.h", "native/maca/metax/runtime_.h"), + ), + "mars": ( + ("mars", "data_type_.h", "native/hpcc/mars/data_type_.h"), + ("mars", "device_.h", "native/hpcc/mars/device_.h"), + ("mars", "runtime_.h", "native/hpcc/mars/runtime_.h"), ), "moore": ( ("moore", "data_type_.h", "native/cuda/moore/data_type_.h"), @@ -55,6 +61,7 @@ "iluvatar": "Device::Type::kIluvatar", "hygon": "Device::Type::kHygon", "metax": "Device::Type::kMetax", + "mars": "Device::Type::kMars", "moore": "Device::Type::kMoore", "cambricon": "Device::Type::kCambricon", "ascend": "Device::Type::kAscend", @@ -65,6 +72,7 @@ "iluvatar", "hygon", "metax", + "mars", "moore", "cambricon", "ascend", @@ -663,6 +671,9 @@ def main(): include_root = pathlib.Path(args.output_dir) source_root = pathlib.Path(args.runtime_header).parent + if include_root.exists(): + shutil.rmtree(include_root) + _write_detail_headers(include_root, source_root, devices) for device in devices: for wrapper_device, header_name, target in _DEVICE_HEADERS[device]: diff --git a/scripts/run_performance_tests.py b/scripts/run_performance_tests.py index 4e08c0d..630413e 100644 --- a/scripts/run_performance_tests.py +++ b/scripts/run_performance_tests.py @@ -12,6 +12,7 @@ ("WITH_ILUVATAR", "iluvatar"), ("WITH_HYGON", "hygon"), ("WITH_METAX", "metax"), + ("WITH_MARS", "mars"), ("WITH_MOORE", "moore"), ("WITH_CAMBRICON", "cambricon"), ("WITH_ASCEND", "ascend"), diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 2091963..f755ca6 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -64,12 +64,19 @@ if(WITH_HYGON) endif() if(WITH_METAX) - target_compile_definitions(infinirt PRIVATE WITH_METAX=1) + target_compile_definitions(infinirt PUBLIC WITH_METAX=1) _infinirt_use_backend_runtime( "${MACA_PATH}/include" "${MACA_RUNTIME_LIB}") endif() +if(WITH_MARS) + target_compile_definitions(infinirt PUBLIC WITH_MARS=1) + _infinirt_use_backend_runtime( + "${HPCC_PATH}/include" + "${HPCC_RUNTIME_LIB}") +endif() + if(WITH_MOORE) target_compile_definitions(infinirt PRIVATE WITH_MOORE=1) _infinirt_use_backend_runtime( diff --git a/src/device.h b/src/device.h index 482b2fb..d1f04d1 100644 --- a/src/device.h +++ b/src/device.h @@ -21,6 +21,7 @@ class Device { kMoore = 5, kIluvatar = 6, kHygon = 7, + kMars = 8, kCount }; @@ -64,6 +65,7 @@ class Device { {Type::kMoore, "moore"}, {Type::kIluvatar, "iluvatar"}, {Type::kHygon, "hygon"}, + {Type::kMars, "mars"}, }}}; static constexpr ConstexprMap; + Device::Type::kIluvatar, Device::Type::kHygon, Device::Type::kMars>; // Deferred computation of active devices. The `Filter` and `FilterList` // evaluation are nested inside a class template so that `DeviceEnabled` diff --git a/src/native/hpcc/mars/data_type_.h b/src/native/hpcc/mars/data_type_.h new file mode 100644 index 0000000..708c154 --- /dev/null +++ b/src/native/hpcc/mars/data_type_.h @@ -0,0 +1,25 @@ +#ifndef INFINI_RT_MARS_DATA_TYPE__H_ +#define INFINI_RT_MARS_DATA_TYPE__H_ + +#include +#include +#include + +#include "data_type.h" +#include "native/hpcc/mars/device_.h" + +namespace infini::rt { + +template <> +struct TypeMap { + using type = __half; +}; + +template <> +struct TypeMap { + using type = __maca_bfloat16; +}; + +} // namespace infini::rt + +#endif diff --git a/src/native/hpcc/mars/device_.h b/src/native/hpcc/mars/device_.h new file mode 100644 index 0000000..7ec4f05 --- /dev/null +++ b/src/native/hpcc/mars/device_.h @@ -0,0 +1,13 @@ +#ifndef INFINI_RT_MARS_DEVICE__H_ +#define INFINI_RT_MARS_DEVICE__H_ + +#include "device.h" + +namespace infini::rt { + +template <> +struct DeviceEnabled : std::true_type {}; + +} // namespace infini::rt + +#endif diff --git a/src/native/hpcc/mars/runtime_.h b/src/native/hpcc/mars/runtime_.h new file mode 100644 index 0000000..cd7ee5a --- /dev/null +++ b/src/native/hpcc/mars/runtime_.h @@ -0,0 +1,131 @@ +#ifndef INFINI_RT_MARS_RUNTIME__H_ +#define INFINI_RT_MARS_RUNTIME__H_ + +#include + +#include + +#include "native/hpcc/mars/device_.h" +#include "native/hpcc/runtime_.h" + +namespace infini::rt::runtime { + +template <> +struct Runtime + : GraphRuntime, + HpccRuntime>> { + using Error = hcError_t; + using Stream = hcStream_t; + using Graph = hcGraph_t; + using GraphExec = hcGraphExec_t; + using Event = hcEvent_t; + using StreamCaptureMode = hcStreamCaptureMode; + + static constexpr Device::Type kDeviceType = Device::Type::kMars; + static constexpr Error kSuccess = hcSuccess; + static constexpr auto SetDevice = hcSetDevice; + static constexpr auto GetDevice = hcGetDevice; + static constexpr auto GetDeviceCount = hcGetDeviceCount; + static constexpr auto DeviceSynchronize = hcDeviceSynchronize; + + static constexpr auto Malloc = [](auto&&... args) { + return hcMalloc(std::forward(args)...); + }; + static constexpr auto MallocHost = [](auto&&... args) { + return hcMallocHost(std::forward(args)...); + }; + static constexpr auto MallocAsync = [](auto&&... args) { + return hcMallocAsync(std::forward(args)...); + }; + static constexpr auto Free = [](auto&&... args) { + return hcFree(std::forward(args)...); + }; + static constexpr auto FreeHost = [](auto&&... args) { + return hcFreeHost(std::forward(args)...); + }; + static constexpr auto FreeAsync = [](auto&&... args) { + return hcFreeAsync(std::forward(args)...); + }; + static constexpr auto MemGetInfo = [](auto&&... args) { + return hcMemGetInfo(std::forward(args)...); + }; + static constexpr auto Memcpy = [](auto&&... args) { + return hcMemcpy(std::forward(args)...); + }; + static constexpr auto MemcpyAsync = [](auto&&... args) { + return hcMemcpyAsync(std::forward(args)...); + }; + + static constexpr auto kMemcpyHostToHost = hcMemcpyHostToHost; + static constexpr auto kMemcpyHostToDevice = hcMemcpyHostToDevice; + static constexpr auto kMemcpyDeviceToHost = hcMemcpyDeviceToHost; + static constexpr auto kMemcpyDeviceToDevice = hcMemcpyDeviceToDevice; + static constexpr auto Memset = hcMemset; + + static constexpr auto MemsetAsync = [](auto&&... args) { + return hcMemsetAsync(std::forward(args)...); + }; + static constexpr auto StreamCreate = [](auto&&... args) { + return hcStreamCreate(std::forward(args)...); + }; + static constexpr auto StreamDestroy = [](auto&&... args) { + return hcStreamDestroy(std::forward(args)...); + }; + static constexpr auto StreamSynchronize = [](auto&&... args) { + return hcStreamSynchronize(std::forward(args)...); + }; + static constexpr auto StreamWaitEvent = [](auto&&... args) { + return hcStreamWaitEvent(std::forward(args)...); + }; + static constexpr auto EventCreate = [](auto&&... args) { + return hcEventCreate(std::forward(args)...); + }; + static constexpr auto EventCreateWithFlags = [](auto&&... args) { + return hcEventCreateWithFlags(std::forward(args)...); + }; + static constexpr auto EventRecord = [](auto&&... args) { + return hcEventRecord(std::forward(args)...); + }; + static constexpr auto EventQuery = [](auto&&... args) { + return hcEventQuery(std::forward(args)...); + }; + static constexpr auto EventSynchronize = [](auto&&... args) { + return hcEventSynchronize(std::forward(args)...); + }; + static constexpr auto EventDestroy = [](auto&&... args) { + return hcEventDestroy(std::forward(args)...); + }; + static constexpr auto EventElapsedTime = [](auto&&... args) { + return hcEventElapsedTime(std::forward(args)...); + }; + + static constexpr auto kStreamCaptureModeGlobal = hcStreamCaptureModeGlobal; + static constexpr auto kStreamCaptureModeThreadLocal = + hcStreamCaptureModeThreadLocal; + static constexpr auto kStreamCaptureModeRelaxed = hcStreamCaptureModeRelaxed; + + static constexpr auto StreamBeginCapture = [](auto&&... args) { + return hcStreamBeginCapture(std::forward(args)...); + }; + static constexpr auto StreamEndCapture = [](auto&&... args) { + return hcStreamEndCapture(std::forward(args)...); + }; + static constexpr auto GraphDestroy = [](auto&&... args) { + return hcGraphDestroy(std::forward(args)...); + }; + static Error GraphInstantiate(GraphExec* graph_exec, Graph graph) { + return hcGraphInstantiate(graph_exec, graph, nullptr, nullptr, 0); + } + static constexpr auto GraphExecDestroy = [](auto&&... args) { + return hcGraphExecDestroy(std::forward(args)...); + }; + static constexpr auto GraphLaunch = [](auto&&... args) { + return hcGraphLaunch(std::forward(args)...); + }; +}; + +static_assert(Runtime::Validate()); + +} // namespace infini::rt::runtime + +#endif diff --git a/src/native/hpcc/runtime_.h b/src/native/hpcc/runtime_.h new file mode 100644 index 0000000..cd1a31a --- /dev/null +++ b/src/native/hpcc/runtime_.h @@ -0,0 +1,29 @@ +#ifndef INFINI_RT_HPCC_RUNTIME_H_ +#define INFINI_RT_HPCC_RUNTIME_H_ + +#include + +#include "runtime.h" + +namespace infini::rt::runtime { + +template +struct HpccRuntime : DeviceRuntime { + static constexpr bool Validate() { + DeviceRuntime::Validate(); + static_assert( + std::is_invocable_v, + "`Runtime::Memcpy` must accept the HPCC memcpy kind."); + static_assert( + std::is_invocable_v, + "`Runtime::MemcpyAsync` must accept the HPCC memcpy kind and stream."); + return true; + } +}; + +} // namespace infini::rt::runtime + +#endif diff --git a/src/native/cuda/metax/data_type_.h b/src/native/maca/metax/data_type_.h similarity index 79% rename from src/native/cuda/metax/data_type_.h rename to src/native/maca/metax/data_type_.h index 88d3e24..7e08eff 100644 --- a/src/native/cuda/metax/data_type_.h +++ b/src/native/maca/metax/data_type_.h @@ -6,14 +6,10 @@ #include #include "data_type.h" -#include "native/cuda/metax/device_.h" +#include "native/maca/metax/device_.h" namespace infini::rt { -using cuda_bfloat16 = maca_bfloat16; - -using cuda_bfloat162 = maca_bfloat162; - template <> struct TypeMap { using type = __half; diff --git a/src/native/cuda/metax/device_.h b/src/native/maca/metax/device_.h similarity index 100% rename from src/native/cuda/metax/device_.h rename to src/native/maca/metax/device_.h diff --git a/src/native/cuda/metax/runtime_.h b/src/native/maca/metax/runtime_.h similarity index 73% rename from src/native/cuda/metax/runtime_.h rename to src/native/maca/metax/runtime_.h index 462f5f1..1982fe8 100644 --- a/src/native/cuda/metax/runtime_.h +++ b/src/native/maca/metax/runtime_.h @@ -5,125 +5,123 @@ #include -#include "native/cuda/metax/device_.h" -#include "native/cuda/runtime_.h" +#include "native/maca/metax/device_.h" +#include "native/maca/runtime_.h" namespace infini::rt::runtime { template <> struct Runtime - : CudaRuntime> { + : GraphRuntime, + MacaRuntime>> { using Error = mcError_t; - using Stream = mcStream_t; - + using Graph = mcGraph_t; + using GraphExec = mcGraphExec_t; using Event = mcEvent_t; + using StreamCaptureMode = mcStreamCaptureMode; static constexpr Device::Type kDeviceType = Device::Type::kMetax; - static constexpr Error kSuccess = mcSuccess; - static constexpr auto SetDevice = mcSetDevice; - static constexpr auto GetDevice = mcGetDevice; - static constexpr auto GetDeviceCount = mcGetDeviceCount; - static constexpr auto DeviceSynchronize = mcDeviceSynchronize; static constexpr auto Malloc = [](auto&&... args) { return mcMalloc(std::forward(args)...); }; - static constexpr auto MallocHost = [](auto&&... args) { return mcMallocHost(std::forward(args)...); }; - static constexpr auto MallocAsync = [](auto&&... args) { return mcMallocAsync(std::forward(args)...); }; - static constexpr auto Free = [](auto&&... args) { return mcFree(std::forward(args)...); }; - static constexpr auto FreeHost = [](auto&&... args) { return mcFreeHost(std::forward(args)...); }; - static constexpr auto FreeAsync = [](auto&&... args) { return mcFreeAsync(std::forward(args)...); }; - static constexpr auto MemGetInfo = [](auto&&... args) { return mcMemGetInfo(std::forward(args)...); }; - static constexpr auto Memcpy = [](auto&&... args) { return mcMemcpy(std::forward(args)...); }; - static constexpr auto MemcpyAsync = [](auto&&... args) { return mcMemcpyAsync(std::forward(args)...); }; static constexpr auto kMemcpyHostToHost = mcMemcpyHostToHost; - static constexpr auto kMemcpyHostToDevice = mcMemcpyHostToDevice; - static constexpr auto kMemcpyDeviceToHost = mcMemcpyDeviceToHost; - static constexpr auto kMemcpyDeviceToDevice = mcMemcpyDeviceToDevice; - static constexpr auto Memset = mcMemset; static constexpr auto MemsetAsync = [](auto&&... args) { return mcMemsetAsync(std::forward(args)...); }; - static constexpr auto StreamCreate = [](auto&&... args) { return mcStreamCreate(std::forward(args)...); }; - static constexpr auto StreamDestroy = [](auto&&... args) { return mcStreamDestroy(std::forward(args)...); }; - static constexpr auto StreamSynchronize = [](auto&&... args) { return mcStreamSynchronize(std::forward(args)...); }; - static constexpr auto StreamWaitEvent = [](auto&&... args) { return mcStreamWaitEvent(std::forward(args)...); }; - static constexpr auto EventCreate = [](auto&&... args) { return mcEventCreate(std::forward(args)...); }; - static constexpr auto EventCreateWithFlags = [](auto&&... args) { return mcEventCreateWithFlags(std::forward(args)...); }; - static constexpr auto EventRecord = [](auto&&... args) { return mcEventRecord(std::forward(args)...); }; - static constexpr auto EventQuery = [](auto&&... args) { return mcEventQuery(std::forward(args)...); }; - static constexpr auto EventSynchronize = [](auto&&... args) { return mcEventSynchronize(std::forward(args)...); }; - static constexpr auto EventDestroy = [](auto&&... args) { return mcEventDestroy(std::forward(args)...); }; - static constexpr auto EventElapsedTime = [](auto&&... args) { return mcEventElapsedTime(std::forward(args)...); }; + + static constexpr auto kStreamCaptureModeGlobal = mcStreamCaptureModeGlobal; + static constexpr auto kStreamCaptureModeThreadLocal = + mcStreamCaptureModeThreadLocal; + static constexpr auto kStreamCaptureModeRelaxed = mcStreamCaptureModeRelaxed; + + static constexpr auto StreamBeginCapture = [](auto&&... args) { + return mcStreamBeginCapture(std::forward(args)...); + }; + static constexpr auto StreamEndCapture = [](auto&&... args) { + return mcStreamEndCapture(std::forward(args)...); + }; + static constexpr auto GraphDestroy = [](auto&&... args) { + return mcGraphDestroy(std::forward(args)...); + }; + static Error GraphInstantiate(GraphExec* graph_exec, Graph graph) { + return mcGraphInstantiate(graph_exec, graph, nullptr, nullptr, 0); + } + static constexpr auto GraphExecDestroy = [](auto&&... args) { + return mcGraphExecDestroy(std::forward(args)...); + }; + static constexpr auto GraphLaunch = [](auto&&... args) { + return mcGraphLaunch(std::forward(args)...); + }; }; static_assert(Runtime::Validate()); diff --git a/src/native/maca/runtime_.h b/src/native/maca/runtime_.h new file mode 100644 index 0000000..a9613f7 --- /dev/null +++ b/src/native/maca/runtime_.h @@ -0,0 +1,29 @@ +#ifndef INFINI_RT_MACA_RUNTIME_H_ +#define INFINI_RT_MACA_RUNTIME_H_ + +#include + +#include "runtime.h" + +namespace infini::rt::runtime { + +template +struct MacaRuntime : DeviceRuntime { + static constexpr bool Validate() { + DeviceRuntime::Validate(); + static_assert( + std::is_invocable_v, + "`Runtime::Memcpy` must accept the MACA memcpy kind."); + static_assert( + std::is_invocable_v, + "`Runtime::MemcpyAsync` must accept the MACA memcpy kind and stream."); + return true; + } +}; + +} // namespace infini::rt::runtime + +#endif diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 1dd2048..73052c2 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -42,6 +42,13 @@ endfunction() add_infini_rt_test(test_smoke test_smoke.cc) add_infini_rt_test(test_core test_core.cc) +add_test( + NAME test_ambiguous_metax_family_autodetect + COMMAND "${CMAKE_COMMAND}" + "-DINFINI_RT_SOURCE_DIR=${PROJECT_SOURCE_DIR}" + "-DINFINI_RT_TEST_BINARY_DIR=${CMAKE_CURRENT_BINARY_DIR}/ambiguous_metax_family" + -P "${CMAKE_CURRENT_SOURCE_DIR}/test_ambiguous_metax_family_autodetect.cmake") + set(INFINI_RT_TEST_HAS_RUNTIME_BACKEND OFF) if(WITH_CPU) @@ -85,6 +92,18 @@ if(WITH_METAX) METAX infini::rt::Device::Type::kMetax infini/rt/metax/runtime_.h 1 1 1 1 1 1 1 1) + add_infini_rt_backend_graph_test( + METAX infini::rt::Device::Type::kMetax 1) +endif() + +if(WITH_MARS) + set(INFINI_RT_TEST_HAS_RUNTIME_BACKEND ON) + add_infini_rt_backend_runtime_test( + MARS infini::rt::Device::Type::kMars + infini/rt/mars/runtime_.h + 1 1 1 1 1 1 1 1) + add_infini_rt_backend_graph_test( + MARS infini::rt::Device::Type::kMars 1) endif() if(WITH_MOORE) @@ -134,6 +153,10 @@ if(INFINI_RT_TEST_HAS_RUNTIME_BACKEND) target_compile_definitions(test_runtime_dispatch PRIVATE INFINI_RT_TEST_WITH_METAX=1) endif() + if(WITH_MARS) + target_compile_definitions(test_runtime_dispatch + PRIVATE INFINI_RT_TEST_WITH_MARS=1) + endif() if(WITH_MOORE) target_compile_definitions(test_runtime_dispatch PRIVATE INFINI_RT_TEST_WITH_MOORE=1) @@ -176,6 +199,10 @@ elseif(WITH_METAX) set(INFINI_RT_TEST_CONSUMER_BACKEND METAX) set(INFINI_RT_TEST_BACKEND_ROOT_VARIABLE InfiniRT_METAX_ROOT) set(INFINI_RT_TEST_BACKEND_ROOT "${MACA_PATH}") +elseif(WITH_MARS) + set(INFINI_RT_TEST_CONSUMER_BACKEND MARS) + set(INFINI_RT_TEST_BACKEND_ROOT_VARIABLE InfiniRT_MARS_ROOT) + set(INFINI_RT_TEST_BACKEND_ROOT "${HPCC_PATH}") elseif(WITH_MOORE) set(INFINI_RT_TEST_CONSUMER_BACKEND MOORE) set(INFINI_RT_TEST_BACKEND_ROOT_VARIABLE InfiniRT_MOORE_ROOT) diff --git a/tests/install_consumer/CMakeLists.txt b/tests/install_consumer/CMakeLists.txt index c1a3c33..19a9c54 100644 --- a/tests/install_consumer/CMakeLists.txt +++ b/tests/install_consumer/CMakeLists.txt @@ -24,7 +24,7 @@ if(NOT "${InfiniRT_ENABLED_BACKENDS}" STREQUAL "InfiniRT reported enabled backends '${InfiniRT_ENABLED_BACKENDS}', expected '${_infinirt_expected_backends}'.") endif() -foreach(_infinirt_backend CPU NVIDIA ILUVATAR HYGON METAX MOORE CAMBRICON ASCEND) +foreach(_infinirt_backend CPU NVIDIA ILUVATAR HYGON METAX MARS MOORE CAMBRICON ASCEND) string(TOLOWER "${_infinirt_backend}" _infinirt_backend_lower) list(FIND _infinirt_expected_backends "${_infinirt_backend_lower}" _infinirt_backend_index) diff --git a/tests/install_consumer_smoke.cc b/tests/install_consumer_smoke.cc index 6b72130..02a8e37 100644 --- a/tests/install_consumer_smoke.cc +++ b/tests/install_consumer_smoke.cc @@ -24,6 +24,7 @@ int main() { defined(INFINI_RT_CONSUMER_BACKEND_ILUVATAR) || \ defined(INFINI_RT_CONSUMER_BACKEND_HYGON) || \ defined(INFINI_RT_CONSUMER_BACKEND_METAX) || \ + defined(INFINI_RT_CONSUMER_BACKEND_MARS) || \ defined(INFINI_RT_CONSUMER_BACKEND_MOORE) || \ defined(INFINI_RT_CONSUMER_BACKEND_CAMBRICON) || \ defined(INFINI_RT_CONSUMER_BACKEND_ASCEND) @@ -43,6 +44,9 @@ int main() { #elif defined(INFINI_RT_CONSUMER_BACKEND_METAX) constexpr auto kExpectedDeviceType = infini::rt::Device::Type::kMetax; constexpr bool kExpectAsyncMemcpySuccess = true; +#elif defined(INFINI_RT_CONSUMER_BACKEND_MARS) + constexpr auto kExpectedDeviceType = infini::rt::Device::Type::kMars; + constexpr bool kExpectAsyncMemcpySuccess = true; #elif defined(INFINI_RT_CONSUMER_BACKEND_MOORE) constexpr auto kExpectedDeviceType = infini::rt::Device::Type::kMoore; constexpr bool kExpectAsyncMemcpySuccess = true; diff --git a/tests/performance/CMakeLists.txt b/tests/performance/CMakeLists.txt index e694be2..ad69724 100644 --- a/tests/performance/CMakeLists.txt +++ b/tests/performance/CMakeLists.txt @@ -7,6 +7,8 @@ elseif(WITH_HYGON) set(INFINI_RT_PERF_BACKEND_NAME "hygon") elseif(WITH_METAX) set(INFINI_RT_PERF_BACKEND_NAME "metax") +elseif(WITH_MARS) + set(INFINI_RT_PERF_BACKEND_NAME "mars") elseif(WITH_MOORE) set(INFINI_RT_PERF_BACKEND_NAME "moore") elseif(WITH_CAMBRICON) diff --git a/tests/test_ambiguous_metax_family_autodetect.cmake b/tests/test_ambiguous_metax_family_autodetect.cmake new file mode 100644 index 0000000..94cfb49 --- /dev/null +++ b/tests/test_ambiguous_metax_family_autodetect.cmake @@ -0,0 +1,38 @@ +if(NOT INFINI_RT_SOURCE_DIR OR NOT INFINI_RT_TEST_BINARY_DIR) + message(FATAL_ERROR "InfiniRT source and test binary directories are required.") +endif() + +set(_hpcc_root "${INFINI_RT_TEST_BINARY_DIR}/hpcc") +set(_maca_root "${INFINI_RT_TEST_BINARY_DIR}/maca") +set(_child_build "${INFINI_RT_TEST_BINARY_DIR}/build") +file(REMOVE_RECURSE "${INFINI_RT_TEST_BINARY_DIR}") +file(MAKE_DIRECTORY "${_hpcc_root}/include/hcr" "${_maca_root}/include/mcr") +file(WRITE "${_hpcc_root}/include/hcr/hc_runtime.h" "") +file(WRITE "${_maca_root}/include/mcr/mc_runtime.h" "") + +execute_process( + COMMAND "${CMAKE_COMMAND}" -E env + "HPCC_PATH=${_hpcc_root}" + "MACA_PATH=${_maca_root}" + "${CMAKE_COMMAND}" + -S "${INFINI_RT_SOURCE_DIR}" + -B "${_child_build}" + -DAUTO_DETECT_DEVICES=ON + -DINFINI_RT_BUILD_TESTING=OFF + RESULT_VARIABLE _configure_result + OUTPUT_VARIABLE _configure_stdout + ERROR_VARIABLE _configure_stderr) + +if(_configure_result EQUAL 0) + message(FATAL_ERROR "Ambiguous MetaX-family auto-detection unexpectedly succeeded.") +endif() + +set(_configure_output "${_configure_stdout}\n${_configure_stderr}") +string(FIND + "${_configure_output}" + "Both HPCC_PATH and MACA_PATH identify valid SDKs" + _expected_error_index) +if(_expected_error_index EQUAL -1) + message(FATAL_ERROR + "Ambiguous auto-detection failed for the wrong reason:\n${_configure_output}") +endif() diff --git a/tests/test_runtime_dispatch.cc b/tests/test_runtime_dispatch.cc index be49b3b..75a9ab1 100644 --- a/tests/test_runtime_dispatch.cc +++ b/tests/test_runtime_dispatch.cc @@ -425,6 +425,11 @@ int main() { {true, true, true, true, true, true, true, true}); #endif +#if defined(INFINI_RT_TEST_WITH_MARS) + TestDispatch(&context, infini::rt::Device::Type::kMars, "MARS", + {true, true, true, true, true, true, true, true}); +#endif + #if defined(INFINI_RT_TEST_WITH_MOORE) TestDispatch(&context, infini::rt::Device::Type::kMoore, "MOORE", {true, true, false, true, true, true, true, true});