From 8d788497c086a4d090879124f926a09bea58671b Mon Sep 17 00:00:00 2001 From: gongchensu Date: Tue, 14 Jul 2026 09:08:21 +0000 Subject: [PATCH 1/5] feat(metax): add graph capture and replay support --- docs/backends.md | 2 +- src/native/cuda/metax/runtime_.h | 42 +++++++++++++++++++++++++++++++- tests/CMakeLists.txt | 2 ++ 3 files changed, 44 insertions(+), 2 deletions(-) diff --git a/docs/backends.md b/docs/backends.md index 26ad1ff..8ecb8f8 100644 --- a/docs/backends.md +++ b/docs/backends.md @@ -33,7 +33,7 @@ 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 | | 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/src/native/cuda/metax/runtime_.h b/src/native/cuda/metax/runtime_.h index 462f5f1..63d4d24 100644 --- a/src/native/cuda/metax/runtime_.h +++ b/src/native/cuda/metax/runtime_.h @@ -12,13 +12,20 @@ namespace infini::rt::runtime { template <> struct Runtime - : CudaRuntime> { + : GraphRuntime, + CudaRuntime>> { 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; @@ -124,6 +131,39 @@ struct Runtime 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/tests/CMakeLists.txt b/tests/CMakeLists.txt index 1dd2048..42cdf6b 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -85,6 +85,8 @@ 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_MOORE) From a62bf0bec4a5b6c8be09737a8dc6dfd0e4e62078 Mon Sep 17 00:00:00 2001 From: gongchensu Date: Fri, 17 Jul 2026 14:24:39 +0800 Subject: [PATCH 2/5] feat(metax): support HPCC runtime API --- CMakeLists.txt | 21 +++-- cmake/InfiniRTConfig.cmake.in | 28 ++++-- src/CMakeLists.txt | 3 + src/native/cuda/metax/data_type_.h | 4 + src/native/cuda/metax/runtime_.h | 137 +++++++++++++++++++---------- 5 files changed, 134 insertions(+), 59 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 65e63ec..c23aa65 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -220,16 +220,27 @@ endif() if(WITH_METAX) add_compile_definitions(WITH_METAX=1) - # Normally can be found at: `/opt/maca/`. + # MACA is normally installed under /opt/maca, while HPCC uses /opt/hpcc. set(MACA_PATH $ENV{MACA_PATH}) - set(CMAKE_C_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) - set(CMAKE_CXX_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) + if(EXISTS "${MACA_PATH}/include/hcr/hc_runtime.h" AND + EXISTS "${MACA_PATH}/htgpu_llvm/bin/htcc") + set(INFINI_RT_METAX_USE_HPCC ON) + message(STATUS "MetaX: using HPCC SDK at ${MACA_PATH}") + else() + set(INFINI_RT_METAX_USE_HPCC OFF) + set(CMAKE_C_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) + set(CMAKE_CXX_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) + message(STATUS "MetaX: using MACA SDK at ${MACA_PATH}") + endif() 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) + if(INFINI_RT_METAX_USE_HPCC) + find_library(MACA_RUNTIME_LIB NAMES hcruntime HINTS "${MACA_PATH}/lib" REQUIRED) + else() + find_library(MACA_RUNTIME_LIB NAMES mcruntime HINTS "${MACA_PATH}/lib" REQUIRED) + endif() endif() if(WITH_MOORE) diff --git a/cmake/InfiniRTConfig.cmake.in b/cmake/InfiniRTConfig.cmake.in index 09bce5a..6164059 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_METAX_USE_HPCC @INFINI_RT_METAX_USE_HPCC@) set(InfiniRT_WITH_MOORE @WITH_MOORE@) set(InfiniRT_WITH_CAMBRICON @WITH_CAMBRICON@) set(InfiniRT_WITH_ASCEND @WITH_ASCEND@) @@ -34,14 +35,25 @@ if(_InfiniRT_VENDOR_BACKEND STREQUAL "METAX") if(DEFINED ENV{MACA_PATH}) list(APPEND _InfiniRT_ROOT_HINTS "$ENV{MACA_PATH}") endif() - find_path(_InfiniRT_BACKEND_INCLUDE_DIR - NAMES mcr/mc_runtime.h - HINTS ${_InfiniRT_ROOT_HINTS} - PATH_SUFFIXES include) - find_library(_InfiniRT_BACKEND_RUNTIME_LIBRARY - NAMES mcruntime - HINTS ${_InfiniRT_ROOT_HINTS} - PATH_SUFFIXES lib lib64) + if(InfiniRT_METAX_USE_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) + else() + find_path(_InfiniRT_BACKEND_INCLUDE_DIR + NAMES mcr/mc_runtime.h + HINTS ${_InfiniRT_ROOT_HINTS} + PATH_SUFFIXES include) + find_library(_InfiniRT_BACKEND_RUNTIME_LIBRARY + NAMES mcruntime + HINTS ${_InfiniRT_ROOT_HINTS} + PATH_SUFFIXES lib lib64) + endif() if(NOT _InfiniRT_BACKEND_INCLUDE_DIR OR NOT _InfiniRT_BACKEND_RUNTIME_LIBRARY) set(InfiniRT_FOUND FALSE) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 2091963..8206c02 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -65,6 +65,9 @@ endif() if(WITH_METAX) target_compile_definitions(infinirt PRIVATE WITH_METAX=1) + if(INFINI_RT_METAX_USE_HPCC) + target_compile_definitions(infinirt PUBLIC INFINIRT_METAX_USE_HPCC=1) + endif() _infinirt_use_backend_runtime( "${MACA_PATH}/include" "${MACA_RUNTIME_LIB}") diff --git a/src/native/cuda/metax/data_type_.h b/src/native/cuda/metax/data_type_.h index 88d3e24..b26917c 100644 --- a/src/native/cuda/metax/data_type_.h +++ b/src/native/cuda/metax/data_type_.h @@ -3,7 +3,11 @@ #include #include +#if defined(INFINIRT_METAX_USE_HPCC) || defined(USE_HPCC) +#include +#else #include +#endif #include "data_type.h" #include "native/cuda/metax/device_.h" diff --git a/src/native/cuda/metax/runtime_.h b/src/native/cuda/metax/runtime_.h index 63d4d24..932a587 100644 --- a/src/native/cuda/metax/runtime_.h +++ b/src/native/cuda/metax/runtime_.h @@ -1,7 +1,19 @@ #ifndef INFINI_RT_METAX_RUNTIME__H_ #define INFINI_RT_METAX_RUNTIME__H_ +#if defined(INFINIRT_METAX_USE_HPCC) || defined(USE_HPCC) +#include +#define INFINIRT_METAX_RUNTIME_API_(prefix, name) prefix##name +#define INFINIRT_METAX_RUNTIME_API(prefix, name) \ + INFINIRT_METAX_RUNTIME_API_(prefix, name) +#define INFINIRT_METAX_API(name) INFINIRT_METAX_RUNTIME_API(hc, name) +#else #include +#define INFINIRT_METAX_RUNTIME_API_(prefix, name) prefix##name +#define INFINIRT_METAX_RUNTIME_API(prefix, name) \ + INFINIRT_METAX_RUNTIME_API_(prefix, name) +#define INFINIRT_METAX_API(name) INFINIRT_METAX_RUNTIME_API(mc, name) +#endif #include @@ -14,155 +26,184 @@ template <> struct Runtime : GraphRuntime, CudaRuntime>> { - using Error = mcError_t; + using Error = INFINIRT_METAX_API(Error_t); - using Stream = mcStream_t; + using Stream = INFINIRT_METAX_API(Stream_t); - using Graph = mcGraph_t; + using Graph = INFINIRT_METAX_API(Graph_t); - using GraphExec = mcGraphExec_t; + using GraphExec = INFINIRT_METAX_API(GraphExec_t); - using Event = mcEvent_t; + using Event = INFINIRT_METAX_API(Event_t); - using StreamCaptureMode = mcStreamCaptureMode; + using StreamCaptureMode = INFINIRT_METAX_API(StreamCaptureMode); static constexpr Device::Type kDeviceType = Device::Type::kMetax; - static constexpr Error kSuccess = mcSuccess; + static constexpr Error kSuccess = INFINIRT_METAX_API(Success); - static constexpr auto SetDevice = mcSetDevice; + static constexpr auto SetDevice = INFINIRT_METAX_API(SetDevice); - static constexpr auto GetDevice = mcGetDevice; + static constexpr auto GetDevice = INFINIRT_METAX_API(GetDevice); - static constexpr auto GetDeviceCount = mcGetDeviceCount; + static constexpr auto GetDeviceCount = INFINIRT_METAX_API(GetDeviceCount); - static constexpr auto DeviceSynchronize = mcDeviceSynchronize; + static constexpr auto DeviceSynchronize = + INFINIRT_METAX_API(DeviceSynchronize); static constexpr auto Malloc = [](auto&&... args) { - return mcMalloc(std::forward(args)...); + return INFINIRT_METAX_API(Malloc)(std::forward(args)...); }; static constexpr auto MallocHost = [](auto&&... args) { - return mcMallocHost(std::forward(args)...); + return INFINIRT_METAX_API(MallocHost)( + std::forward(args)...); }; static constexpr auto MallocAsync = [](auto&&... args) { - return mcMallocAsync(std::forward(args)...); + return INFINIRT_METAX_API(MallocAsync)( + std::forward(args)...); }; static constexpr auto Free = [](auto&&... args) { - return mcFree(std::forward(args)...); + return INFINIRT_METAX_API(Free)(std::forward(args)...); }; static constexpr auto FreeHost = [](auto&&... args) { - return mcFreeHost(std::forward(args)...); + return INFINIRT_METAX_API(FreeHost)( + std::forward(args)...); }; static constexpr auto FreeAsync = [](auto&&... args) { - return mcFreeAsync(std::forward(args)...); + return INFINIRT_METAX_API(FreeAsync)( + std::forward(args)...); }; static constexpr auto MemGetInfo = [](auto&&... args) { - return mcMemGetInfo(std::forward(args)...); + return INFINIRT_METAX_API(MemGetInfo)( + std::forward(args)...); }; static constexpr auto Memcpy = [](auto&&... args) { - return mcMemcpy(std::forward(args)...); + return INFINIRT_METAX_API(Memcpy)(std::forward(args)...); }; static constexpr auto MemcpyAsync = [](auto&&... args) { - return mcMemcpyAsync(std::forward(args)...); + return INFINIRT_METAX_API(MemcpyAsync)( + std::forward(args)...); }; - static constexpr auto kMemcpyHostToHost = mcMemcpyHostToHost; + static constexpr auto kMemcpyHostToHost = + INFINIRT_METAX_API(MemcpyHostToHost); - static constexpr auto kMemcpyHostToDevice = mcMemcpyHostToDevice; + static constexpr auto kMemcpyHostToDevice = + INFINIRT_METAX_API(MemcpyHostToDevice); - static constexpr auto kMemcpyDeviceToHost = mcMemcpyDeviceToHost; + static constexpr auto kMemcpyDeviceToHost = + INFINIRT_METAX_API(MemcpyDeviceToHost); - static constexpr auto kMemcpyDeviceToDevice = mcMemcpyDeviceToDevice; + static constexpr auto kMemcpyDeviceToDevice = + INFINIRT_METAX_API(MemcpyDeviceToDevice); - static constexpr auto Memset = mcMemset; + static constexpr auto Memset = INFINIRT_METAX_API(Memset); static constexpr auto MemsetAsync = [](auto&&... args) { - return mcMemsetAsync(std::forward(args)...); + return INFINIRT_METAX_API(MemsetAsync)( + std::forward(args)...); }; static constexpr auto StreamCreate = [](auto&&... args) { - return mcStreamCreate(std::forward(args)...); + return INFINIRT_METAX_API(StreamCreate)( + std::forward(args)...); }; static constexpr auto StreamDestroy = [](auto&&... args) { - return mcStreamDestroy(std::forward(args)...); + return INFINIRT_METAX_API(StreamDestroy)( + std::forward(args)...); }; static constexpr auto StreamSynchronize = [](auto&&... args) { - return mcStreamSynchronize(std::forward(args)...); + return INFINIRT_METAX_API(StreamSynchronize)( + std::forward(args)...); }; static constexpr auto StreamWaitEvent = [](auto&&... args) { - return mcStreamWaitEvent(std::forward(args)...); + return INFINIRT_METAX_API(StreamWaitEvent)( + std::forward(args)...); }; static constexpr auto EventCreate = [](auto&&... args) { - return mcEventCreate(std::forward(args)...); + return INFINIRT_METAX_API(EventCreate)( + std::forward(args)...); }; static constexpr auto EventCreateWithFlags = [](auto&&... args) { - return mcEventCreateWithFlags(std::forward(args)...); + return INFINIRT_METAX_API(EventCreateWithFlags)( + std::forward(args)...); }; static constexpr auto EventRecord = [](auto&&... args) { - return mcEventRecord(std::forward(args)...); + return INFINIRT_METAX_API(EventRecord)( + std::forward(args)...); }; static constexpr auto EventQuery = [](auto&&... args) { - return mcEventQuery(std::forward(args)...); + return INFINIRT_METAX_API(EventQuery)( + std::forward(args)...); }; static constexpr auto EventSynchronize = [](auto&&... args) { - return mcEventSynchronize(std::forward(args)...); + return INFINIRT_METAX_API(EventSynchronize)( + std::forward(args)...); }; static constexpr auto EventDestroy = [](auto&&... args) { - return mcEventDestroy(std::forward(args)...); + return INFINIRT_METAX_API(EventDestroy)( + std::forward(args)...); }; static constexpr auto EventElapsedTime = [](auto&&... args) { - return mcEventElapsedTime(std::forward(args)...); + return INFINIRT_METAX_API(EventElapsedTime)( + std::forward(args)...); }; static constexpr auto kStreamCaptureModeGlobal = - mcStreamCaptureModeGlobal; + INFINIRT_METAX_API(StreamCaptureModeGlobal); static constexpr auto kStreamCaptureModeThreadLocal = - mcStreamCaptureModeThreadLocal; + INFINIRT_METAX_API(StreamCaptureModeThreadLocal); static constexpr auto kStreamCaptureModeRelaxed = - mcStreamCaptureModeRelaxed; + INFINIRT_METAX_API(StreamCaptureModeRelaxed); static constexpr auto StreamBeginCapture = [](auto&&... args) { - return mcStreamBeginCapture(std::forward(args)...); + return INFINIRT_METAX_API(StreamBeginCapture)( + std::forward(args)...); }; static constexpr auto StreamEndCapture = [](auto&&... args) { - return mcStreamEndCapture(std::forward(args)...); + return INFINIRT_METAX_API(StreamEndCapture)( + std::forward(args)...); }; static constexpr auto GraphDestroy = [](auto&&... args) { - return mcGraphDestroy(std::forward(args)...); + return INFINIRT_METAX_API(GraphDestroy)( + std::forward(args)...); }; static Error GraphInstantiate(GraphExec* graph_exec, Graph graph) { - return mcGraphInstantiate(graph_exec, graph, nullptr, nullptr, 0); + return INFINIRT_METAX_API(GraphInstantiate)( + graph_exec, graph, nullptr, nullptr, 0); } static constexpr auto GraphExecDestroy = [](auto&&... args) { - return mcGraphExecDestroy(std::forward(args)...); + return INFINIRT_METAX_API(GraphExecDestroy)( + std::forward(args)...); }; static constexpr auto GraphLaunch = [](auto&&... args) { - return mcGraphLaunch(std::forward(args)...); + return INFINIRT_METAX_API(GraphLaunch)( + std::forward(args)...); }; }; @@ -170,4 +211,8 @@ static_assert(Runtime::Validate()); } // namespace infini::rt::runtime +#undef INFINIRT_METAX_API +#undef INFINIRT_METAX_RUNTIME_API +#undef INFINIRT_METAX_RUNTIME_API_ + #endif From af636869cace44b374c8e01670fcd9b6aeaebb11 Mon Sep 17 00:00:00 2001 From: gongchensu Date: Fri, 17 Jul 2026 16:14:26 +0800 Subject: [PATCH 3/5] style(metax): apply clang-format --- src/native/cuda/metax/runtime_.h | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/src/native/cuda/metax/runtime_.h b/src/native/cuda/metax/runtime_.h index 932a587..41abec1 100644 --- a/src/native/cuda/metax/runtime_.h +++ b/src/native/cuda/metax/runtime_.h @@ -70,13 +70,11 @@ struct Runtime }; static constexpr auto FreeHost = [](auto&&... args) { - return INFINIRT_METAX_API(FreeHost)( - std::forward(args)...); + return INFINIRT_METAX_API(FreeHost)(std::forward(args)...); }; static constexpr auto FreeAsync = [](auto&&... args) { - return INFINIRT_METAX_API(FreeAsync)( - std::forward(args)...); + return INFINIRT_METAX_API(FreeAsync)(std::forward(args)...); }; static constexpr auto MemGetInfo = [](auto&&... args) { @@ -192,8 +190,8 @@ struct Runtime }; static Error GraphInstantiate(GraphExec* graph_exec, Graph graph) { - return INFINIRT_METAX_API(GraphInstantiate)( - graph_exec, graph, nullptr, nullptr, 0); + return INFINIRT_METAX_API(GraphInstantiate)(graph_exec, graph, nullptr, + nullptr, 0); } static constexpr auto GraphExecDestroy = [](auto&&... args) { From 1ddb3e15a548d98b9b173e8f0b601c91ae037896 Mon Sep 17 00:00:00 2001 From: gongchensu Date: Tue, 21 Jul 2026 16:04:53 +0800 Subject: [PATCH 4/5] refactor: split MetaX and Mars runtime backends --- CMakeLists.txt | 66 +++--- cmake/InfiniRTConfig.cmake.in | 54 +++-- docs/backends.md | 2 + docs/build.md | 1 + scripts/generate_public_headers.py | 17 +- scripts/run_performance_tests.py | 1 + src/CMakeLists.txt | 12 +- src/device.h | 5 +- src/native/cuda/metax/runtime_.h | 216 ------------------- src/native/hpcc/mars/data_type_.h | 25 +++ src/native/hpcc/mars/device_.h | 13 ++ src/native/hpcc/mars/runtime_.h | 131 +++++++++++ src/native/hpcc/runtime_.h | 29 +++ src/native/{cuda => maca}/metax/data_type_.h | 10 +- src/native/{cuda => maca}/metax/device_.h | 0 src/native/maca/metax/runtime_.h | 131 +++++++++++ src/native/maca/runtime_.h | 29 +++ tests/CMakeLists.txt | 18 ++ tests/performance/CMakeLists.txt | 2 + tests/test_runtime_dispatch.cc | 5 + 20 files changed, 490 insertions(+), 277 deletions(-) delete mode 100644 src/native/cuda/metax/runtime_.h create mode 100644 src/native/hpcc/mars/data_type_.h create mode 100644 src/native/hpcc/mars/device_.h create mode 100644 src/native/hpcc/mars/runtime_.h create mode 100644 src/native/hpcc/runtime_.h rename src/native/{cuda => maca}/metax/data_type_.h (68%) rename src/native/{cuda => maca}/metax/device_.h (100%) create mode 100644 src/native/maca/metax/runtime_.h create mode 100644 src/native/maca/runtime_.h diff --git a/CMakeLists.txt b/CMakeLists.txt index c23aa65..58891e8 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,15 @@ if(AUTO_DETECT_DEVICES) endif() endif() - if(DEFINED ENV{MACA_PATH}) + if((DEFINED ENV{HPCC_PATH} AND + EXISTS "$ENV{HPCC_PATH}/include/hcr/hc_runtime.h") OR + EXISTS "/opt/hpcc/include/hcr/hc_runtime.h") + set(WITH_MARS ON) + message(STATUS "Auto-detected Mars HPCC environment") + elseif(DEFINED ENV{MACA_PATH} AND + EXISTS "$ENV{MACA_PATH}/include/mcr/mc_runtime.h") set(WITH_METAX ON) - message(STATUS "Auto-detected MetaX environment from MACA_PATH") + message(STATUS "Auto-detected MetaX MACA environment from MACA_PATH") else() execute_process( COMMAND sh -c "grep -h 9999 /sys/bus/pci/devices/*/vendor 2>/dev/null" @@ -89,8 +96,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 +124,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 +141,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,27 +226,36 @@ endif() if(WITH_METAX) add_compile_definitions(WITH_METAX=1) - # MACA is normally installed under /opt/maca, while HPCC uses /opt/hpcc. - set(MACA_PATH $ENV{MACA_PATH}) - if(EXISTS "${MACA_PATH}/include/hcr/hc_runtime.h" AND - EXISTS "${MACA_PATH}/htgpu_llvm/bin/htcc") - set(INFINI_RT_METAX_USE_HPCC ON) - message(STATUS "MetaX: using HPCC SDK at ${MACA_PATH}") - else() - set(INFINI_RT_METAX_USE_HPCC OFF) - set(CMAKE_C_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) - set(CMAKE_CXX_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/mxcc_wrapper.sh) - message(STATUS "MetaX: using MACA SDK at ${MACA_PATH}") + 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") + find_library(MACA_RUNTIME_LIB NAMES mcruntime HINTS "${MACA_PATH}/lib" REQUIRED) +endif() - if(INFINI_RT_METAX_USE_HPCC) - find_library(MACA_RUNTIME_LIB NAMES hcruntime HINTS "${MACA_PATH}/lib" REQUIRED) - else() - find_library(MACA_RUNTIME_LIB NAMES mcruntime HINTS "${MACA_PATH}/lib" REQUIRED) +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) @@ -305,7 +320,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() @@ -326,6 +341,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() @@ -358,7 +376,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/cmake/InfiniRTConfig.cmake.in b/cmake/InfiniRTConfig.cmake.in index 6164059..137cdba 100644 --- a/cmake/InfiniRTConfig.cmake.in +++ b/cmake/InfiniRTConfig.cmake.in @@ -10,7 +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_METAX_USE_HPCC @INFINI_RT_METAX_USE_HPCC@) +set(InfiniRT_WITH_MARS @WITH_MARS@) set(InfiniRT_WITH_MOORE @WITH_MOORE@) set(InfiniRT_WITH_CAMBRICON @WITH_CAMBRICON@) set(InfiniRT_WITH_ASCEND @WITH_ASCEND@) @@ -35,25 +35,15 @@ if(_InfiniRT_VENDOR_BACKEND STREQUAL "METAX") if(DEFINED ENV{MACA_PATH}) list(APPEND _InfiniRT_ROOT_HINTS "$ENV{MACA_PATH}") endif() - if(InfiniRT_METAX_USE_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) - else() - find_path(_InfiniRT_BACKEND_INCLUDE_DIR - NAMES mcr/mc_runtime.h - HINTS ${_InfiniRT_ROOT_HINTS} - PATH_SUFFIXES include) - find_library(_InfiniRT_BACKEND_RUNTIME_LIBRARY - NAMES mcruntime - HINTS ${_InfiniRT_ROOT_HINTS} - PATH_SUFFIXES lib lib64) - endif() + list(APPEND _InfiniRT_ROOT_HINTS /opt/maca) + find_path(_InfiniRT_BACKEND_INCLUDE_DIR + NAMES mcr/mc_runtime.h + HINTS ${_InfiniRT_ROOT_HINTS} + PATH_SUFFIXES include) + find_library(_InfiniRT_BACKEND_RUNTIME_LIBRARY + NAMES mcruntime + HINTS ${_InfiniRT_ROOT_HINTS} + PATH_SUFFIXES lib lib64) if(NOT _InfiniRT_BACKEND_INCLUDE_DIR OR NOT _InfiniRT_BACKEND_RUNTIME_LIBRARY) set(InfiniRT_FOUND FALSE) @@ -64,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/backends.md b/docs/backends.md index 8ecb8f8..3796816 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`. | @@ -34,6 +35,7 @@ Current test expectations are: | NVIDIA | Yes | Yes | Yes | Yes | Yes | Yes | Yes | | Iluvatar | Yes | Yes | Yes | Yes | Yes | Yes | Yes | | 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 8206c02..f755ca6 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -64,15 +64,19 @@ if(WITH_HYGON) endif() if(WITH_METAX) - target_compile_definitions(infinirt PRIVATE WITH_METAX=1) - if(INFINI_RT_METAX_USE_HPCC) - target_compile_definitions(infinirt PUBLIC INFINIRT_METAX_USE_HPCC=1) - endif() + 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/cuda/metax/runtime_.h b/src/native/cuda/metax/runtime_.h deleted file mode 100644 index 41abec1..0000000 --- a/src/native/cuda/metax/runtime_.h +++ /dev/null @@ -1,216 +0,0 @@ -#ifndef INFINI_RT_METAX_RUNTIME__H_ -#define INFINI_RT_METAX_RUNTIME__H_ - -#if defined(INFINIRT_METAX_USE_HPCC) || defined(USE_HPCC) -#include -#define INFINIRT_METAX_RUNTIME_API_(prefix, name) prefix##name -#define INFINIRT_METAX_RUNTIME_API(prefix, name) \ - INFINIRT_METAX_RUNTIME_API_(prefix, name) -#define INFINIRT_METAX_API(name) INFINIRT_METAX_RUNTIME_API(hc, name) -#else -#include -#define INFINIRT_METAX_RUNTIME_API_(prefix, name) prefix##name -#define INFINIRT_METAX_RUNTIME_API(prefix, name) \ - INFINIRT_METAX_RUNTIME_API_(prefix, name) -#define INFINIRT_METAX_API(name) INFINIRT_METAX_RUNTIME_API(mc, name) -#endif - -#include - -#include "native/cuda/metax/device_.h" -#include "native/cuda/runtime_.h" - -namespace infini::rt::runtime { - -template <> -struct Runtime - : GraphRuntime, - CudaRuntime>> { - using Error = INFINIRT_METAX_API(Error_t); - - using Stream = INFINIRT_METAX_API(Stream_t); - - using Graph = INFINIRT_METAX_API(Graph_t); - - using GraphExec = INFINIRT_METAX_API(GraphExec_t); - - using Event = INFINIRT_METAX_API(Event_t); - - using StreamCaptureMode = INFINIRT_METAX_API(StreamCaptureMode); - - static constexpr Device::Type kDeviceType = Device::Type::kMetax; - - static constexpr Error kSuccess = INFINIRT_METAX_API(Success); - - static constexpr auto SetDevice = INFINIRT_METAX_API(SetDevice); - - static constexpr auto GetDevice = INFINIRT_METAX_API(GetDevice); - - static constexpr auto GetDeviceCount = INFINIRT_METAX_API(GetDeviceCount); - - static constexpr auto DeviceSynchronize = - INFINIRT_METAX_API(DeviceSynchronize); - - static constexpr auto Malloc = [](auto&&... args) { - return INFINIRT_METAX_API(Malloc)(std::forward(args)...); - }; - - static constexpr auto MallocHost = [](auto&&... args) { - return INFINIRT_METAX_API(MallocHost)( - std::forward(args)...); - }; - - static constexpr auto MallocAsync = [](auto&&... args) { - return INFINIRT_METAX_API(MallocAsync)( - std::forward(args)...); - }; - - static constexpr auto Free = [](auto&&... args) { - return INFINIRT_METAX_API(Free)(std::forward(args)...); - }; - - static constexpr auto FreeHost = [](auto&&... args) { - return INFINIRT_METAX_API(FreeHost)(std::forward(args)...); - }; - - static constexpr auto FreeAsync = [](auto&&... args) { - return INFINIRT_METAX_API(FreeAsync)(std::forward(args)...); - }; - - static constexpr auto MemGetInfo = [](auto&&... args) { - return INFINIRT_METAX_API(MemGetInfo)( - std::forward(args)...); - }; - - static constexpr auto Memcpy = [](auto&&... args) { - return INFINIRT_METAX_API(Memcpy)(std::forward(args)...); - }; - - static constexpr auto MemcpyAsync = [](auto&&... args) { - return INFINIRT_METAX_API(MemcpyAsync)( - std::forward(args)...); - }; - - static constexpr auto kMemcpyHostToHost = - INFINIRT_METAX_API(MemcpyHostToHost); - - static constexpr auto kMemcpyHostToDevice = - INFINIRT_METAX_API(MemcpyHostToDevice); - - static constexpr auto kMemcpyDeviceToHost = - INFINIRT_METAX_API(MemcpyDeviceToHost); - - static constexpr auto kMemcpyDeviceToDevice = - INFINIRT_METAX_API(MemcpyDeviceToDevice); - - static constexpr auto Memset = INFINIRT_METAX_API(Memset); - - static constexpr auto MemsetAsync = [](auto&&... args) { - return INFINIRT_METAX_API(MemsetAsync)( - std::forward(args)...); - }; - - static constexpr auto StreamCreate = [](auto&&... args) { - return INFINIRT_METAX_API(StreamCreate)( - std::forward(args)...); - }; - - static constexpr auto StreamDestroy = [](auto&&... args) { - return INFINIRT_METAX_API(StreamDestroy)( - std::forward(args)...); - }; - - static constexpr auto StreamSynchronize = [](auto&&... args) { - return INFINIRT_METAX_API(StreamSynchronize)( - std::forward(args)...); - }; - - static constexpr auto StreamWaitEvent = [](auto&&... args) { - return INFINIRT_METAX_API(StreamWaitEvent)( - std::forward(args)...); - }; - - static constexpr auto EventCreate = [](auto&&... args) { - return INFINIRT_METAX_API(EventCreate)( - std::forward(args)...); - }; - - static constexpr auto EventCreateWithFlags = [](auto&&... args) { - return INFINIRT_METAX_API(EventCreateWithFlags)( - std::forward(args)...); - }; - - static constexpr auto EventRecord = [](auto&&... args) { - return INFINIRT_METAX_API(EventRecord)( - std::forward(args)...); - }; - - static constexpr auto EventQuery = [](auto&&... args) { - return INFINIRT_METAX_API(EventQuery)( - std::forward(args)...); - }; - - static constexpr auto EventSynchronize = [](auto&&... args) { - return INFINIRT_METAX_API(EventSynchronize)( - std::forward(args)...); - }; - - static constexpr auto EventDestroy = [](auto&&... args) { - return INFINIRT_METAX_API(EventDestroy)( - std::forward(args)...); - }; - - static constexpr auto EventElapsedTime = [](auto&&... args) { - return INFINIRT_METAX_API(EventElapsedTime)( - std::forward(args)...); - }; - - static constexpr auto kStreamCaptureModeGlobal = - INFINIRT_METAX_API(StreamCaptureModeGlobal); - - static constexpr auto kStreamCaptureModeThreadLocal = - INFINIRT_METAX_API(StreamCaptureModeThreadLocal); - - static constexpr auto kStreamCaptureModeRelaxed = - INFINIRT_METAX_API(StreamCaptureModeRelaxed); - - static constexpr auto StreamBeginCapture = [](auto&&... args) { - return INFINIRT_METAX_API(StreamBeginCapture)( - std::forward(args)...); - }; - - static constexpr auto StreamEndCapture = [](auto&&... args) { - return INFINIRT_METAX_API(StreamEndCapture)( - std::forward(args)...); - }; - - static constexpr auto GraphDestroy = [](auto&&... args) { - return INFINIRT_METAX_API(GraphDestroy)( - std::forward(args)...); - }; - - static Error GraphInstantiate(GraphExec* graph_exec, Graph graph) { - return INFINIRT_METAX_API(GraphInstantiate)(graph_exec, graph, nullptr, - nullptr, 0); - } - - static constexpr auto GraphExecDestroy = [](auto&&... args) { - return INFINIRT_METAX_API(GraphExecDestroy)( - std::forward(args)...); - }; - - static constexpr auto GraphLaunch = [](auto&&... args) { - return INFINIRT_METAX_API(GraphLaunch)( - std::forward(args)...); - }; -}; - -static_assert(Runtime::Validate()); - -} // namespace infini::rt::runtime - -#undef INFINIRT_METAX_API -#undef INFINIRT_METAX_RUNTIME_API -#undef INFINIRT_METAX_RUNTIME_API_ - -#endif 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 68% rename from src/native/cuda/metax/data_type_.h rename to src/native/maca/metax/data_type_.h index b26917c..7e08eff 100644 --- a/src/native/cuda/metax/data_type_.h +++ b/src/native/maca/metax/data_type_.h @@ -3,21 +3,13 @@ #include #include -#if defined(INFINIRT_METAX_USE_HPCC) || defined(USE_HPCC) -#include -#else #include -#endif #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/maca/metax/runtime_.h b/src/native/maca/metax/runtime_.h new file mode 100644 index 0000000..1982fe8 --- /dev/null +++ b/src/native/maca/metax/runtime_.h @@ -0,0 +1,131 @@ +#ifndef INFINI_RT_METAX_RUNTIME__H_ +#define INFINI_RT_METAX_RUNTIME__H_ + +#include + +#include + +#include "native/maca/metax/device_.h" +#include "native/maca/runtime_.h" + +namespace infini::rt::runtime { + +template <> +struct Runtime + : 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()); + +} // namespace infini::rt::runtime + +#endif 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 42cdf6b..c59f65e 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -89,6 +89,16 @@ if(WITH_METAX) 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) set(INFINI_RT_TEST_HAS_RUNTIME_BACKEND ON) add_infini_rt_backend_runtime_test( @@ -136,6 +146,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) @@ -178,6 +192,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/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_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}); From 671412fd4b8f069758fbc334cbb4f68ebf702e09 Mon Sep 17 00:00:00 2001 From: gongchensu Date: Thu, 23 Jul 2026 09:40:17 +0800 Subject: [PATCH 5/5] fix(mars): harden backend detection and package tests Distinguish HPCC and MACA SDK signals during automatic backend selection. Reject ambiguous hosts unless WITH_MARS or WITH_METAX is selected explicitly, while preserving PCI fallback when neither SDK is present. Expose Mars through installed package consumer metadata, document the device selection rules, and cover dual-SDK detection with a CMake regression test. --- CMakeLists.txt | 45 ++++++++++++++++--- CONTRIBUTING.md | 4 +- docs/api/core-types.md | 1 + docs/backends.md | 3 ++ tests/CMakeLists.txt | 7 +++ tests/install_consumer/CMakeLists.txt | 2 +- tests/install_consumer_smoke.cc | 4 ++ ...st_ambiguous_metax_family_autodetect.cmake | 38 ++++++++++++++++ 8 files changed, 96 insertions(+), 8 deletions(-) create mode 100644 tests/test_ambiguous_metax_family_autodetect.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index 58891e8..8982717 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -77,15 +77,48 @@ if(AUTO_DETECT_DEVICES) endif() endif() - if((DEFINED ENV{HPCC_PATH} AND - EXISTS "$ENV{HPCC_PATH}/include/hcr/hc_runtime.h") OR - EXISTS "/opt/hpcc/include/hcr/hc_runtime.h") + 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") - elseif(DEFINED ENV{MACA_PATH} AND - EXISTS "$ENV{MACA_PATH}/include/mcr/mc_runtime.h") + message(STATUS "Auto-detected Mars HPCC environment from HPCC_PATH") + elseif(_infinirt_maca_env_root) set(WITH_METAX ON) 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" 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/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 3796816..a55566c 100644 --- a/docs/backends.md +++ b/docs/backends.md @@ -21,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 diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index c59f65e..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) 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/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()