PaddlePaddle · guangzhixie · Jun 28, 2018 · Jun 28, 2018 · Jun 28, 2018 · Jun 28, 2018
diff --git a/cmake/compiler_options.cmake b/cmake/compiler_options.cmake
@@ -39,6 +39,9 @@ anakin_add_compile_option(-Wshadow)
 anakin_add_compile_option(-fpermissive)
 anakin_add_compile_option(-Wsign-promo)
 anakin_add_compile_option(-fdiagnostics-show-option)
+if(USE_BM_PLACE)
+    anakin_add_compile_option(-lbmlib-asic)
+endif()
 
 if(ENABLE_NOISY_WARNINGS)
 	anakin_add_compile_option(-Wcast-align)

diff --git a/cmake/find_modules.cmake b/cmake/find_modules.cmake
@@ -357,24 +357,22 @@ macro(anakin_find_openmp)
 endmacro()
 
 macro(anakin_find_bmlib)
-    find_path(BM_ROOT include/bmdnn/bmdnn_api.h ${CMAKE_SOURCE_DIR}/third-party/bm_lib/ $ENV{BM_ROOT}/) 
-    find_path(BM_ROOT_INCLUDE_DNN bmdnn_api.h ${BM_ROOT}/include/bmdnn) 
-    find_path(BM_ROOT_INCLUDE_RT bmruntime.h ${BM_ROOT}/include/bmruntime) 
-    find_path(BM_ROOT_INCLUDE_LIB bmlib_runtime.h ${BM_ROOT}/include/bmlib) 
-    if(BM_ROOT_INCLUDE_DNN AND BM_ROOT_INCLUDE_RT AND BM_ROOT_INCLUDE_LIB) 
-        set(BM_FOUND TRUE) 
-    endif() 
-    if(BM_FOUND) 
-        message(STATUS " Found bm_lib in ${BM_ROOT}  ${BM_ROOT_INCLUDE_DNN} ${BM_ROOT_INCLUDE_RT} ${BM_ROOT_INCLUDE_LIB}")
-        include_directories(${BM_ROOT_INCLUDE_DNN})
-        include_directories(${BM_ROOT_INCLUDE_RT}) 
-        include_directories(${BM_ROOT_INCLUDE_LIB}) 
-        set(BM_LIBRARIES "") 
-        list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/app/libbmdnn_device.so) 
-        list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/app/libbmlib_device.so)
-        list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/app/libbmrt.so) 
-        list(APPEND ANAKIN_LINKER_LIBS ${BM_LIBRARIES}) 
-    else() 
-        message(FATAL_ERROR "Could not found bm_lib") 
-    endif()
+	find_path(BM_ROOT include/bmlib/bmlib_runtime.h /usr/local/include/bm/ $ENV{BM_ROOT}/)
+	if(BM_ROOT)
+		set(BM_FOUND TRUE)
+	endif()
+	if(BM_FOUND)
+		message(STATUS " Found bm_lib in ${BM_ROOT}")
+		anakin_fetch_include_recursively(${BM_ROOT}/include)
+		set(BM_LIBRARIES "")
+		list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/device/bmlib.a)
+		list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/device/common-arm.a)
+		list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/device/common.a)
+		list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/device/fw-arm.a)
+		list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/device/fw-top.a)
+		list(APPEND BM_LIBRARIES ${BM_ROOT}/lib/device/libbmlib-asic.so)
+		list(APPEND ANAKIN_LINKER_LIBS ${BM_LIBRARIES})
+	else()
+		message(FATAL_ERROR "Could not found bm_lib")
+	endif()
 endmacro()
diff --git a/framework/core/net/worker.cpp b/framework/core/net/worker.cpp
@@ -106,15 +106,16 @@ Worker<Ttype, Ptype, RunType>::sync_prediction(std::vector<Tensor4d<typename tar
             d_tensor_in_p->reshape(ins[i].valid_shape());
             d_tensor_in_p->copy_from(ins[i]);
             d_tensor_in_p->set_seq_offset(ins[i].get_seq_offset());
-        } 
+        }
+#ifdef NVIDIA_GPU 
         Context<NV> ctx(0, 0, 0); 
         saber::SaberTimer<NV> my_time; 
         my_time.start(ctx);
 #ifdef ENABLE_OP_TIMER
         Context<NV> ctx(0, 0, 0); 
         saber::SaberTimer<NV> my_time;
         my_time.start(ctx);
-#endif
+#endif // ENABLE_OP_TIMER
         net.prediction(); 
 
         my_time.end(ctx); 
@@ -127,7 +128,8 @@ Worker<Ttype, Ptype, RunType>::sync_prediction(std::vector<Tensor4d<typename tar
             _thead_id_to_prediction_times_vec_in_ms[std::this_thread::get_id()].push_back(my_time.get_average_ms());
             LOG(ERROR) << " exec  << time: " << my_time.get_average_ms() << " ms ";
         }
-#endif
+#endif // ENABLE_OP_TIMER
+#endif // NVIDIA_GPU
         // get outputs of graph
         std::vector<Tensor4d<typename target_host<Ttype>::type>> ret;
         ret.resize(_outputs_in_order.size());
@@ -226,6 +228,16 @@ template class Worker<NV, Precision::FP16, OpRunType::SYNC>;
 template class Worker<NV, Precision::INT8, OpRunType::SYNC>;
 #endif
 
+#ifdef AMD_GPU
+template class Worker<AMD, Precision::FP32, OpRunType::ASYNC>;
+template class Worker<AMD, Precision::FP16, OpRunType::ASYNC>;
+template class Worker<AMD, Precision::INT8, OpRunType::ASYNC>;
+
+template class Worker<AMD, Precision::FP32, OpRunType::SYNC>;
+template class Worker<AMD, Precision::FP16, OpRunType::SYNC>;
+template class Worker<AMD, Precision::INT8, OpRunType::SYNC>;
+#endif
+
 #ifdef USE_X86_PLACE
 template class Worker<X86, Precision::FP32, OpRunType::ASYNC>;
 template class Worker<X86, Precision::FP16, OpRunType::ASYNC>;
@@ -256,4 +268,3 @@ template class Worker<ARM, Precision::INT8, OpRunType::SYNC>;
 #endif
 
 } /* namespace */
-
diff --git a/saber/CMakeLists.txt b/saber/CMakeLists.txt
@@ -92,6 +92,31 @@ if(USE_CUDA)
 				      ${WHOLE_ARCHIVE_END})	
 endif()
 
+if(USE_BM_PLACE)
+    set(BIN_NAME bmkernel_bin)
+    set(LINK_CONFIG link/bm1682_ddr.lds)
+    add_custom_command(OUTPUT bm_kernel_tmp
+        COMMAND arm-none-eabi-gcc ${ANAKIN_SABER}/funcs/impl/bm/device/bmkernel_base.c -mcpu=arm926ej-s -mfpu=vfp -fno-short-enums -std=gnu99 -O2 -Wall -Werror -ffunction-sections -fdata-sections -nostdlib -DENABLE_PRINT -I${BM_ROOT}/include/config -I${BM_ROOT}/include/common -I${BM_ROOT}/include/c_model -I${BM_ROOT}/include/firmware_core -I${BM_ROOT}/include/bmlib -c -o ${BIN_NAME}.o
+        COMMAND arm-none-eabi-gcc -T ${BM_ROOT}/${LINK_CONFIG} -mcpu=arm926ej-s -mfpu=vfp -fno-short-enums -Wl,--check-sections -Wl,--gc-sections -Wl,--unresolved-symbols=report-all -Wl,--no-enum-size-warning -o ${BIN_NAME}.elf -Wl,--start-group -lc -lm ${BIN_NAME}.o ${BM_ROOT}/lib/device/fw-top.a ${BM_ROOT}/lib/device/fw-arm.a -Wl,--end-group
+        COMMAND arm-none-eabi-objcopy -O binary -R *.slow* ${BIN_NAME}.elf ${BIN_NAME}_itcm.bin
+        COMMAND hexdump -v -e '1/4 \"%08x\\n\"' ${BIN_NAME}_itcm.bin > ${BIN_NAME}_itcm.hex.sim
+        COMMAND arm-none-eabi-objcopy -O binary -j *.slow* ${BIN_NAME}.elf ${BIN_NAME}_ddr.bin
+        COMMAND hexdump -v -e '1/4 \"%08x\\n\"' ${BIN_NAME}_ddr.bin > ${BIN_NAME}_ddr.hex.sim
+        COMMAND printf "%x" 0xAABBCCDD > ${BIN_NAME}.bin
+        COMMAND printf "%x" 0x0 >> ${BIN_NAME}.bin
+        COMMAND printf "%x" 0x0 >> ${BIN_NAME}.bin
+        COMMAND printf "%x" 0x0 >> ${BIN_NAME}.bin
+
+        COMMAND printf \"%x\" `wc -c < ${BIN_NAME}_itcm.hex.sim` >> ${BIN_NAME}.bin
+
+        COMMAND cat ${BIN_NAME}_itcm.hex.sim >> ${BIN_NAME}.bin
+        COMMAND cat ${BIN_NAME}_ddr.hex.sim >> ${BIN_NAME}.bin
+        COMMAND cp ${ANAKIN_ROOT}/build/saber/bmkernel_bin.bin /var/tmp/${BIN_NAME}.bin
+        COMMENT "BM Kernel compilation..."
+    )
+    add_custom_target(ANAKIN ALL DEPENDS bm_kernel_tmp)
+endif()
+
 # add saber library to static
 if(UNIX OR APPLE)
     if (USE_ARM_PLACE)

diff --git a/saber/core/common.h b/saber/core/common.h
@@ -179,11 +179,9 @@ const char* cudnn_get_errorstring(cudnnStatus_t status);
 #ifdef USE_BM_PLACE 
 
 #include "bmlib_runtime.h"
-#include "bmdnn_api.h"
-#include "bmdnn_ext_api.h"
 #include "bmlib_utils.h"
 
-#define BMDNN_CHECK(condition) \
+#define BM_CHECK(condition) \
   do { \
     bm_status_t error = condition; \
     CHECK_EQ(error, BM_SUCCESS) << " Failed with error code:" << error; \

diff --git a/saber/core/context.h b/saber/core/context.h
@@ -35,7 +35,7 @@ class Context final{
      * @param compute_stream_id
      */
     Context(int device_id = 0, int data_stream_id = 0, int compute_stream_id = 0){
-#ifdef USE_BM        
+#ifdef USE_BM_PLACE        
         if(std::is_same<TargetType, BM>::value){
             LOG(INFO) << "context init for BM";
             int dev_count = 0;
@@ -69,7 +69,7 @@ class Context final{
     }
 
     Context(const Context<TargetType>& ctx){
-#ifdef USE_BM
+#ifdef USE_BM_PLACE
         if(std::is_same<TargetType, BM>::value){
             LOG(INFO) << "context init for BM";
             _bm_handle = ctx._bm_handle;
@@ -98,7 +98,7 @@ class Context final{
         this->_act_ids = ctx._act_ids;
         this->_mode = ctx._mode;
 #endif
-#ifdef USE_BM
+#ifdef USE_BM_PLACE
         this->_bm_handle = ctx._bm_handle;
 #endif
         return *this;
@@ -109,7 +109,7 @@ class Context final{
         comp_eq = comp_eq && (_device_id == right._device_id);
         comp_eq = comp_eq && (_data_stream_id == right._data_stream_id);
         comp_eq = comp_eq && (_compute_stream_id == right._compute_stream_id);
-#ifdef USE_BM
+#ifdef USE_BM_PLACE
         comp_eq = comp_eq && (_bm_handle == right._bm_handle);
 #endif
         return comp_eq;
@@ -151,7 +151,7 @@ class Context final{
     //std::vector<int> get_act_ids();
 #endif
 
-#ifdef USE_BM
+#ifdef USE_BM_PLACE
     bm_handle_t get_handle() {
         return _bm_handle;
     }
@@ -170,7 +170,7 @@ class Context final{
     PowerMode _mode{SABER_POWER_HIGH};
     std::vector<int> _act_ids{0};
 #endif
-#ifdef USE_BM
+#ifdef USE_BM_PLACE
     bm_handle_t _bm_handle;
 #endif
 };

diff --git a/saber/core/data_traits.h b/saber/core/data_traits.h
@@ -20,7 +20,6 @@
 
 #ifdef USE_BM_PLACE 
 #include "bmlib_runtime.h"
-#include "bmdnn_api.h"
 #include "bmlib_utils.h"
 #endif
 

diff --git a/saber/core/impl/bm/bm_impl.cpp b/saber/core/impl/bm/bm_impl.cpp
@@ -45,19 +45,19 @@ typedef TargetWrapper<BM, __device_target> BM_API;
 
 // Init handle only once in the lifetime
 static bm_handle_t handle;
-static bm_status_t init_handle{bmdnn_init(&handle)};
+static bm_status_t init_handle{bmlib_kernel_init(&handle)};
 
 bm_handle_t BM_API::get_handle() {
     return handle;
 };
 
 void BM_API::get_device_count(int& count) {
-    BMDNN_CHECK(bm_dev_getcount(&count));
+    BM_CHECK(bm_dev_getcount(&count));
 }
 
 void BM_API::set_device(int id) {
     //(bm_handle_t &handle, bool bmkernel_used, int id){
-    //BMDNN_CHECK(bm_dev_request(&handle, 0, id));
+    //BM_CHECK(bm_dev_request(&handle, 0, id));
 }
 
 //TODO: Do we have this functionality?
@@ -69,39 +69,39 @@ void BM_API::mem_alloc(TPtr* ptr, size_t n) {
     /* bm_device_mem_t *mem = reinterpret_cast<struct bm_mem_desc *>(*ptr); */
     //    bm_device_mem_t *mem = new bm_device_mem_t();
     bm_device_mem_t mem;
-    BMDNN_CHECK(bm_malloc_device_byte(handle, &mem, n));
+    BM_CHECK(bm_malloc_device_byte(handle, &mem, n));
     *ptr = TPtr(mem);
 }
 
 void BM_API::mem_free(TPtr ptr) {
-    if ((ptr != BM_MEM_NULL)) {
+    if (bm_mem_get_type(ptr) == BM_MEM_TYPE_SYSTEM) {
         bm_free_device(handle, ptr);
         //        delete ptr;
     }
 }
 
 void BM_API::mem_set(TPtr ptr, int value, size_t n) {
     //(bm_handle_t handle, const int value, bm_device_mem_t mem){
-    BMDNN_CHECK(bm_memset_device(handle, value, ptr));
+    BM_CHECK(bm_memset_device(handle, value, ptr));
     //bm_device_mem_t* pmem = (struct bm_mem_desc *)(ptr);
-    //BMDNN_CHECK(bm_memset_device(handle, value, *pmem));
+    //BM_CHECK(bm_memset_device(handle, value, *pmem));
 }
 
 void BM_API::sync_memcpy(TPtr dst, size_t dst_offset, int dst_id, \
         const TPtr src, size_t src_offset, int src_id, \
         size_t count, __DtoD) {
     if(count==0)
         return;
-    //BMDNN_CHECK(bm_memcpy_d2d(handle, bm_mem_from_device(dst), dst_id, bm_mem_from_device(src), src_id, count));
-    BMDNN_CHECK(bm_memcpy_d2d(handle, dst, dst_offset, src, src_offset, count));
+    //BM_CHECK(bm_memcpy_d2d(handle, bm_mem_from_device(dst), dst_id, bm_mem_from_device(src), src_id, count));
+    BM_CHECK(bm_memcpy_d2d(handle, dst, dst_offset, src, src_offset, count));
 };
 
 void BM_API::sync_memcpy(TPtr dst, size_t dst_offset, int dst_id, \
         const void* src, size_t src_offset, int src_id, \
         size_t count, __HtoD) {
     if(count==0)
         return;
-    BMDNN_CHECK(bm_memcpy_s2d(handle, dst+dst_offset, bm_mem_from_system(const_cast<void*>(src)+src_offset)));
+    BM_CHECK(bm_memcpy_s2d(handle, dst+dst_offset, bm_mem_from_system(const_cast<void*>(src)+src_offset)));
 
 #ifdef DEBUG
 
@@ -118,7 +118,7 @@ void BM_API::sync_memcpy(void* dst, size_t dst_offset, int dst_id, \
     if(count==0)
         return;
 //    LOG(INFO)<<"host ptr = "<<(dst)<<",dst_offset = "<<dst_offset<<", dev ptr = "<<(src)<<",dev offset = "<<src_offset;
-    BMDNN_CHECK(bm_memcpy_d2s(handle, bm_mem_from_system(dst+dst_offset), src+src_offset));
+    BM_CHECK(bm_memcpy_d2s(handle, bm_mem_from_system(dst+dst_offset), src+src_offset));
 
 #ifdef DEBUG
 

diff --git a/saber/core/impl/bm/tensor_op_bm.cpp b/saber/core/impl/bm/tensor_op_bm.cpp
@@ -6,29 +6,33 @@ namespace saber {
 template<>
 void fill_tensor_const<BM>(Tensor<BM>& tensor, float value,
                            typename Tensor<BM>::API::stream_t stream = NULL) {
-    Tensor<X86> temp_tensor(tensor.valid_shape(),tensor.get_dtype());
+    Tensor<X86> temp_tensor(tensor.shape(), tensor.get_dtype());
+    temp_tensor.set_shape(tensor.valid_shape());
     fill_tensor_const(temp_tensor, value);
     tensor.copy_from(temp_tensor);
 }
 template<>
 void fill_tensor_rand<BM>(Tensor<BM>& tensor, typename Tensor<BM>::API::stream_t stream = NULL) {
-    Tensor<X86> temp_tensor(tensor.valid_shape(),tensor.get_dtype());
+    Tensor<X86> temp_tensor(tensor.shape(), tensor.get_dtype());
+    temp_tensor.set_shape(tensor.valid_shape());
     fill_tensor_rand(temp_tensor);
     tensor.copy_from(temp_tensor);
 }
 
 template<>
 void fill_tensor_rand<BM>(Tensor<BM>& tensor, float vstart, float vend,
                           typename Tensor<BM>::API::stream_t stream = NULL) {
-    Tensor<X86> temp_tensor(tensor.valid_shape(),tensor.get_dtype());
+    Tensor<X86> temp_tensor(tensor.shape(), tensor.get_dtype());
+    temp_tensor.set_shape(tensor.valid_shape());
     fill_tensor_rand(temp_tensor, vstart, vend);
     tensor.copy_from(temp_tensor);
 }
 
 template<>
 void print_tensor<BM>(Tensor<BM>& tensor, typename Tensor<BM>::API::stream_t stream = NULL) {
-    LOG(INFO) << "device tensor data";
-    Tensor<X86> temp_tensor(tensor.valid_shape(),tensor.get_dtype());
+    LOG(INFO) << "BM device tensor data:";
+    Tensor<X86> temp_tensor(tensor.shape(), tensor.get_dtype());
+    temp_tensor.set_shape(tensor.valid_shape());
     temp_tensor.copy_from(tensor);
     print_tensor(temp_tensor);
 }
@@ -41,15 +45,17 @@ void print_tensor_valid<BM>(Tensor<BM>& tensor, typename Tensor<BM>::API::stream
 
 template<>
 double tensor_mean_value<BM>(Tensor<BM>& tensor, typename Tensor<BM>::API::stream_t stream = NULL) {
-    Tensor<X86> temp_tensor(tensor.valid_shape(),tensor.get_dtype());
+    Tensor<X86> temp_tensor(tensor.shape(), tensor.get_dtype());
+    temp_tensor.set_shape(tensor.valid_shape());
     temp_tensor.copy_from(tensor);
     return tensor_mean_value(temp_tensor);
 }
 
 template<>
 double tensor_mean_value_valid<BM>(Tensor<BM>& tensor,
                                    typename Tensor<BM>::API::stream_t stream = NULL) {
-    Tensor<X86> temp_tensor(tensor.valid_shape(),tensor.get_dtype());
+    Tensor<X86> temp_tensor(tensor.shape(), tensor.get_dtype());
+    temp_tensor.set_shape(tensor.valid_shape());
     temp_tensor.copy_from(tensor);
     return tensor_mean_value(temp_tensor);
 }

diff --git a/saber/core/target_wrapper.h b/saber/core/target_wrapper.h
@@ -450,7 +450,7 @@ struct TargetWrapper<BM, __device_target> {
     static bm_handle_t get_handle();
 
 };
-#endif //USE_BM
+#endif //USE_BM_PLACE
 
 #ifdef AMD_GPU