From 51d9134b46f1423265a69c94b4065580194561aa Mon Sep 17 00:00:00 2001 From: Ceng23333 <441651826@qq.com> Date: Wed, 19 Aug 2026 12:49:05 +0800 Subject: [PATCH 1/2] feat(mars): add native Mars kernels and htcc wrapper Port MetaX CUDA kernels to native/cuda/mars, compile .cc via htcc_wrapper, and link HPCC libhtomp so InfiniOps can build WITH_MARS. --- CMakeLists.txt | 40 ++++++++-- examples/runtime_api.h | 5 ++ scripts/generate_wrappers.py | 4 +- scripts/htcc_wrapper.sh | 56 +++++++++++++ src/CMakeLists.txt | 32 ++++++++ src/device.h | 2 +- src/native/cuda/mars/blas.h | 51 ++++++++++++ src/native/cuda/mars/blas_utils.h | 30 +++++++ src/native/cuda/mars/caster.cuh | 78 +++++++++++++++++++ src/native/cuda/mars/data_type_.h | 13 ++++ src/native/cuda/mars/device_.h | 6 ++ src/native/cuda/mars/device_property.h | 11 +++ src/native/cuda/mars/ops/add/kernel.h | 21 +++++ .../cuda/mars/ops/add_rms_norm/kernel.h | 21 +++++ .../cuda/mars/ops/causal_softmax/kernel.h | 21 +++++ .../mars/ops/causal_softmax_infinilm/kernel.h | 22 ++++++ src/native/cuda/mars/ops/conv1d/kernel.h | 21 +++++ src/native/cuda/mars/ops/conv2d/kernel.h | 21 +++++ src/native/cuda/mars/ops/conv3d/kernel.h | 21 +++++ .../cuda/mars/ops/conv_infinilm/kernel.h | 21 +++++ src/native/cuda/mars/ops/convolution/kernel.h | 21 +++++ src/native/cuda/mars/ops/copy/kernel.h | 21 +++++ src/native/cuda/mars/ops/embedding/kernel.h | 21 +++++ src/native/cuda/mars/ops/fill/kernel.h | 21 +++++ .../cuda/mars/ops/fused_add_rms_norm/kernel.h | 21 +++++ src/native/cuda/mars/ops/gelu/kernel.h | 21 +++++ .../cuda/mars/ops/gelu_infinilm/kernel.h | 21 +++++ .../cuda/mars/ops/gelutanh_infinilm/kernel.h | 22 ++++++ src/native/cuda/mars/ops/gemm/hcblas.h | 18 +++++ .../mars/ops/kv_caching_infinilm/kernel.h | 22 ++++++ src/native/cuda/mars/ops/mul/kernel.h | 21 +++++ .../ops/paged_attention_infinilm/kernel.h | 22 ++++++ .../paged_attention_prefill_infinilm/kernel.h | 22 ++++++ .../mars/ops/paged_caching_infinilm/kernel.h | 22 ++++++ .../mars/ops/random_sample_infinilm/kernel.h | 22 ++++++ .../cuda/mars/ops/rearrange_infinilm/kernel.h | 22 ++++++ src/native/cuda/mars/ops/relu/kernel.h | 21 +++++ .../cuda/mars/ops/relu_infinilm/kernel.h | 21 +++++ .../mars/ops/reshape_and_cache_flash/kernel.h | 22 ++++++ src/native/cuda/mars/ops/rms_norm/kernel.h | 21 +++++ .../cuda/mars/ops/rotary_embedding/kernel.h | 21 +++++ .../ops/rotary_embedding_infinilm/kernel.h | 22 ++++++ src/native/cuda/mars/ops/sigmoid/kernel.h | 21 +++++ .../cuda/mars/ops/sigmoid_infinilm/kernel.h | 21 +++++ src/native/cuda/mars/ops/silu/kernel.h | 21 +++++ .../cuda/mars/ops/silu_and_mul/kernel.h | 21 +++++ .../mars/ops/silu_and_mul_infinilm/kernel.h | 22 ++++++ src/native/cuda/mars/ops/softmax/kernel.h | 21 +++++ .../cuda/mars/ops/softmax_infinilm/kernel.h | 21 +++++ src/native/cuda/mars/ops/swiglu/kernel.h | 21 +++++ .../cuda/mars/ops/topk_softmax/kernel.h | 21 +++++ .../mars/ops/topksoftmax_infinilm/kernel.h | 22 ++++++ .../cuda/mars/ops/zeros_infinilm/kernel.h | 21 +++++ src/native/cuda/mars/runtime_.h | 9 +++ src/native/cuda/mars/runtime_utils.h | 15 ++++ src/pybind11_utils.h | 1 + 56 files changed, 1215 insertions(+), 8 deletions(-) create mode 100755 scripts/htcc_wrapper.sh create mode 100644 src/native/cuda/mars/blas.h create mode 100644 src/native/cuda/mars/blas_utils.h create mode 100644 src/native/cuda/mars/caster.cuh create mode 100644 src/native/cuda/mars/data_type_.h create mode 100644 src/native/cuda/mars/device_.h create mode 100644 src/native/cuda/mars/device_property.h create mode 100644 src/native/cuda/mars/ops/add/kernel.h create mode 100644 src/native/cuda/mars/ops/add_rms_norm/kernel.h create mode 100644 src/native/cuda/mars/ops/causal_softmax/kernel.h create mode 100644 src/native/cuda/mars/ops/causal_softmax_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/conv1d/kernel.h create mode 100644 src/native/cuda/mars/ops/conv2d/kernel.h create mode 100644 src/native/cuda/mars/ops/conv3d/kernel.h create mode 100644 src/native/cuda/mars/ops/conv_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/convolution/kernel.h create mode 100644 src/native/cuda/mars/ops/copy/kernel.h create mode 100644 src/native/cuda/mars/ops/embedding/kernel.h create mode 100644 src/native/cuda/mars/ops/fill/kernel.h create mode 100644 src/native/cuda/mars/ops/fused_add_rms_norm/kernel.h create mode 100644 src/native/cuda/mars/ops/gelu/kernel.h create mode 100644 src/native/cuda/mars/ops/gelu_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/gelutanh_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/gemm/hcblas.h create mode 100644 src/native/cuda/mars/ops/kv_caching_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/mul/kernel.h create mode 100644 src/native/cuda/mars/ops/paged_attention_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/paged_attention_prefill_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/paged_caching_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/random_sample_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/rearrange_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/relu/kernel.h create mode 100644 src/native/cuda/mars/ops/relu_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/reshape_and_cache_flash/kernel.h create mode 100644 src/native/cuda/mars/ops/rms_norm/kernel.h create mode 100644 src/native/cuda/mars/ops/rotary_embedding/kernel.h create mode 100644 src/native/cuda/mars/ops/rotary_embedding_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/sigmoid/kernel.h create mode 100644 src/native/cuda/mars/ops/sigmoid_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/silu/kernel.h create mode 100644 src/native/cuda/mars/ops/silu_and_mul/kernel.h create mode 100644 src/native/cuda/mars/ops/silu_and_mul_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/softmax/kernel.h create mode 100644 src/native/cuda/mars/ops/softmax_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/swiglu/kernel.h create mode 100644 src/native/cuda/mars/ops/topk_softmax/kernel.h create mode 100644 src/native/cuda/mars/ops/topksoftmax_infinilm/kernel.h create mode 100644 src/native/cuda/mars/ops/zeros_infinilm/kernel.h create mode 100644 src/native/cuda/mars/runtime_.h create mode 100644 src/native/cuda/mars/runtime_utils.h diff --git a/CMakeLists.txt b/CMakeLists.txt index e9a0cac20..636748aea 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -17,6 +17,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) @@ -114,7 +115,10 @@ if(AUTO_DETECT_DEVICES) message(STATUS "Auto-detected Hygon environment.") endif() - if(DEFINED ENV{MACA_PATH}) + if(DEFINED ENV{HPCC_PATH}) + set(WITH_MARS ON) + message(STATUS "Auto-detected Mars environment from HPCC_PATH") + elseif(DEFINED ENV{MACA_PATH}) set(WITH_METAX ON) message(STATUS "Auto-detected MetaX environment from MACA_PATH") else() @@ -156,7 +160,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() @@ -328,14 +332,14 @@ endif() # 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_NINETOOTHED AND NOT WITH_NVIDIA) @@ -472,6 +476,30 @@ if(WITH_METAX) find_library(MACA_BLAS_LIB NAMES mcblas HINTS "${MACA_PATH}/lib" REQUIRED) endif() +if(WITH_MARS) + add_compile_definitions(WITH_MARS=1) + + # Normally can be found at: `/opt/hpcc/`. + set(HPCC_PATH $ENV{HPCC_PATH}) + if(NOT HPCC_PATH) + set(HPCC_PATH "/opt/hpcc") + endif() + set(CMAKE_C_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/htcc_wrapper.sh) + set(CMAKE_CXX_COMPILER ${CMAKE_CURRENT_SOURCE_DIR}/scripts/htcc_wrapper.sh) + + include_directories("${HPCC_PATH}/include") + include_directories("${HPCC_PATH}/tools/cu-bridge/include") + link_directories("${HPCC_PATH}/lib") + link_directories("${HPCC_PATH}/htgpu_llvm/lib") + + # Libraries: hcruntime / hcdnn / hcblas. HPCC is the toolkit; Mars is the device. + find_library(HPCC_RUNTIME_LIB NAMES hcruntime HINTS "${HPCC_PATH}/lib" REQUIRED) + find_library(HPCC_DNN_LIB NAMES hcdnn HINTS "${HPCC_PATH}/lib" REQUIRED) + find_library(HPCC_BLAS_LIB NAMES hcblas HINTS "${HPCC_PATH}/lib" REQUIRED) + # htcc emits LLVM OpenMP (__kmpc_*), not GNU libgomp. + find_library(HPCC_OMP_LIB NAMES omp iomp5 HINTS "${HPCC_PATH}/htgpu_llvm/lib" REQUIRED) +endif() + if(WITH_MOORE) add_compile_definitions(WITH_MOORE=1) @@ -539,11 +567,11 @@ 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) add_compile_definitions(WITH_CPU=1) endif() -if(WITH_TORCH OR WITH_METAX OR WITH_MOORE) +if(WITH_TORCH OR WITH_METAX OR WITH_MARS OR WITH_MOORE) set(PYBIND11_ENABLE_EXTRAS OFF) endif() diff --git a/examples/runtime_api.h b/examples/runtime_api.h index 1bbfdc1f3..7f754026d 100644 --- a/examples/runtime_api.h +++ b/examples/runtime_api.h @@ -14,6 +14,9 @@ #elif WITH_METAX #include "native/cuda/metax/ops/gemm/mcblas.h" #include "native/cuda/metax/runtime_.h" +#elif WITH_MARS +#include "native/cuda/mars/ops/gemm/hcblas.h" +#include "native/cuda/mars/runtime_.h" #elif WITH_CAMBRICON #include "native/cambricon/ops/gemm/cnblas.h" #include "native/cambricon/runtime_.h" @@ -39,6 +42,8 @@ using DefaultRuntimeUtils = Runtime; using DefaultRuntimeUtils = Runtime; #elif WITH_METAX using DefaultRuntimeUtils = Runtime; +#elif WITH_MARS +using DefaultRuntimeUtils = Runtime; #elif WITH_CAMBRICON using DefaultRuntimeUtils = Runtime; #elif WITH_MOORE diff --git a/scripts/generate_wrappers.py b/scripts/generate_wrappers.py index 0a188b2a0..6cf1a8188 100644 --- a/scripts/generate_wrappers.py +++ b/scripts/generate_wrappers.py @@ -1721,8 +1721,10 @@ def _device_marker_headers(devices): "cambricon": "infini/rt/cambricon/device_.h", "ascend": "infini/rt/ascend/device_.h", "metax": "infini/rt/metax/device_.h", + "mars": "infini/rt/mars/device_.h", "moore": "infini/rt/moore/device_.h", "iluvatar": "infini/rt/iluvatar/device_.h", + "hygon": "infini/rt/hygon/device_.h", } return [paths[device] for device in devices if device in paths] @@ -2108,7 +2110,7 @@ def _dispatch_gen_batch_size(): nargs="+", default="cpu", type=str, - help="Devices to use. Please pick from `cpu`, `nvidia`, `cambricon`, `ascend`, `metax`, `moore`, `iluvatar`, and `hygon`. (default: `cpu`)", + help="Devices to use. Please pick from `cpu`, `nvidia`, `cambricon`, `ascend`, `metax`, `mars`, `moore`, `iluvatar`, and `hygon`. (default: `cpu`)", ) parser.add_argument( diff --git a/scripts/htcc_wrapper.sh b/scripts/htcc_wrapper.sh new file mode 100755 index 000000000..19d11743f --- /dev/null +++ b/scripts/htcc_wrapper.sh @@ -0,0 +1,56 @@ +#!/bin/bash +# HPCC analog of mxcc_wrapper.sh. +# htcc only enables device compilation for `.cu`. Rewrite compile inputs. +ARGS=() +skip_next=0 +has_compile=0 +for arg in "$@"; do + if [ $skip_next -eq 1 ]; then + skip_next=0 + ARGS+=("$arg") + continue + fi + case "$arg" in + -pthread) + ;; + -B) + skip_next=1 + ;; + -B*) + ;; + -c) + has_compile=1 + ARGS+=("$arg") + ;; + *) + ARGS+=("$arg") + ;; + esac +done + +if [ "$has_compile" -eq 1 ]; then + rewritten=() + mkdir -p /tmp/htcc-wrap + for arg in "${ARGS[@]}"; do + case "$arg" in + *.cc|*.cpp|*.cxx) + if [ -f "$arg" ]; then + abs=$(readlink -f "$arg") + key=$(printf '%s' "$abs" | md5sum | awk '{print $1}') + cu="/tmp/htcc-wrap/${key}.cu" + ln -sfn "$abs" "$cu" + rewritten+=("$cu") + else + rewritten+=("$arg") + fi + ;; + *) + rewritten+=("$arg") + ;; + esac + done + ARGS=("${rewritten[@]}") +fi + +HPCC_PATH="${HPCC_PATH:-/opt/hpcc}" +exec "${HPCC_PATH}/htgpu_llvm/bin/htcc" "${ARGS[@]}" diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 36f946d8d..8f33d1dee 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -291,6 +291,38 @@ if(WITH_METAX) list(APPEND DEVICE_LIST "metax") endif() +if(WITH_MARS) + set(MARS_PATTERNS + "native/cuda/*.cc" + "native/cuda/*.cpp" + "native/cuda/mars/*.cc" + "native/cuda/mars/*.cu" + ) + + file(GLOB_RECURSE MARS_SOURCES CONFIGURE_DEPENDS ${MARS_PATTERNS}) + + set_source_files_properties(${MARS_SOURCES} PROPERTIES LANGUAGE CXX) + + target_compile_definitions(infiniops PRIVATE WITH_MARS=1) + target_sources(infiniops PRIVATE ${MARS_SOURCES}) + + target_include_directories(infiniops PUBLIC + "${HPCC_PATH}/include" + "${HPCC_PATH}/tools/cu-bridge/include") + target_link_libraries(infiniops PUBLIC + ${HPCC_RUNTIME_LIB} + ${HPCC_DNN_LIB} + ${HPCC_BLAS_LIB} + ${HPCC_OMP_LIB} + ) + set_property(TARGET infiniops APPEND PROPERTY + INSTALL_RPATH "${HPCC_PATH}/htgpu_llvm/lib") + set_property(TARGET infiniops APPEND PROPERTY + BUILD_RPATH "${HPCC_PATH}/htgpu_llvm/lib") + + list(APPEND DEVICE_LIST "mars") +endif() + if(WITH_MOORE) set(MOORE_PATTERNS "native/cuda/*.cc" diff --git a/src/device.h b/src/device.h index 59bbeb9c3..e049d0f5b 100644 --- a/src/device.h +++ b/src/device.h @@ -15,7 +15,7 @@ using DeviceEnabled = infini::rt::DeviceEnabled; using AllDeviceTypes = List; + Device::Type::kIluvatar, Device::Type::kHygon, Device::Type::kMars>; template struct ActiveDevicesImpl { diff --git a/src/native/cuda/mars/blas.h b/src/native/cuda/mars/blas.h new file mode 100644 index 000000000..25ced1dd0 --- /dev/null +++ b/src/native/cuda/mars/blas.h @@ -0,0 +1,51 @@ +#ifndef INFINI_OPS_MARS_BLAS_H_ +#define INFINI_OPS_MARS_BLAS_H_ + +#include + +// clang-format off +#include +// clang-format on + +#include "data_type.h" +#include "native/cuda/blas.h" +#include "native/cuda/mars/blas_utils.h" +#include "native/cuda/mars/runtime_.h" + +namespace infini::ops { + +template <> +struct Blas : public Runtime { + using BlasHandle = hcblasHandle_t; + + static constexpr auto BLAS_OP_N = HCBLAS_OP_N; + + static constexpr auto BLAS_OP_T = HCBLAS_OP_T; + + static constexpr auto R_16F = HPCC_R_16F; + + static constexpr auto R_16BF = HPCC_R_16BF; + + static constexpr auto R_32F = HPCC_R_32F; + + static constexpr auto BLAS_COMPUTE_32F = HCBLAS_COMPUTE_32F; + + static constexpr auto BLAS_COMPUTE_32F_FAST_TF32 = + HCBLAS_COMPUTE_32F_FAST_TF32; + + static constexpr auto BLAS_GEMM_DEFAULT = HCBLAS_GEMM_DEFAULT; + + static constexpr auto BlasCreate = hcblasCreate; + + static constexpr auto BlasSetStream = hcblasSetStream; + + static constexpr auto BlasDestroy = hcblasDestroy; + + static constexpr auto BlasGemmStridedBatchedEx = [](auto&&... args) { + return hcblasGemmStridedBatchedEx(std::forward(args)...); + }; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/blas_utils.h b/src/native/cuda/mars/blas_utils.h new file mode 100644 index 000000000..b6e4bc38c --- /dev/null +++ b/src/native/cuda/mars/blas_utils.h @@ -0,0 +1,30 @@ +#ifndef INFINI_OPS_MARS_BLAS_UTILS_H_ +#define INFINI_OPS_MARS_BLAS_UTILS_H_ + +// clang-format off +#include +// clang-format on + +#include "data_type.h" +#include "native/cuda/blas_utils.h" + +namespace infini::ops { + +template <> +struct BlasUtils { + static auto GetDataType(DataType dtype) { + if (dtype == DataType::kFloat16) return HPCC_R_16F; + if (dtype == DataType::kBFloat16) return HPCC_R_16BF; + return HPCC_R_32F; + } + + static auto GetComputeType(DataType dtype) { + if (dtype == DataType::kFloat16 || dtype == DataType::kBFloat16) + return HCBLAS_COMPUTE_32F; + return HCBLAS_COMPUTE_32F_FAST_TF32; + } +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/caster.cuh b/src/native/cuda/mars/caster.cuh new file mode 100644 index 000000000..6d51a151a --- /dev/null +++ b/src/native/cuda/mars/caster.cuh @@ -0,0 +1,78 @@ +#ifndef INFINI_OPS_MARS_CASTER__H_ +#define INFINI_OPS_MARS_CASTER__H_ + +#include "native/cuda/caster.cuh" +#include "native/cuda/mars/data_type_.h" + +namespace infini::ops { + +namespace detail { + +template <> +struct ToFloat { + __host__ __device__ float operator()(__half x) { return __half2float(x); } +}; + +template <> +struct ToFloat { + __host__ __device__ float operator()(__hpcc_bfloat16 x) { + return __bfloat162float(x); + } +}; + +template <> +struct FromFloat { + __host__ __device__ __half operator()(float f) { return __float2half(f); } +}; + +template <> +struct FromFloat { + __host__ __device__ __hpcc_bfloat16 operator()(float f) { + return __float2bfloat16(f); + } +}; + +template <> +struct HardwareCast { + inline static constexpr bool kSupported = true; + __host__ __device__ __hpcc_bfloat16 operator()(int x) { + return __int2bfloat16_rn(x); + } +}; + +template <> +struct HardwareCast { + inline static constexpr bool kSupported = true; + __host__ __device__ __half operator()(int x) { return __int2half_rn(x); } +}; + +template <> +struct HardwareCast { + inline static constexpr bool kSupported = true; + __host__ __device__ __hpcc_bfloat16 operator()(double x) { + return __double2bfloat16(x); + } +}; + +template <> +struct HardwareCast { + inline static constexpr bool kSupported = true; + __host__ __device__ __half operator()(double x) { return __double2half(x); } +}; + +template <> +struct HardwareCast { + inline static constexpr bool kSupported = true; + __host__ __device__ __half operator()(__hpcc_bfloat16 x) { + return __float2half_rn(__bfloat162float(x)); + } +}; + +} // namespace detail + +template <> +struct Caster : CudaCasterImpl {}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/data_type_.h b/src/native/cuda/mars/data_type_.h new file mode 100644 index 000000000..140e300c8 --- /dev/null +++ b/src/native/cuda/mars/data_type_.h @@ -0,0 +1,13 @@ +#ifndef INFINI_OPS_MARS_DATA_TYPE__H_ +#define INFINI_OPS_MARS_DATA_TYPE__H_ + +#include + +namespace infini::ops { + +using infini::rt::cuda_bfloat16; +using infini::rt::cuda_bfloat162; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/device_.h b/src/native/cuda/mars/device_.h new file mode 100644 index 000000000..eb7158d1a --- /dev/null +++ b/src/native/cuda/mars/device_.h @@ -0,0 +1,6 @@ +#ifndef INFINI_OPS_MARS_DEVICE__H_ +#define INFINI_OPS_MARS_DEVICE__H_ + +#include + +#endif diff --git a/src/native/cuda/mars/device_property.h b/src/native/cuda/mars/device_property.h new file mode 100644 index 000000000..0700a39a6 --- /dev/null +++ b/src/native/cuda/mars/device_property.h @@ -0,0 +1,11 @@ +#ifndef INFINI_OPS_MARS_DEVICE_PROPERTY_H_ +#define INFINI_OPS_MARS_DEVICE_PROPERTY_H_ + +namespace infini::ops { + +// TODO: Add HCR device properties query for Mars. +inline int QueryMaxThreadsPerBlock() { return 256; } + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/add/kernel.h b/src/native/cuda/mars/ops/add/kernel.h new file mode 100644 index 000000000..f94613cda --- /dev/null +++ b/src/native/cuda/mars/ops/add/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_ADD_KERNEL_H_ +#define INFINI_OPS_MARS_ADD_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/add/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaAdd> { + public: + using CudaAdd>::CudaAdd; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/add_rms_norm/kernel.h b/src/native/cuda/mars/ops/add_rms_norm/kernel.h new file mode 100644 index 000000000..5d7bf0b6a --- /dev/null +++ b/src/native/cuda/mars/ops/add_rms_norm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_ADD_RMS_NORM_KERNEL_H_ +#define INFINI_OPS_MARS_ADD_RMS_NORM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/add_rms_norm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaAddRmsNorm> { + public: + using CudaAddRmsNorm>::CudaAddRmsNorm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/causal_softmax/kernel.h b/src/native/cuda/mars/ops/causal_softmax/kernel.h new file mode 100644 index 000000000..df3b5579e --- /dev/null +++ b/src/native/cuda/mars/ops/causal_softmax/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_CAUSAL_SOFTMAX_KERNEL_H_ +#define INFINI_OPS_MARS_CAUSAL_SOFTMAX_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/causal_softmax/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaCausalSoftmax> { + public: + using CudaCausalSoftmax>::CudaCausalSoftmax; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/causal_softmax_infinilm/kernel.h b/src/native/cuda/mars/ops/causal_softmax_infinilm/kernel.h new file mode 100644 index 000000000..f97239ffb --- /dev/null +++ b/src/native/cuda/mars/ops/causal_softmax_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_CAUSAL_SOFTMAX_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_CAUSAL_SOFTMAX_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/causal_softmax_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaCausalSoftmaxInfinilm> { + public: + using CudaCausalSoftmaxInfinilm< + Runtime>::CudaCausalSoftmaxInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/conv1d/kernel.h b/src/native/cuda/mars/ops/conv1d/kernel.h new file mode 100644 index 000000000..f855b7fea --- /dev/null +++ b/src/native/cuda/mars/ops/conv1d/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_CONV1D_KERNEL_H_ +#define INFINI_OPS_MARS_CONV1D_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/convolution/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaConv, Conv1d> { + public: + using CudaConv, Conv1d>::CudaConv; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/conv2d/kernel.h b/src/native/cuda/mars/ops/conv2d/kernel.h new file mode 100644 index 000000000..e3b9ef5c0 --- /dev/null +++ b/src/native/cuda/mars/ops/conv2d/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_CONV2D_KERNEL_H_ +#define INFINI_OPS_MARS_CONV2D_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/convolution/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaConv, Conv2d> { + public: + using CudaConv, Conv2d>::CudaConv; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/conv3d/kernel.h b/src/native/cuda/mars/ops/conv3d/kernel.h new file mode 100644 index 000000000..02ff7b5cf --- /dev/null +++ b/src/native/cuda/mars/ops/conv3d/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_CONV3D_KERNEL_H_ +#define INFINI_OPS_MARS_CONV3D_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/convolution/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaConv, Conv3d> { + public: + using CudaConv, Conv3d>::CudaConv; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/conv_infinilm/kernel.h b/src/native/cuda/mars/ops/conv_infinilm/kernel.h new file mode 100644 index 000000000..f5010866f --- /dev/null +++ b/src/native/cuda/mars/ops/conv_infinilm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_CONV_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_CONV_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/conv_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaConvInfinilm> { + public: + using CudaConvInfinilm>::CudaConvInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/convolution/kernel.h b/src/native/cuda/mars/ops/convolution/kernel.h new file mode 100644 index 000000000..a69289ffc --- /dev/null +++ b/src/native/cuda/mars/ops/convolution/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_CONVOLUTION_KERNEL_H_ +#define INFINI_OPS_MARS_CONVOLUTION_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/convolution/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaConv, Convolution> { + public: + using CudaConv, Convolution>::CudaConv; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/copy/kernel.h b/src/native/cuda/mars/ops/copy/kernel.h new file mode 100644 index 000000000..36653d910 --- /dev/null +++ b/src/native/cuda/mars/ops/copy/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_COPY_KERNEL_H_ +#define INFINI_OPS_MARS_COPY_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/copy/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaCopy> { + public: + using CudaCopy>::CudaCopy; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/embedding/kernel.h b/src/native/cuda/mars/ops/embedding/kernel.h new file mode 100644 index 000000000..f98ac3da7 --- /dev/null +++ b/src/native/cuda/mars/ops/embedding/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_EMBEDDING_KERNEL_H_ +#define INFINI_OPS_MARS_EMBEDDING_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/embedding/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaEmbedding> { + public: + using CudaEmbedding>::CudaEmbedding; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/fill/kernel.h b/src/native/cuda/mars/ops/fill/kernel.h new file mode 100644 index 000000000..90d40c485 --- /dev/null +++ b/src/native/cuda/mars/ops/fill/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_FILL_KERNEL_H_ +#define INFINI_OPS_MARS_FILL_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/fill/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaFill> { + public: + using CudaFill>::CudaFill; +}; + +} // namespace infini::ops + +#endif // INFINI_OPS_MARS_FILL_KERNEL_H_ diff --git a/src/native/cuda/mars/ops/fused_add_rms_norm/kernel.h b/src/native/cuda/mars/ops/fused_add_rms_norm/kernel.h new file mode 100644 index 000000000..666d94474 --- /dev/null +++ b/src/native/cuda/mars/ops/fused_add_rms_norm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_FUSED_ADD_RMS_NORM_KERNEL_H_ +#define INFINI_OPS_MARS_FUSED_ADD_RMS_NORM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/fused_add_rms_norm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaFusedAddRmsNorm> { + public: + using CudaFusedAddRmsNorm>::CudaFusedAddRmsNorm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/gelu/kernel.h b/src/native/cuda/mars/ops/gelu/kernel.h new file mode 100644 index 000000000..67b88702d --- /dev/null +++ b/src/native/cuda/mars/ops/gelu/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_GELU_KERNEL_H_ +#define INFINI_OPS_MARS_GELU_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/gelu/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaGelu> { + public: + using CudaGelu>::CudaGelu; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/gelu_infinilm/kernel.h b/src/native/cuda/mars/ops/gelu_infinilm/kernel.h new file mode 100644 index 000000000..6272fb8b8 --- /dev/null +++ b/src/native/cuda/mars/ops/gelu_infinilm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_GELU_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_GELU_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/gelu_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaGeluInfinilm> { + public: + using CudaGeluInfinilm>::CudaGeluInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/gelutanh_infinilm/kernel.h b/src/native/cuda/mars/ops/gelutanh_infinilm/kernel.h new file mode 100644 index 000000000..9ae8bce99 --- /dev/null +++ b/src/native/cuda/mars/ops/gelutanh_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_GELUTANH_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_GELUTANH_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/gelutanh_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaGelutanhInfinilm> { + public: + using CudaGelutanhInfinilm< + Runtime>::CudaGelutanhInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/gemm/hcblas.h b/src/native/cuda/mars/ops/gemm/hcblas.h new file mode 100644 index 000000000..be00e99de --- /dev/null +++ b/src/native/cuda/mars/ops/gemm/hcblas.h @@ -0,0 +1,18 @@ +#ifndef INFINI_OPS_MARS_GEMM_HCBLAS_H_ +#define INFINI_OPS_MARS_GEMM_HCBLAS_H_ + +#include "native/cuda/mars/blas.h" +#include "native/cuda/ops/gemm/blas.h" + +namespace infini::ops { + +template <> +class Operator + : public BlasGemm> { + public: + using BlasGemm>::BlasGemm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/kv_caching_infinilm/kernel.h b/src/native/cuda/mars/ops/kv_caching_infinilm/kernel.h new file mode 100644 index 000000000..a46addbd9 --- /dev/null +++ b/src/native/cuda/mars/ops/kv_caching_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_KV_CACHING_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_KV_CACHING_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/kv_caching_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaKvCachingInfinilm> { + public: + using CudaKvCachingInfinilm< + Runtime>::CudaKvCachingInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/mul/kernel.h b/src/native/cuda/mars/ops/mul/kernel.h new file mode 100644 index 000000000..d4d2520c3 --- /dev/null +++ b/src/native/cuda/mars/ops/mul/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_MUL_KERNEL_H_ +#define INFINI_OPS_MARS_MUL_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/mul/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaMul> { + public: + using CudaMul>::CudaMul; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/paged_attention_infinilm/kernel.h b/src/native/cuda/mars/ops/paged_attention_infinilm/kernel.h new file mode 100644 index 000000000..4ac03add9 --- /dev/null +++ b/src/native/cuda/mars/ops/paged_attention_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_PAGED_ATTENTION_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_PAGED_ATTENTION_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/paged_attention_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaPagedAttentionInfinilm> { + public: + using CudaPagedAttentionInfinilm< + Runtime>::CudaPagedAttentionInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/paged_attention_prefill_infinilm/kernel.h b/src/native/cuda/mars/ops/paged_attention_prefill_infinilm/kernel.h new file mode 100644 index 000000000..3825d1002 --- /dev/null +++ b/src/native/cuda/mars/ops/paged_attention_prefill_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_PAGED_ATTENTION_PREFILL_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_PAGED_ATTENTION_PREFILL_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/paged_attention_prefill_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaPagedAttentionPrefillInfinilm> { + public: + using CudaPagedAttentionPrefillInfinilm< + Runtime>::CudaPagedAttentionPrefillInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/paged_caching_infinilm/kernel.h b/src/native/cuda/mars/ops/paged_caching_infinilm/kernel.h new file mode 100644 index 000000000..b95b3e252 --- /dev/null +++ b/src/native/cuda/mars/ops/paged_caching_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_PAGED_CACHING_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_PAGED_CACHING_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/paged_caching_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaPagedCachingInfinilm> { + public: + using CudaPagedCachingInfinilm< + Runtime>::CudaPagedCachingInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/random_sample_infinilm/kernel.h b/src/native/cuda/mars/ops/random_sample_infinilm/kernel.h new file mode 100644 index 000000000..2861368c1 --- /dev/null +++ b/src/native/cuda/mars/ops/random_sample_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_RANDOM_SAMPLE_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_RANDOM_SAMPLE_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/random_sample_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaRandomSampleInfinilm> { + public: + using CudaRandomSampleInfinilm< + Runtime>::CudaRandomSampleInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/rearrange_infinilm/kernel.h b/src/native/cuda/mars/ops/rearrange_infinilm/kernel.h new file mode 100644 index 000000000..85b9e235d --- /dev/null +++ b/src/native/cuda/mars/ops/rearrange_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_REARRANGE_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_REARRANGE_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/rearrange_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaRearrangeInfinilm> { + public: + using CudaRearrangeInfinilm< + Runtime>::CudaRearrangeInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/relu/kernel.h b/src/native/cuda/mars/ops/relu/kernel.h new file mode 100644 index 000000000..8fa6b8c76 --- /dev/null +++ b/src/native/cuda/mars/ops/relu/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_RELU_KERNEL_H_ +#define INFINI_OPS_MARS_RELU_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/relu/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaRelu> { + public: + using CudaRelu>::CudaRelu; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/relu_infinilm/kernel.h b/src/native/cuda/mars/ops/relu_infinilm/kernel.h new file mode 100644 index 000000000..e7810c07b --- /dev/null +++ b/src/native/cuda/mars/ops/relu_infinilm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_RELU_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_RELU_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/relu_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaReluInfinilm> { + public: + using CudaReluInfinilm>::CudaReluInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/reshape_and_cache_flash/kernel.h b/src/native/cuda/mars/ops/reshape_and_cache_flash/kernel.h new file mode 100644 index 000000000..c8a022a77 --- /dev/null +++ b/src/native/cuda/mars/ops/reshape_and_cache_flash/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_RESHAPE_AND_CACHE_FLASH_KERNEL_H_ +#define INFINI_OPS_MARS_RESHAPE_AND_CACHE_FLASH_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/reshape_and_cache_flash/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaReshapeAndCacheFlash> { + public: + using CudaReshapeAndCacheFlash< + Runtime>::CudaReshapeAndCacheFlash; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/rms_norm/kernel.h b/src/native/cuda/mars/ops/rms_norm/kernel.h new file mode 100644 index 000000000..3bc9f79e5 --- /dev/null +++ b/src/native/cuda/mars/ops/rms_norm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_RMS_NORM_KERNEL_H_ +#define INFINI_OPS_MARS_RMS_NORM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/rms_norm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaRmsNorm> { + public: + using CudaRmsNorm>::CudaRmsNorm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/rotary_embedding/kernel.h b/src/native/cuda/mars/ops/rotary_embedding/kernel.h new file mode 100644 index 000000000..fe5d73659 --- /dev/null +++ b/src/native/cuda/mars/ops/rotary_embedding/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_ROTARY_EMBEDDING_KERNEL_H_ +#define INFINI_OPS_MARS_ROTARY_EMBEDDING_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/rotary_embedding/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaRotaryEmbedding> { + public: + using CudaRotaryEmbedding>::CudaRotaryEmbedding; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/rotary_embedding_infinilm/kernel.h b/src/native/cuda/mars/ops/rotary_embedding_infinilm/kernel.h new file mode 100644 index 000000000..38629b6ae --- /dev/null +++ b/src/native/cuda/mars/ops/rotary_embedding_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_ROTARY_EMBEDDING_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_ROTARY_EMBEDDING_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/rotary_embedding_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaRotaryEmbeddingInfinilm> { + public: + using CudaRotaryEmbeddingInfinilm< + Runtime>::CudaRotaryEmbeddingInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/sigmoid/kernel.h b/src/native/cuda/mars/ops/sigmoid/kernel.h new file mode 100644 index 000000000..fbb6367eb --- /dev/null +++ b/src/native/cuda/mars/ops/sigmoid/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_SIGMOID_KERNEL_H_ +#define INFINI_OPS_MARS_SIGMOID_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/sigmoid/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSigmoid> { + public: + using CudaSigmoid>::CudaSigmoid; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/sigmoid_infinilm/kernel.h b/src/native/cuda/mars/ops/sigmoid_infinilm/kernel.h new file mode 100644 index 000000000..cde257122 --- /dev/null +++ b/src/native/cuda/mars/ops/sigmoid_infinilm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_SIGMOID_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_SIGMOID_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/sigmoid_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSigmoidInfinilm> { + public: + using CudaSigmoidInfinilm>::CudaSigmoidInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/silu/kernel.h b/src/native/cuda/mars/ops/silu/kernel.h new file mode 100644 index 000000000..2336ac20d --- /dev/null +++ b/src/native/cuda/mars/ops/silu/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_SILU_KERNEL_H_ +#define INFINI_OPS_MARS_SILU_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/silu/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSilu> { + public: + using CudaSilu>::CudaSilu; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/silu_and_mul/kernel.h b/src/native/cuda/mars/ops/silu_and_mul/kernel.h new file mode 100644 index 000000000..bed125e31 --- /dev/null +++ b/src/native/cuda/mars/ops/silu_and_mul/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_SILU_AND_MUL_KERNEL_H_ +#define INFINI_OPS_MARS_SILU_AND_MUL_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/silu_and_mul/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSiluAndMul> { + public: + using CudaSiluAndMul>::CudaSiluAndMul; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/silu_and_mul_infinilm/kernel.h b/src/native/cuda/mars/ops/silu_and_mul_infinilm/kernel.h new file mode 100644 index 000000000..45e136208 --- /dev/null +++ b/src/native/cuda/mars/ops/silu_and_mul_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_SILU_AND_MUL_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_SILU_AND_MUL_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/silu_and_mul_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSiluAndMulInfinilm> { + public: + using CudaSiluAndMulInfinilm< + Runtime>::CudaSiluAndMulInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/softmax/kernel.h b/src/native/cuda/mars/ops/softmax/kernel.h new file mode 100644 index 000000000..8b7ffa2a9 --- /dev/null +++ b/src/native/cuda/mars/ops/softmax/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_SOFTMAX_KERNEL_H_ +#define INFINI_OPS_MARS_SOFTMAX_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/softmax/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSoftmax> { + public: + using CudaSoftmax>::CudaSoftmax; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/softmax_infinilm/kernel.h b/src/native/cuda/mars/ops/softmax_infinilm/kernel.h new file mode 100644 index 000000000..f402e9d3c --- /dev/null +++ b/src/native/cuda/mars/ops/softmax_infinilm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_SOFTMAX_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_SOFTMAX_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/softmax_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSoftmaxInfinilm> { + public: + using CudaSoftmaxInfinilm>::CudaSoftmaxInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/swiglu/kernel.h b/src/native/cuda/mars/ops/swiglu/kernel.h new file mode 100644 index 000000000..4155454cc --- /dev/null +++ b/src/native/cuda/mars/ops/swiglu/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_SWIGLU_KERNEL_H_ +#define INFINI_OPS_MARS_SWIGLU_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/swiglu/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaSwiglu> { + public: + using CudaSwiglu>::CudaSwiglu; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/topk_softmax/kernel.h b/src/native/cuda/mars/ops/topk_softmax/kernel.h new file mode 100644 index 000000000..a4468c735 --- /dev/null +++ b/src/native/cuda/mars/ops/topk_softmax/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_TOPK_SOFTMAX_KERNEL_H_ +#define INFINI_OPS_MARS_TOPK_SOFTMAX_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/topk_softmax/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaTopkSoftmax> { + public: + using CudaTopkSoftmax>::CudaTopkSoftmax; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/topksoftmax_infinilm/kernel.h b/src/native/cuda/mars/ops/topksoftmax_infinilm/kernel.h new file mode 100644 index 000000000..687675804 --- /dev/null +++ b/src/native/cuda/mars/ops/topksoftmax_infinilm/kernel.h @@ -0,0 +1,22 @@ +#ifndef INFINI_OPS_MARS_TOPKSOFTMAX_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_TOPKSOFTMAX_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/topksoftmax_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaTopksoftmaxInfinilm> { + public: + using CudaTopksoftmaxInfinilm< + Runtime>::CudaTopksoftmaxInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/ops/zeros_infinilm/kernel.h b/src/native/cuda/mars/ops/zeros_infinilm/kernel.h new file mode 100644 index 000000000..9b677ac7e --- /dev/null +++ b/src/native/cuda/mars/ops/zeros_infinilm/kernel.h @@ -0,0 +1,21 @@ +#ifndef INFINI_OPS_MARS_ZEROS_INFINILM_KERNEL_H_ +#define INFINI_OPS_MARS_ZEROS_INFINILM_KERNEL_H_ + +#include + +#include "native/cuda/mars/caster.cuh" +#include "native/cuda/mars/runtime_.h" +#include "native/cuda/ops/zeros_infinilm/kernel.h" + +namespace infini::ops { + +template <> +class Operator + : public CudaZerosInfinilm> { + public: + using CudaZerosInfinilm>::CudaZerosInfinilm; +}; + +} // namespace infini::ops + +#endif diff --git a/src/native/cuda/mars/runtime_.h b/src/native/cuda/mars/runtime_.h new file mode 100644 index 000000000..ee0bbcdbd --- /dev/null +++ b/src/native/cuda/mars/runtime_.h @@ -0,0 +1,9 @@ +#ifndef INFINI_OPS_MARS_RUNTIME_H_ +#define INFINI_OPS_MARS_RUNTIME_H_ + +#include + +#include "native/cuda/mars/runtime_utils.h" +#include "runtime.h" + +#endif diff --git a/src/native/cuda/mars/runtime_utils.h b/src/native/cuda/mars/runtime_utils.h new file mode 100644 index 000000000..35c6c02bc --- /dev/null +++ b/src/native/cuda/mars/runtime_utils.h @@ -0,0 +1,15 @@ +#ifndef INFINI_OPS_MARS_RUNTIME_UTILS_H_ +#define INFINI_OPS_MARS_RUNTIME_UTILS_H_ + +#include "native/cuda/mars/device_property.h" +#include "native/cuda/runtime_utils.h" + +namespace infini::ops { + +template <> +struct RuntimeUtils + : CudaRuntimeUtils {}; + +} // namespace infini::ops + +#endif diff --git a/src/pybind11_utils.h b/src/pybind11_utils.h index b1bfe344f..42d2b219e 100644 --- a/src/pybind11_utils.h +++ b/src/pybind11_utils.h @@ -203,6 +203,7 @@ inline std::optional TryDeviceTypeFromString( {"cambricon", Device::Type::kCambricon}, {"ascend", Device::Type::kAscend}, {"metax", Device::Type::kMetax}, + {"mars", Device::Type::kMars}, {"moore", Device::Type::kMoore}, {"iluvatar", Device::Type::kIluvatar}, {"hygon", Device::Type::kHygon}, From 9805393dbd2f34e98edf2b98cb704270ef7f0066 Mon Sep 17 00:00:00 2001 From: Ceng23333 <441651826@qq.com> Date: Fri, 21 Aug 2026 17:19:42 +0800 Subject: [PATCH 2/2] fix(mars): enable smoke build with system g++ torch and Mars torch device Compile Mars torch sources with system g++ (USE_HPCC) like MetaX/Moore, map TorchDeviceName to cuda, instantiate Mars torch backends, and keep htcc .cu rewrites beside sources so bindings includes resolve correctly. Co-authored-by: Cursor --- .gitignore | 5 ++++ scripts/generate_torch_ops.py | 2 ++ scripts/htcc_wrapper.sh | 7 +++-- src/CMakeLists.txt | 15 ++++++++-- src/torch/device_.h | 5 ++++ src/torch/mars/c10.h | 28 +++++++++++++++++++ src/torch/ops/add/add.cc | 1 + src/torch/ops/gemm/gemm.cc | 1 + .../reshape_and_cache/reshape_and_cache.cc | 1 + .../ops/rotary_embedding/rotary_embedding.cc | 1 + .../scaled_dot_product_attention.cc | 1 + 11 files changed, 61 insertions(+), 6 deletions(-) create mode 100644 src/torch/mars/c10.h diff --git a/.gitignore b/.gitignore index c6501fb2d..6050e9104 100644 --- a/.gitignore +++ b/.gitignore @@ -262,3 +262,8 @@ __marimo__/ # Streamlit .streamlit/secrets.toml + +# htcc_wrapper sibling symlinks for device compilation +*.cc.cu +*.cpp.cu +*.cxx.cu diff --git a/scripts/generate_torch_ops.py b/scripts/generate_torch_ops.py index 482e0f5b3..5a6788cf6 100644 --- a/scripts/generate_torch_ops.py +++ b/scripts/generate_torch_ops.py @@ -75,6 +75,7 @@ "kMoore", "kIluvatar", "kHygon", + "kMars", ) _C10_DEVICE_TYPES = ( @@ -83,6 +84,7 @@ "kMetax", "kMoore", "kIluvatar", + "kMars", ) # YAML scalar-type tokens → C++ types. Reference types (e.g. `const Scalar&`) diff --git a/scripts/htcc_wrapper.sh b/scripts/htcc_wrapper.sh index 19d11743f..69ab3a431 100755 --- a/scripts/htcc_wrapper.sh +++ b/scripts/htcc_wrapper.sh @@ -1,6 +1,9 @@ #!/bin/bash # HPCC analog of mxcc_wrapper.sh. # htcc only enables device compilation for `.cu`. Rewrite compile inputs. +# Place the `.cu` symlink next to the source so same-directory `#include` +# resolution (e.g. bindings `mul.h` vs C-API `generated/include/mul.h`) still +# works — unlike rewriting into `/tmp`. ARGS=() skip_next=0 has_compile=0 @@ -30,14 +33,12 @@ done if [ "$has_compile" -eq 1 ]; then rewritten=() - mkdir -p /tmp/htcc-wrap for arg in "${ARGS[@]}"; do case "$arg" in *.cc|*.cpp|*.cxx) if [ -f "$arg" ]; then abs=$(readlink -f "$arg") - key=$(printf '%s' "$abs" | md5sum | awk '{print $1}') - cu="/tmp/htcc-wrap/${key}.cu" + cu="${abs}.cu" ln -sfn "$abs" "$cu" rewritten+=("$cu") else diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 8f33d1dee..c39ad6bef 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -889,9 +889,10 @@ if(TORCH_SOURCES) torch_compile=${INFINI_OPS_TORCH_COMPILE_JOBS}) endif() - if(WITH_METAX OR WITH_MOORE) - # Vendor compilers (`mxcc`/`mcc`) cannot compile vendor-forked `torch` - # headers. Compile `torch` sources with the system C++ compiler instead. + if(WITH_METAX OR WITH_MOORE OR WITH_MARS) + # Vendor compilers (`mxcc`/`mcc`/`htcc`) cannot compile vendor-forked + # `torch` headers. Compile `torch` sources with the system C++ compiler + # instead. find_program(SYSTEM_CXX NAMES g++ c++) if(NOT SYSTEM_CXX) @@ -910,6 +911,11 @@ if(TORCH_SOURCES) "-I${MACA_PATH}/include/mcr" "-I${MACA_PATH}/tools/cu-bridge/include") endif() + if(WITH_MARS) + list(APPEND _torch_vendor_include_flags + "-I${HPCC_PATH}/include" + "-I${HPCC_PATH}/tools/cu-bridge/include") + endif() if(WITH_MOORE) list(APPEND _torch_vendor_include_flags "-I${MUSA_ROOT}/include") execute_process( @@ -945,6 +951,9 @@ if(TORCH_SOURCES) if(WITH_METAX) list(APPEND _torch_extra_flags "-DUSE_MACA=1" "-DWITH_METAX=1") endif() + if(WITH_MARS) + list(APPEND _torch_extra_flags "-DUSE_HPCC=1" "-DWITH_MARS=1") + endif() if(WITH_MOORE) list(APPEND _torch_extra_flags "-DWITH_MOORE=1") endif() diff --git a/src/torch/device_.h b/src/torch/device_.h index dfd3cae6b..3cec42dc1 100644 --- a/src/torch/device_.h +++ b/src/torch/device_.h @@ -50,6 +50,11 @@ struct TorchDeviceName { static constexpr std::string_view kValue{"cuda"}; }; +template <> +struct TorchDeviceName { + static constexpr std::string_view kValue{"cuda"}; +}; + } // namespace infini::ops::detail #endif diff --git a/src/torch/mars/c10.h b/src/torch/mars/c10.h new file mode 100644 index 000000000..f0ead156d --- /dev/null +++ b/src/torch/mars/c10.h @@ -0,0 +1,28 @@ +#ifndef INFINI_OPS_TORCH_MARS_C10_H_ +#define INFINI_OPS_TORCH_MARS_C10_H_ + +#include +#include +#include + +#include "torch/c10.h" + +namespace infini::ops { + +template <> +struct C10 { + static constexpr Device::Type kDeviceType = Device::Type::kMars; + + using StreamGuard = c10::cuda::CUDAStreamGuard; + + static c10::cuda::CUDAStream GetStreamFromExternal(void* stream, + int device_index) { + return c10::cuda::getStreamFromExternal( + reinterpret_cast(stream), + static_cast(device_index)); + } +}; + +} // namespace infini::ops + +#endif // INFINI_OPS_TORCH_MARS_C10_H_ diff --git a/src/torch/ops/add/add.cc b/src/torch/ops/add/add.cc index cc36f5bc0..07134d5f2 100644 --- a/src/torch/ops/add/add.cc +++ b/src/torch/ops/add/add.cc @@ -40,5 +40,6 @@ template class Operator; template class Operator; template class Operator; template class Operator; +template class Operator; } // namespace infini::ops diff --git a/src/torch/ops/gemm/gemm.cc b/src/torch/ops/gemm/gemm.cc index 01a4a2a8f..499dbedcf 100644 --- a/src/torch/ops/gemm/gemm.cc +++ b/src/torch/ops/gemm/gemm.cc @@ -84,5 +84,6 @@ template class Operator; template class Operator; template class Operator; template class Operator; +template class Operator; } // namespace infini::ops diff --git a/src/torch/ops/reshape_and_cache/reshape_and_cache.cc b/src/torch/ops/reshape_and_cache/reshape_and_cache.cc index 99ca3c82a..60e76d7a7 100644 --- a/src/torch/ops/reshape_and_cache/reshape_and_cache.cc +++ b/src/torch/ops/reshape_and_cache/reshape_and_cache.cc @@ -85,5 +85,6 @@ template class Operator; template class Operator; template class Operator; template class Operator; +template class Operator; } // namespace infini::ops diff --git a/src/torch/ops/rotary_embedding/rotary_embedding.cc b/src/torch/ops/rotary_embedding/rotary_embedding.cc index 06658ea3a..a74920438 100644 --- a/src/torch/ops/rotary_embedding/rotary_embedding.cc +++ b/src/torch/ops/rotary_embedding/rotary_embedding.cc @@ -83,5 +83,6 @@ template class Operator; template class Operator; template class Operator; template class Operator; +template class Operator; } // namespace infini::ops diff --git a/src/torch/ops/scaled_dot_product_attention/scaled_dot_product_attention.cc b/src/torch/ops/scaled_dot_product_attention/scaled_dot_product_attention.cc index 81489cf8c..db5cf10f4 100644 --- a/src/torch/ops/scaled_dot_product_attention/scaled_dot_product_attention.cc +++ b/src/torch/ops/scaled_dot_product_attention/scaled_dot_product_attention.cc @@ -78,5 +78,6 @@ template class Operator; template class Operator; template class Operator; template class Operator; +template class Operator; } // namespace infini::ops