Merge pull request #322 from YdrMaster/main

issue/291/style: 所有 maca 改为 metax

Author: PanZezhong1725

src/infiniccl/infiniccl.cc

@@ -2,7 +2,7 @@
 
 #include "./ascend/infiniccl_ascend.h"
 #include "./cuda/infiniccl_cuda.h"
-#include "./maca/infiniccl_maca.h"
+#include "./metax/infiniccl_metax.h"
 
 __C infiniStatus_t infinicclCommInitAll(
     infiniDevice_t device_type,
@@ -17,7 +17,7 @@ __C infiniStatus_t infinicclCommInitAll(
     switch (device_type) {
         COMM_INIT_ALL(INFINI_DEVICE_NVIDIA, cuda)
         COMM_INIT_ALL(INFINI_DEVICE_ASCEND, ascend)
-        COMM_INIT_ALL(INFINI_DEVICE_METAX, maca)
+        COMM_INIT_ALL(INFINI_DEVICE_METAX, metax)
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
     }
@@ -37,7 +37,7 @@ __C infiniStatus_t infinicclCommDestroy(infinicclComm_t comm) {
     switch (comm->device_type) {
         COMM_DESTROY(INFINI_DEVICE_NVIDIA, cuda)
         COMM_DESTROY(INFINI_DEVICE_ASCEND, ascend)
-        COMM_DESTROY(INFINI_DEVICE_METAX, maca)
+        COMM_DESTROY(INFINI_DEVICE_METAX, metax)
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;
@@ -65,7 +65,7 @@ __C infiniStatus_t infinicclAllReduce(
     switch (comm->device_type) {
         ALL_REDUCE(INFINI_DEVICE_NVIDIA, cuda)
         ALL_REDUCE(INFINI_DEVICE_ASCEND, ascend)
-        ALL_REDUCE(INFINI_DEVICE_METAX, maca)
+        ALL_REDUCE(INFINI_DEVICE_METAX, metax)
 
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;

src/infiniccl/maca/infiniccl_maca.h

@@ -1,4 +1,4 @@
-#include "infiniccl_maca.h"
+#include "infiniccl_metax.h"
 
 #include "../../utils.h"
 
@@ -51,7 +51,7 @@ inline hcclComm_t getHcclComm(infinicclComm_t comm) {
     return static_cast<hcclComm_t>(comm->comm);
 }
 
-namespace infiniccl::maca {
+namespace infiniccl::metax {
 
 infiniStatus_t commInitAll(
     infinicclComm_t *comms,
@@ -92,4 +92,4 @@ infiniStatus_t allReduce(
 
     return INFINI_STATUS_SUCCESS;
 }
-} // namespace infiniccl::maca
+} // namespace infiniccl::metax

src/infiniccl/maca/infiniccl_maca.cc renamed to src/infiniccl/metax/infiniccl_metax.cc

@@ -0,0 +1,12 @@
+#ifndef INFINICCL_METAX_H_
+#define INFINICCL_METAX_H_
+
+#include "../infiniccl_impl.h"
+
+#if defined(ENABLE_METAX_API) && defined(ENABLE_CCL)
+INFINICCL_DEVICE_API_IMPL(metax)
+#else
+INFINICCL_DEVICE_API_NOOP(metax)
+#endif
+
+#endif /* INFINICCL_METAX_H_ */

src/infiniccl/metax/infiniccl_metax.h

@@ -21,7 +21,7 @@
 #include "kunlun/kunlun_handle.h"
 #endif
 #ifdef ENABLE_METAX_API
-#include "maca/maca_handle.h"
+#include "metax/metax_handle.h"
 #endif
 
 __C infiniStatus_t infiniopCreateHandle(infiniopHandle_t *handle_ptr) {
@@ -57,7 +57,7 @@ __C infiniStatus_t infiniopCreateHandle(infiniopHandle_t *handle_ptr) {
         CREATE(INFINI_DEVICE_KUNLUN, kunlun);
 #endif
 #ifdef ENABLE_METAX_API
-        CREATE(INFINI_DEVICE_METAX, maca);
+        CREATE(INFINI_DEVICE_METAX, metax);
 #endif
 
     default:
@@ -94,7 +94,7 @@ __C infiniStatus_t infiniopDestroyHandle(infiniopHandle_t handle) {
         DELETE(INFINI_DEVICE_KUNLUN, kunlun);
 #endif
 #ifdef ENABLE_METAX_API
-        DELETE(INFINI_DEVICE_METAX, maca);
+        DELETE(INFINI_DEVICE_METAX, metax);
 #endif
     default:
         return INFINI_STATUS_DEVICE_TYPE_NOT_SUPPORTED;

src/infiniop/devices/handle.cc

@@ -1,14 +1,14 @@
 #include "../../../utils.h"
 #include "../pool.h"
-#include "maca_handle.h"
+#include "metax_handle.h"
 #include <hcblas/hcblas.h>
 #include <hcdnn/hcdnn.h>
 #include <memory>
 
 #define CHECK_MCBLAS(API) CHECK_INTERNAL(API, HCBLAS_STATUS_SUCCESS)
 #define CHECK_MCDNN(API) CHECK_INTERNAL(API, HCDNN_STATUS_SUCCESS)
 
-namespace device::maca {
+namespace device::metax {
 
 class Handle::Internal {
     Pool<hcblasHandle_t> mcblas_handles;
@@ -39,4 +39,4 @@ class Handle::Internal {
 
 hcdnnDataType_t getHcdnnDtype(infiniDtype_t dt);
 
-} // namespace device::maca
+} // namespace device::metax

src/infiniop/devices/maca/common_maca.h renamed to src/infiniop/devices/metax/metax_common.h

@@ -1,6 +1,6 @@
-#include "common_maca.h"
+#include "metax_common.h"
 
-namespace device::maca {
+namespace device::metax {
 Handle::Handle(infiniDevice_t device, int device_id)
     : InfiniopHandle{device, device_id},
       _internal(std::make_shared<Handle::Internal>(device_id)) {}
@@ -83,4 +83,4 @@ infiniStatus_t Handle::create(InfiniopHandle **handle_ptr, int device_id) {
     return INFINI_STATUS_SUCCESS;
 }
 
-} // namespace device::maca
+} // namespace device::metax

src/infiniop/devices/maca/maca_handle.cc renamed to src/infiniop/devices/metax/metax_handle.cc

@@ -1,10 +1,10 @@
-#ifndef __INFINIOP_MACA_HANDLE_H__
-#define __INFINIOP_MACA_HANDLE_H__
+#ifndef __INFINIOP_METAX_HANDLE_H__
+#define __INFINIOP_METAX_HANDLE_H__
 
 #include "../../handle.h"
 #include <memory>
 
-namespace device::maca {
+namespace device::metax {
 struct Handle : public InfiniopHandle {
     Handle(int device_id);
     class Internal;
@@ -20,6 +20,6 @@ struct Handle : public InfiniopHandle {
     std::shared_ptr<Internal> _internal;
 };
 
-} // namespace device::maca
+} // namespace device::metax
 
-#endif // __INFINIOP_MACA_HANDLE_H__
+#endif // __INFINIOP_METAX_HANDLE_H__

src/infiniop/devices/maca/maca_handle.h renamed to src/infiniop/devices/metax/metax_handle.h

@@ -1,16 +1,16 @@
-#define INFINIOP_MACA_KERNEL __global__ void
+#define INFINIOP_METAX_KERNEL __global__ void
 
-// Posible maximum number of threads per block for MACA architectures
+// Posible maximum number of threads per block for METAX architectures
 // Used for picking correct kernel launch configuration
-#define MACA_BLOCK_SIZE_1024 1024
-#define MACA_BLOCK_SIZE_512 512
+#define METAX_BLOCK_SIZE_1024 1024
+#define METAX_BLOCK_SIZE_512 512
 
-#define CHECK_MACA(API) CHECK_INTERNAL(API, hcSuccess)
+#define CHECK_METAX(API) CHECK_INTERNAL(API, hcSuccess)
 
 using cuda_bfloat16 = hpcc_bfloat16;
 using cuda_bfloat162 = hpcc_bfloat162;
 
-namespace device::maca {
+namespace device::metax {
 
 // return the memory offset of original tensor, given the flattened index of broadcasted tensor
 __forceinline__ __device__ __host__ size_t
@@ -41,7 +41,7 @@ indexToOffset(
     }
     return res;
 }
-} // namespace device::maca
+} // namespace device::metax
 
 __forceinline__ __device__ float
 exp_(const float val) {

src/infiniop/devices/maca/maca_kernel_common.h renamed to src/infiniop/devices/metax/metax_kernel_common.h

@@ -1,20 +1,20 @@
-#ifndef __INFINIOP_ELEMENTWISE_MACA_H__
-#define __INFINIOP_ELEMENTWISE_MACA_H__
+#ifndef __INFINIOP_ELEMENTWISE_METAX_H__
+#define __INFINIOP_ELEMENTWISE_METAX_H__
 
 #include "../../../utils.h"
-#include "../../devices/maca/common_maca.h"
-#include "../../devices/maca/maca_kernel_common.h"
-#include "elementwise_maca_api.h"
+#include "../../devices/metax/metax_common.h"
+#include "../../devices/metax/metax_kernel_common.h"
+#include "elementwise_metax_api.h"
 
-namespace op::elementwise::maca {
+namespace op::elementwise::metax {
 template <typename T>
 __device__ __forceinline__ const T *typedInputPtr(const void *ptr) {
     return reinterpret_cast<const T *>(ptr);
 }
 
 __device__ __forceinline__ size_t getOutputIndex(size_t idx, bool is_contiguous, size_t ndim,
                                                  const size_t *shape, const ptrdiff_t *strides) {
-    return is_contiguous ? idx : device::maca::indexToOffset(idx, ndim, shape, strides);
+    return is_contiguous ? idx : device::metax::indexToOffset(idx, ndim, shape, strides);
 }
 
 struct InputIndexer {
@@ -30,8 +30,8 @@ struct InputIndexer {
         return input_contiguous[input_id]
                  ? idx
                  : (input_broadcasted[input_id]
-                        ? device::maca::indexToReducedOffset(idx, ndim, output_strides, input_strides + input_id * ndim)
-                        : device::maca::indexToOffset(idx, ndim, input_shapes + input_id * ndim, input_strides + input_id * ndim));
+                        ? device::metax::indexToReducedOffset(idx, ndim, output_strides, input_strides + input_id * ndim)
+                        : device::metax::indexToOffset(idx, ndim, input_shapes + input_id * ndim, input_strides + input_id * ndim));
     }
 };
 
@@ -41,7 +41,7 @@ __device__ __forceinline__ void unpackInputsAndApply(F &&f, std::index_sequence<
 }
 
 template <size_t N, typename Op, typename Tdata, typename... Args>
-INFINIOP_MACA_KERNEL elementwiseKernel(
+INFINIOP_METAX_KERNEL elementwiseKernel(
     size_t output_size,
     size_t ndim,
     bool output_contiguous,
@@ -72,7 +72,7 @@ INFINIOP_MACA_KERNEL elementwiseKernel(
 }
 
 template <typename Op, typename Tout, typename... Tin>
-INFINIOP_MACA_KERNEL elementwiseKernel(
+INFINIOP_METAX_KERNEL elementwiseKernel(
     size_t output_size,
     size_t ndim,
     bool output_contiguous,
@@ -102,9 +102,9 @@ INFINIOP_MACA_KERNEL elementwiseKernel(
 }
 
 struct DeviceImpl::Opaque {
-    std::shared_ptr<device::maca::Handle::Internal> internal;
+    std::shared_ptr<device::metax::Handle::Internal> internal;
 
-    Opaque(const std::shared_ptr<device::maca::Handle::Internal> &internal)
+    Opaque(const std::shared_ptr<device::metax::Handle::Internal> &internal)
         : internal(internal) {}
 
     template <uint32_t BLOCK_SIZE, size_t N, typename Op, typename Tdata, typename... Args>
@@ -159,8 +159,8 @@ struct DeviceImpl::Opaque {
         const int8_t *d_meta_start = reinterpret_cast<int8_t *>(workspace) + input_arr_size;
 
         // copy the input pointer array and meta to device
-        CHECK_MACA(hcMemcpyAsync(workspace, h_inputs_arr, input_arr_size, hcMemcpyHostToDevice, stream));
-        CHECK_MACA(hcMemcpyAsync((void *)d_meta_start, info_meta_start, info.getMetaMemSize(), hcMemcpyHostToDevice, stream));
+        CHECK_METAX(hcMemcpyAsync(workspace, h_inputs_arr, input_arr_size, hcMemcpyHostToDevice, stream));
+        CHECK_METAX(hcMemcpyAsync((void *)d_meta_start, info_meta_start, info.getMetaMemSize(), hcMemcpyHostToDevice, stream));
 
         // offset/assign the pointers
         d_inputs_arr = reinterpret_cast<const void **>(workspace);
@@ -259,6 +259,6 @@ infiniStatus_t DeviceImpl::calculate(const op::elementwise::ElementwiseInfo &inf
         std::forward<Args>(args)...);
 }
 
-} // namespace op::elementwise::maca
+} // namespace op::elementwise::metax
 
 #endif