[Feature](mlu-ops): add mluApplyAdamW.

bangc_helper_dtype.h

@@ -0,0 +1,113 @@
+/*************************************************************************
+ * Copyright (C) [2024] by Cambricon, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the
+ * "Software"), to deal in the Software without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Software, and to
+ * permit persons to whom the Software is furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+ * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+ *************************************************************************/
+#pragma once
+
+/**
+ * Provides `BANG_WRAP_T(ptr_arg)` for .cc and `BANG_UNWRAP_T(ptr_arg)` for .mlu
+ * to bridge Eigen:: type and BANGC type
+ */
+
+#include <type_traits>
+
+struct bang_half_t;
+struct bang_bfloat16_t;
+
+namespace detail {
+/*
+ * `bang_wrap_data` and `bang_unwrap_data` could be the same thing,
+ *  but should be used in different scope
+ *
+ *  handle 'const DType', 'Dtype *',
+ *  could be implemented by SFINAE or just specialization
+ */
+template <typename DType, template <typename> class Impl,
+          typename RawType = DType>
+struct bang_trans_impl_ {
+  static_assert(std::is_same_v<RawType, DType>);
+  typedef DType type;
+};
+
+template <typename DType, template <typename> class Impl, typename RawType>
+struct bang_trans_impl_<DType*, Impl, RawType> {
+  typedef typename Impl<DType>::type* type;
+};
+
+template <typename DType, template <typename> class Impl, typename RawType>
+struct bang_trans_impl_<const DType, Impl, RawType> {
+  typedef const typename Impl<DType>::type type;
+};
+
+}  // namespace detail
+
+#define BANG_TRANS_TYPE_FROM_TO(TOKEN, From, To) \
+  template <>                                    \
+  struct TOKEN<From> {                           \
+    typedef To type;                             \
+  }
+
+/* For .cc/.cpp trans unknown type to wrapped type */
+#if !defined(__BANG__)
+
+namespace Eigen {
+struct half;
+struct bfloat16;
+}  // namespace Eigen
+
+template <typename DType>
+struct bang_wrap_data {
+  using type = typename detail::bang_trans_impl_<DType, bang_wrap_data>::type;
+};
+
+#define BANG_WRAP_TYPE_FROM_TO(From, To) \
+  BANG_TRANS_TYPE_FROM_TO(bang_wrap_data, From, To)
+
+BANG_WRAP_TYPE_FROM_TO(Eigen::half, bang_half_t);
+BANG_WRAP_TYPE_FROM_TO(Eigen::bfloat16, bang_bfloat16_t);
+
+template <typename T>
+using bang_wrap_data_t = typename bang_wrap_data<T>::type;
+
+#define BANG_WRAP_T(a) reinterpret_cast<bang_wrap_data_t<decltype(a)>>(a)
+
+#endif  // !defined(__BANG__)
+
+/* For .mlu trans intermediate type to mlu's underlying type */
+
+#if __BANG__
+template <typename DType>
+struct bang_unwrap_data {
+  using type = typename detail::bang_trans_impl_<DType, bang_unwrap_data>::type;
+};
+
+#define BANG_UNWRAP_TYPE_FROM_TO(From, To) \
+  BANG_TRANS_TYPE_FROM_TO(bang_unwrap_data, From, To)
+
+BANG_UNWRAP_TYPE_FROM_TO(bang_half_t, half);
+BANG_UNWRAP_TYPE_FROM_TO(bang_bfloat16_t, bfloat16_t);
+
+template <typename T>
+using bang_unwrap_data_t = typename bang_unwrap_data<T>::type;
+
+#define BANG_UNWRAP_T(a) reinterpret_cast<bang_unwrap_data_t<decltype(a)>>(a)
+
+#endif  // __BANG__

bangc_kernels.h

@@ -0,0 +1,95 @@
+/*************************************************************************
+ * Copyright (C) [2024] by Cambricon, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the
+ * "Software"), to deal in the Software without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Software, and to
+ * permit persons to whom the Software is furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
+ * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+ *************************************************************************/
+#ifndef BANGC_KERNELS_H_
+#define BANGC_KERNELS_H_
+
+#ifndef NAMESPACE_BANGC_KERNELS_GEGIN
+#define NAMESPACE_BANGC_KERNELS_GEGIN namespace bangc_kernels {
+#endif
+
+NAMESPACE_BANGC_KERNELS_GEGIN
+
+#ifndef BANGC_KERNELS_WIN_API
+#ifdef _WIN32
+#define BANGC_KERNELS_WIN_API __stdcall
+#else
+#define BANGC_KERNELS_WIN_API
+#endif
+#endif
+
+typedef enum {
+  BANGC_KERNELS_STATUS_SUCCESS =
+      0, /*!< The operation is successfully completed. */
+  BANGC_KERNELS_STATUS_ALLOC_FAILED = 1,
+  /*!< This error occurs when the resource allocation fails, which is usually
+     caused by failing to call cnMallocHost due to exceeded memory usage. Make
+     sure that the memory allocated previously is deallocated as much as
+     possible. */
+  BANGC_KERNELS_STATUS_BAD_PARAM = 2,
+  /*!< Invalid value or parameters are passed to the function, including data
+     type, layout, dimensions, etc. */
+  BANGC_KERNELS_STATUS_INTERNAL_ERROR = 3,
+  /*!< An error occurs inside of the function, which may indicate an internal
+     error or bug in the library. This error is usually caused by failing to
+     call cnrtMemcpyAsync. Check whether the memory passed to the function is
+     deallocated before the completion of the routine. */
+  BANGC_KERNELS_STATUS_ARCH_MISMATCH = 4,
+  /*!< Invalid MLU device which is not supported by current function. */
+  BANGC_KERNELS_STATUS_EXECUTION_FAILED = 5,
+  /*!< An error occurs when the function fails to be executed on MLU device due
+     to multiple reasons. You can check whether the hardware environment, driver
+     version and other prerequisite libraries are correctly installed. */
+  BANGC_KERNELS_STATUS_NOT_SUPPORTED = 6,
+  /*!< An error occurs when the requested functionality is not supported in this
+     version but would be supported in the future. */
+  BANGC_KERNELS_STATUS_NUMERICAL_OVERFLOW = 7,
+  /*!< A numerical overflow occurs when executing the function, which is usually
+     due to large scale or inappropriate range of value of input tensor. */
+} bangcKernelsStatus_t;
+
+template <typename T>
+bangcKernelsStatus_t BANGC_KERNELS_WIN_API
+mluApplyAdamW(const cnrtQueue_t queue,
+              const float lr,
+              const float beta1,
+              const float beta2,
+              const float bias1,
+              const float bias2,
+              const float epsilon,
+              const float weight_decay,
+              const float scale,
+              const bool use_nesterov,
+              const size_t size,
+              T *param_h,
+              T *grad,
+              void *param,
+              void *momentum,
+              void *velocity);
+
+#ifndef NAMESPACE_BANGC_KERNELS_END
+#define NAMESPACE_BANGC_KERNELS_END }
+#endif
+
+NAMESPACE_BANGC_KERNELS_END
+
+#endif  // BANGC_KERNELS_H_

independent_build.sh

@@ -8,6 +8,7 @@ MLUOP_TARGET_CPU_ARCH=`uname -m`
 GEN_SYMBOL_VIS_FILE_PY="./scripts/gen_symbol_visibility_map.py"
 MLUOP_SYMBOL_VIS_FILE="symbol_visibility.map"
 TARGET_SYMBOL_FILE="mlu_op.h"
+TARGET_SYMBOL_FILE_LITE="bangc_kernels.h"
 PACKAGE_EXTRACT_DIR="dep_libs_extract"
 
 PROG_NAME=$(basename $0)  # current script filename, DO NOT EDIT
@@ -421,7 +422,7 @@ if [ "$OS_RELEASE_ID" = "centos" -a "$OS_RELEASE_VERSION_ID" = "7" ]; then
   fi
 fi
 
-if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" < "5" ]]; then
+if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" -lt "5" ]]; then
   prog_log_note "we do not support g++<5, try to activate devtoolset-8 env"
   source /opt/rh/devtoolset-8/enable && prog_log_warn "devtoolset-8 activated" \
     || ( prog_log_warn "source devtoolset-8 failed, ignore this info if you have set env TOOLCHAIN_ROOT, TARGET_C_COMPILER, TARGET_CXX_COMPILER properly (see more details in README.md)" && sleep 4 ) # I hope user will see it
@@ -459,8 +460,8 @@ export PATH=${NEUWARE_HOME}/bin:$PATH
 export LD_LIBRARY_PATH=${NEUWARE_HOME}/lib64:$LD_LIBRARY_PATH
 
 prog_log_info "generate ${MLUOP_SYMBOL_VIS_FILE} file."
-prog_log_info "python3 ${GEN_SYMBOL_VIS_FILE_PY} ${BUILD_PATH}/${MLUOP_SYMBOL_VIS_FILE} ${TARGET_SYMBOL_FILE}"
-python3 ${GEN_SYMBOL_VIS_FILE_PY} ${BUILD_PATH}/${MLUOP_SYMBOL_VIS_FILE} ${TARGET_SYMBOL_FILE}
+prog_log_info "python3 ${GEN_SYMBOL_VIS_FILE_PY} ${BUILD_PATH}/${MLUOP_SYMBOL_VIS_FILE} ${TARGET_SYMBOL_FILE} ${TARGET_SYMBOL_FILE_LITE}"
+python3 ${GEN_SYMBOL_VIS_FILE_PY} ${BUILD_PATH}/${MLUOP_SYMBOL_VIS_FILE} ${TARGET_SYMBOL_FILE} ${TARGET_SYMBOL_FILE_LITE}
 
 pushd ${BUILD_PATH} > /dev/null
   prog_log_info "Rmove cmake cache ${PWD}"

kernels/adam_w/adam_w_union1.mlu

@@ -21,12 +21,14 @@
  * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
  *************************************************************************/
 
+#include "adam_w.h"
+
 #include <mlu.h>
-#include <cmath>
-#include <algorithm>
-#include "core/logging.h"
-#include "kernels/adam_w/adam_w.h"
-#include "kernels/utils/common.h"
+
+#include "bangc_helper_dtype.h"
+#include "bangc_kernels.h"
+
+#define MIN(a, b) ((a) < (b) ? (a) : (b))
 
 #define SIZE_NRAM_PER_REGION PAD_DOWN((MAX_NRAM_SIZE / 12), NFU_ALIGN_SIZE)
 #define HIGH_PRECISION_MODE 1
@@ -193,23 +195,22 @@ __mlu_global__ void unionApplyAdamW(T *param_h, T *grad, float *param,
                 ddr_param - 2 * num_x * param_flag, ddr_momentum - 2 * num_x,
                 ddr_velocity - 2 * num_x, nbuf_paramh, nbuf_param,
                 nbuf_momentum, nbuf_velocity,
-                std::min(num_x, (int)(num_task - (i - 2) * num_x)),
+                MIN(num_x, (int)(num_task - (i - 2) * num_x)),
                 (i - 2) % 2 * pong);
     }
     // load data
     if (i <= num_iter - 1) {
       loadData(nbuf_paramh, (T *)(nbuf_grad + pong / 2), nbuf_param,
                nbuf_momentum, nbuf_velocity, ddr_paramh, ddr_grad, ddr_param,
                ddr_momentum, ddr_velocity,
-               std::min(num_x, (int)(num_task - i * num_x)), i % 2 * pong);
+               MIN(num_x, (int)(num_task - i * num_x)), i % 2 * pong);
     }
     // compute
     if (i >= 1 && i <= num_iter) {
       computeAdamW(nbuf_paramh, (T *)(nbuf_grad + pong / 2), nbuf_param,
                    nbuf_grad, nbuf_momentum, nbuf_velocity, temp_1, temp_2, lr,
                    beta1, beta2, bias1, bias2, epsilon, weight_decay, scale,
-                   use_nesterov,
-                   std::min(num_x, (int)(num_task - (i - 1) * num_x)),
+                   use_nesterov, MIN(num_x, (int)(num_task - (i - 1) * num_x)),
                    (i - 1) % 2 * pong, param_flag);
     }
     ddr_paramh += num_x * paramh_flag;
@@ -228,16 +229,48 @@ mluOpStatus_t MLUOP_WIN_API KernelApplyAdamW(
     void *momentum, void *velocity, float lr, float beta1, float beta2,
     float bias1, float bias2, float epsilon, float weight_decay, float scale,
     bool use_nesterov, size_t size, mluOpDataType_t k_data_type) {
-  switch (k_data_type) {
-    default: {
-      LOG(ERROR) << "Not Implemented.";
-    }
-    case MLUOP_DTYPE_BFLOAT16: {
-      KERNEL_CHECK(unionApplyAdamW<bfloat16_t><<<k_dim, k_type, queue>>>(
-          (bfloat16_t *)param_h, (bfloat16_t *)grad, (float *)param,
-          (float *)momentum, (float *)velocity, lr, beta1, beta2, bias1, bias2,
-          epsilon, weight_decay, scale, use_nesterov, size));
-    }; break;
-  }
+  // launch kernel
+  unionApplyAdamW<bfloat16_t><<<k_dim, k_type, queue>>>(
+      (bfloat16_t *)param_h, (bfloat16_t *)grad, (float *)param,
+      (float *)momentum, (float *)velocity, lr, beta1, beta2, bias1, bias2,
+      epsilon, weight_decay, scale, use_nesterov, size);
   return MLUOP_STATUS_SUCCESS;
 }
+
+NAMESPACE_BANGC_KERNELS_GEGIN
+
+template <typename T>
+bangcKernelsStatus_t BANGC_KERNELS_WIN_API
+mluApplyAdamW(const cnrtQueue_t queue, const float lr, const float beta1,
+              const float beta2, const float bias1, const float bias2,
+              const float epsilon, const float weight_decay, const float scale,
+              const bool use_nesterov, size_t size, T *param_h, T *grad,
+              void *param, void *momentum, void *velocity) {
+  // set job type
+  int ordinal = -1;
+  int cluster_num;
+  int core_dim;
+  cnrtGetDevice(&ordinal);
+  cnrtDeviceGetAttribute(&core_dim, cnrtAttrMcorePerCluster, ordinal);
+  cnrtDeviceGetAttribute(&cluster_num, cnrtAttrMaxClusterPerUnionLimitTask,
+                         ordinal);
+  cnrtFunctionType_t k_type = cnrtFuncTypeUnion1;
+  cnrtDim3_t k_dim{.x = (uint32_t)core_dim, .y = (uint32_t)cluster_num, .z = 1};
+
+  // launch kernel
+  unionApplyAdamW<<<k_dim, k_type, queue>>>(
+      BANG_UNWRAP_T(param_h), BANG_UNWRAP_T(grad), (float *)param,
+      (float *)momentum, (float *)velocity, lr, beta1, beta2, bias1, bias2,
+      epsilon, weight_decay, scale, use_nesterov, size);
+  return BANGC_KERNELS_STATUS_SUCCESS;
+}
+
+#define IMPL_MLU_APPLY_ADAMW_KERNEL(DType)                                   \
+  template bangcKernelsStatus_t BANGC_KERNELS_WIN_API mluApplyAdamW(         \
+      const cnrtQueue_t, const float, const float, const float, const float, \
+      const float, const float, const float, const float, const bool,        \
+      const size_t, DType *, DType *, void *, void *, void *)
+
+IMPL_MLU_APPLY_ADAMW_KERNEL(bang_bfloat16_t);
+
+NAMESPACE_BANGC_KERNELS_END

mlu_op.h

@@ -20,8 +20,8 @@
  * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
  * SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
  *************************************************************************/
-#ifndef MLUOP_EXAMPLE_H_
-#define MLUOP_EXAMPLE_H_
+#ifndef MLUOP_H_
+#define MLUOP_H_
 
 /******************************************************************************
  * MLU-OPS: Cambricon Open Source operator library for Network
@@ -14526,4 +14526,4 @@ mluOpLgamma(mluOpHandle_t handle,
 }
 #endif
 
-#endif  // MLUOP_EXAMPLE_H_
+#endif  // MLUOP_H_

samples/mlu-ops/abs_sample/build.sh

@@ -9,7 +9,7 @@ if [ "$OS_RELEASE_ID" = "centos" -a "$OS_RELEASE_VERSION_ID" = "7" ]; then
   fi
 fi
 
-if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" < "5" ]]; then
+if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" -lt "5" ]]; then
   echo "we do not support g++<5, try to activate devtoolset-7 env"
   source /opt/rh/devtoolset-7/enable && echo "devtoolset-7 activated" \
     || ( echo "source devtoolset-7 failed, ignore this info if you have set env TOOLCHAIN_ROOT, TARGET_C_COMPILER, TARGET_CXX_COMPILER properly (see more details in README.md)" && sleep 4 ) # I hope user will see it

samples/mlu-ops/build.sh

@@ -9,7 +9,7 @@ if [ "$OS_RELEASE_ID" = "centos" -a "$OS_RELEASE_VERSION_ID" = "7" ]; then
   fi
 fi
 
-if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" < "5" ]]; then
+if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" -lt "5" ]]; then
   echo "we do not support g++<5, try to activate devtoolset-8 env"
   source /opt/rh/devtoolset-8/enable && echo "devtoolset-8 activated" \
     || ( echo "source devtoolset-8 failed, ignore this info if you have set env TOOLCHAIN_ROOT, TARGET_C_COMPILER, TARGET_CXX_COMPILER properly (see more details in README.md)" && sleep 4 ) # I hope user will see it

samples/mlu-ops/fault_sample/build.sh

@@ -9,7 +9,7 @@ if [ "$OS_RELEASE_ID" = "centos" -a "$OS_RELEASE_VERSION_ID" = "7" ]; then
   fi
 fi
 
-if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" < "5" ]]; then
+if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" -lt "5" ]]; then
   echo "we do not support g++<5, try to activate devtoolset-7 env"
   source /opt/rh/devtoolset-7/enable && echo "devtoolset-7 activated" \
     || ( echo "source devtoolset-7 failed, ignore this info if you have set env TOOLCHAIN_ROOT, TARGET_C_COMPILER, TARGET_CXX_COMPILER properly (see more details in README.md)" && sleep 4 ) # I hope user will see it

samples/mlu-ops/poly_nms_sample/build.sh

@@ -9,7 +9,7 @@ if [ "$OS_RELEASE_ID" = "centos" -a "$OS_RELEASE_VERSION_ID" = "7" ]; then
   fi
 fi
 
-if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" < "5" ]]; then
+if [[ "$(g++ --version | head -n1 | awk '{ print $3 }' | cut -d '.' -f1)" -lt "5" ]]; then
   echo "we do not support g++<5, try to activate devtoolset-7 env"
   source /opt/rh/devtoolset-7/enable && echo "devtoolset-7 activated" \
     || ( echo "source devtoolset-7 failed, ignore this info if you have set env TOOLCHAIN_ROOT, TARGET_C_COMPILER, TARGET_CXX_COMPILER properly (see more details in README.md)" && sleep 4 ) # I hope user will see it