diff --git a/README.md b/README.md new file mode 100644 index 0000000..1e2de76 --- /dev/null +++ b/README.md @@ -0,0 +1,503 @@ +# 芯合原生平台 + +## 1 原生编译器部署 + +### 1.1 基础需求 + +* 硬件：X86通用服务器一台，CPU > 4、 内存 > 8G +* 操作系统：Linux （推荐使用ubuntu20.04 或 ubuntu22.04） + +### 1.2 宿主机配置 + +* docker安装 （推荐使用24以上版本） + + 参考： + +### 1.3 镜像获取 + +移动云盘：链接: +提取码\:gOr2 + +### 1.4 编译器容器部署 + +通过`docker`命令加载本地下载好的编译器镜像。 + +```shell +sudo docker load --input xh-cmcc-compiler-v1.0.tar +``` + +其中，`xh-cmcc-compiler-v1.0.tar`为芯合编译器镜像文件的名称。 + +通过`docker`命令创建新的容器实例。 + +```shell +sudo docker run --privileged=true -itd -v /etc/timezone:/etc/timezone:ro -v /etc/localtime:/etc/localtime:ro --name cmcc-compiler --hostname cmcc-compiler xh-cmcc-compiler:v1.0 bash +``` + +除了下述参数以外，均为`docker`自身运行相关的参数，如需了解更多，可参考`docker`的官方文档。 + +1. `--privileged=true`：为提权命令，固定参数。 +2. `-v /etc/timezone:/etc/timezone:ro`和`-v /etc/localtime:/etc/localtime:ro`为挂载宿主机的时间信息，固定参数。 +3. `--name`：指定容器的名字，此参数用户可根据需要变更。 +4. `--hostname`：指定主机名，此参数用户可根据需要变更。 +5. `xh-cmcc-compiler:v1.0`：指定编译器镜像的名称和版本，固定参数。 +6. `bash`：指定容器内需要运行的进程命令，固定参数。 + +通过`docker`命令进入容器进行操作。 + +```shell +sudo docker exec ${CONTAINER_ID} bash +``` + +### 1.5 如何使用编译器 + +编译器编译命令以及命令参数说明： + +* 原生编译器采用`clang++`编译器对源代码进行编译 + +* 常用参数说明： + + * `-fsycl-device-only`：使用`Clang++`编译器编译`SYCL`程序时，只针对设备（`GPU`）进行优化，该参数为必选参数。 + * `-tensor_compute`：`tensor_compute`为添加编译参数，指定后编译器将执行`Tensor Compute`编译过程，该参数为必选参数；需要注意的是，指定该参数后，需通过`-I`参数指定对应的`mlir`文件所在目录，否则编译过程会报错。 + * `-I`: 用于包含编译过程中所需头文件或指定编译输入的路径，采用该参数指定编译输入的路径时为必选参数，采用该参数包含头文件时根据需要包含的头文件的具体情况而定。 + * `-L`: 链接编译过程中所需的第三方库，例如`opencv`、`ffmpeg`等，该参数的使用根据需要链接的库的情况而定。 + * `-o`: 指定编译好的二进制文件的保存目录和文件名，该参数为必选参数。 + + 注意：其他编译选项参考`LLVM`编译选项的使用。 + +* 编译命令举例 + + ```shell + # 该命令采用clang++编译器针对设备进行优化，启用了tensor_compute功能，将sample.cpp编译为sample二进制可执行文件，并通过-I指定了使用模型的mlir文件 + clang++ -fsycl-device-only -tensor_compute -I/mlir/path/ sample.cpp -o sample + ``` + +## 2 原生运行时部署 + +### 2.1 硬件支持 + +芯合原生运行时当前支持以下3种加速硬件： + +* 英伟达：tasla V100 和 tasla T4 +* 华为： ansend 910 +* 瀚博： VA1A + +### 2.2 瀚博环境运行时部署（VA1A） + +#### 2.2.1 基础需求 + +* 硬件：具备瀚博VA1A加速卡的X86架构服务器一台， CPU > 8、 内存 > 16G (推荐) +* 操作系统：Linux (推荐使用ubuntu20.04) +* `DKMS`版本`1.95`及以上 + +#### 2.2.2 宿主机配置 + +**（1）安装环境依赖包：** + + yum install cairo-devel.x86_64 + yum install libdrm-devel.x86_64 + yum install kmod-devel.x86_64 + yum install libdwarf-devel.x86_64 + yum install libpciaccess-devel.x86_64 + yum install pixman-devel.x86_64 + yum install procps-ng-devel.x86_64 + yum install pkgconf.x86_64 + yum install openssl.x86_64 + yum install cmake + yum install automake.noarch + yum install autoconf.noarch + yum install openssl-devel.x86_64 + yum install libtool.x86_64 + yum install libX11-devel.x86_64 + yum install libXext-devel.x86_64 + yum install libXfixes-devel.x86_64 + +**（2）安装设备驱动：** + + VA1A驱动版本：d3\_0\_v2\_4\_a1\_5\_rc3 + +​ 驱动获取地址： + +​ 驱动安装：`sudo rpm -i vastai-pci-xxx.rpm ` 其中，`xxx`为版本相关信息 + +​ 检查驱动安装状态：`lsmod | grep vastai_pci ` 回显 `vastai_pci` 说明安装成功 + +**（3）安装SDK：** + + 加速库版本：AI Compiler v1.5.0 + +​ SDK获取地址： + +​ SDK安装： `sudo rpm -i VaststreamSdk-xxx.rpm` 其中，`xxx`为版本相关信息 + +​ 检查安装状态： `rpm -qa | grep -i Vaststream ` 回显`VaststreamSdk-xxx`说明安装成功 + +**（4）安装预编译库：** + +​ 预编译库获取地址： + +​ 预编译库安装： `sudo ./vastai_video_install_xxx.bin /opt/vastai/vaststream` + +​ 其中，`/opt/vastai/vaststream`是预编译库安装目录 + +​ 检查预编译库安装状态： + + /opt/vastai/vaststream + ├── bin + ├── env + ├── env.sh + ├── include + ├── lib + ├── patch + ├── python + └── samples + +#### 2.2.3 镜像获取 + +移动云盘链接: +提取码\:cREu + +#### 2.2.4 运行时容器部署 + +通过`docker`命令加载本地下载好的编译器镜像。 + +```shell +sudo docker load --input xh-vastai-runtime-v1.0.tar +``` + +其中，`xh-vastai-runtime-v1.0.tar`为瀚博运行时镜像文件的名称。 + +使用`docker`命令创建新的容器实例。 + +```shell +sudo docker run --privileged=true -itd -v /etc/timezone:/etc/timezone:ro -v /etc/localtime:/etc/localtime:ro --name cmcc-vastai-runtime --hostname cmcc-vastai-runtime xh-vastai-runtime:v1.0 bash +``` + +除了下述参数以外，均为`docker`自身运行相关的参数，如需了解更多，可参考`docker`的官方文档。 + +1. `--privileged=true`：为提权命令，固定参数。 +2. `-v /etc/timezone:/etc/timezone:ro`和`-v /etc/localtime:/etc/localtime:ro`为挂载宿主机的时间信息，固定参数。 +3. `--name`：指定容器的名字，此参数用户可根据需要变更。 +4. `--hostname`：指定主机名，此参数用户可根据需要变更。 +5. `xh-vastai-runtime:v1.0`：指定瀚博运行时镜像的名称和版本，固定参数。 +6. `bash`：指定容器内需要运行的进程命令，固定参数。 + +通过`docker`命令进入容器进行操作。 + +```shell +sudo docker exec ${CONTAINER_ID} bash +``` + +### 2.3 英伟达环境运行时部署（`V100`、`T4`） + +#### 2.3.1 基础需求 + +* 硬件：具备tesla V100或T4 GPU的X86架构服务器一台， CPU > 8、 内存 > 16G (推荐) +* 操作系统：Linux (推荐使用ubuntu20.04) + +#### 2.3.2 宿主机配置 + +* GPU驱动安装： + + 推荐版本：525 +* cuda 安装： + + 推荐版本：11.2 +* cudnn 安装： + + 推荐版本：8.1 +* docker安装 (推荐使用24以上版本) 参考： +* nvidia 容器工具安装及配置： + +#### 2.3.3 镜像获取 + +移动云盘： +v100镜像 链接: +提取码\:DCGu +T4镜像链接: +提取码\:G8gK + +#### 2.3.4 运行时容器部署 + +英伟达运行时镜像区分`V100`以及`T4`版本，需要根据实际的硬件环境进行选择安装。 + +##### 2.3.4.1 `V100`环境部署 + +通过`docker`命令加载本地下载好的编译器镜像。 + +```shell +sudo docker load --input xh-nvidia-runtime-v100-v1.0.tar +``` + +其中，`xh-nvidia-runtime-v100-v1.0.tar`为英伟达`V100`运行时镜像文件的名称。 + +使用`docker`命令创建新的容器实例。 + +```shell +sudo docker run -itd --name cmcc-nvidia-runtime --runtime=nvidia --ipc host --pid host --privileged=true --hostname cmcc-nvidia-runtime -v /etc/timezone:/etc/timezone:ro -v /etc/localtime:/etc/localtime:ro xh-nvidia-runtime-v100:v1.0 bash +``` + +除了下述参数以外，均为`docker`自身运行相关的参数，如需了解更多，可参考`docker`的官方文档。 + +1. `--privileged=true`：为提权命令，固定参数。 +2. `--runtime=nvidia`：用于指定容器运行时为英伟达环境，固定参数。 +3. `--ipc host`和`--pid host`：为容器和宿主机之间共享内存及进程，固定参数。 +4. `-v /etc/timezone:/etc/timezone:ro`和`-v /etc/localtime:/etc/localtime:ro`为挂载宿主机的时间信息，固定参数。 +5. `--name`：指定容器的名字，此参数用户可根据需要变更。 +6. `--hostname`：指定主机名，此参数用户可根据需要变更。 +7. `xh-nvidia-runtime-v100:v1.0`：指定瀚博运行时镜像的名称和版本，固定参数。 +8. `bash`：指定容器内需要运行的进程命令，固定参数。 + +通过`docker`命令进入容器进行操作。 + +```shell +sudo docker exec ${CONTAINER_ID} bash +``` + +##### 2.3.4.2 `T4`环境部署 + +通过`docker`命令加载本地下载好的编译器镜像。 + +```shell +sudo docker load --input xh-vastai-runtime-t4-v1.0.tar +``` + +其中，`xh-nvidia-runtime-t4-v1.0.tar`为英伟达`T4`运行时镜像文件的名称。 + +使用`docker`命令创建新的容器实例。 + +```shell +sudo docker run -itd --name cmcc-nvidia-runtime --runtime=nvidia --ipc host --pid host --privileged=true --hostname cmcc-nvidia-runtime -v /etc/timezone:/etc/timezone:ro -v /etc/localtime:/etc/localtime:ro xh-nvidia-runtime-t4:v1.0 bash +``` + +除了下述参数以外，均为`docker`自身运行相关的参数，如需了解更多，可参考`docker`的官方文档。 + +1. `--privileged=true`：为提权命令，固定参数。 +2. `--runtime=nvidia`：用于指定容器运行时为英伟达环境，固定参数。 +3. `--ipc host`和`--pid host`：为容器和宿主机之间共享内存及进程，固定参数。 +4. `-v /etc/timezone:/etc/timezone:ro`和`-v /etc/localtime:/etc/localtime:ro`为挂载宿主机的时间信息，固定参数。 +5. `--name`：指定容器的名字，此参数用户可根据需要变更。 +6. `--hostname`：指定主机名，此参数用户可根据需要变更。 +7. `xh-nvidia-runtime-t4:v1.0`：指定瀚博运行时镜像的名称和版本，固定参数。 +8. `bash`：指定容器内需要运行的进程命令，固定参数。 + +通过`docker`命令进入容器进行操作。 + +```shell +sudo docker exec ${CONTAINER_ID} bash +``` + +### 2.4 华为环境运行时部署（`ascend910`） + +*详细部署说明和问题解决请参考：* + +** + +#### 2.4.1 基础需求 + +**（1）硬件款型说明：** + +​ `Atlas 800`训练服务器（型号：`9010`） + +**（2）支持操作系统列表：** + +​ `Ubuntu 18.04.1`、`Ubuntu 18.04.5`、`Ubuntu 20.04`、`CentOS 7.6`、`CentOS 8.2`、 + +​ `openEuler 20.03 LTS`、`openEuler 22.03 LTS`、`Kylin V10 SP1`、`BC_Linux 7.6`、 + +​ `Debian 9.9`、`Debian 10.0` + +#### 2.4.2 宿主机配置 + +**（1）安装场景说明：** + +**（2）安装部件说明：** + +① 昇腾`NPU`固件：固件包含昇腾`AI`处理器自带的`OS` 、电源器件和功耗管理器件控制软件，分别用于后续加载到`AI`处理器的模型计算、芯片启动控制和功耗控制。 + +② 昇腾`NPU`驱动：部署在昇腾服务器，用于管理查询昇腾AI处理器，同时为上层`CANN`软件提供芯片控制、资源分配等接口。 + +③ `CANN`：部署在昇腾服务器，包含`Runtime`、算子库、图引擎、媒体数据处理等组件，通过`AscendCL`（`Ascend Computing Language`，昇腾计算语言）对外提供`Device`管理、`Context`管理、`Stream`管理、内存管理、模型加载与执行、算子加载与执行、媒体数据处理等`API`，帮助开发者实现在昇腾软硬件平台上开发和运行`AI`业务。 + +④` Ascend Docker`：`Ascend Docker`（容器引擎插件）本质上是基于`OCI`标准（开放容器倡议标准）实现的`Docker Runtime`（容器运行环境），不修改`Docker`引擎，对`Docker`以插件方式提供`Ascend NPU`适配功能，使用户`AI`作业能够以`Docker`容器的方式平滑运行在昇腾设备上。 + +**（3）安装硬件设备：** + +硬件设备安装请参考： + +**（4）安装OS：** + +各版本OS安装请参考： + +**（5） 下载软件包：** + +软件包下载地址： + + + +**（6）安装NPU驱动和固件：** + +安装NPU驱动： + +`./Ascend-hdk-910-npu-driver_23.0.rc1_linux-x86-64.run --full --install-for-all` + +验证驱动功能可用性：`npu-smi info` + +安装NPU固件：`./Ascend-hdk-910-npu-firmware_6.3.0.1.241.run --full` + +重启系统：`reboot` + +详细部署说明和问题解决请参考： + + + +#### 2.4.3 镜像获取 + +链接: +提取码\:vuXJ + + +#### 2.4.4 运行时容器部署 + +通过`docker`命令加载本地下载好的编译器镜像。 + +```shell +sudo docker load --input xh-ascend-runtime-v1.0.tar +``` + +其中，`xh-ascend-runtime-v1.0.tar`为华为运行时镜像文件的名称。 + +使用`docker`命令创建新的容器实例。 + +```shell +sudo docker run --privileged=true -itd --ipc=host -u root --device=/dev/davinci0 --device=/dev/davinci4 --device=/dev/davinci_manager --device=/dev/devmm_svm --device=/dev/hisi_hdc -v /etc/timezone:/etc/timezone:ro -v /etc/localtime:/etc/localtime:ro -v /usr/local/Ascend/driver:/usr/local/Ascend/driver:ro -v /usr/local/Ascend/ascend-toolkit:/usr/local/Ascend/ascend-toolkit:ro -v /etc/ascend_install.info:/etc/ascend_install.info:ro -v /usr/local/sbin/npu-smi:/usr/local/sbin/npu-smi:ro -v /usr/local/sbin/:/usr/local/sbin/:ro -v /root/ascend/:/root/ascend/:rw --name cmcc-ascend-runtime --hostname cmcc-ascend-runtime xh-ascend-runtime:v1.0 bash +``` + +除了下述参数以外，均为`docker`自身运行相关的参数，如需了解更多，可参考`docker`的官方文档。 + +1. `--privileged=true`：为提权命令，固定参数。 + +2. `--ipc host`：为容器和宿主机之间共享内存，固定参数。 + +3. `-u root`：容器指定使用`root`用户权限运行，固定参数。 + +4. `--device=/dev/davinci0`、`--device=/dev/davinci4`、`-device=/dev/davinci_manager`、`--device=/dev/devmm_svm`以及`--device=/dev/hisi_hdc`：指定硬件描述符，通过下面的方法查询出结果以后，将所有的描述符都添加到命令行中，必要参数。 + + ```shell + # 要求设备已正常挂载并识别： + ll /dev/ | grep davinci + ll /dev/ | grep devmm + ll /dev/ | grep hisi + ``` + +5. `-v /etc/timezone:/etc/timezone:ro`和`-v /etc/localtime:/etc/localtime:ro`为挂载宿主机的时间信息，固定参数。 + +6. `-v /usr/local/Ascend/driver:/usr/local/Ascend/driver:ro`、`-v /usr/local/Ascend/ascend-toolkit:/usr/local/Ascend/ascend-toolkit:ro`、`-v /etc/ascend_install.info:/etc/ascend_install.info:ro`、`-v /usr/local/sbin/npu-smi:/usr/local/sbin/npu-smi:ro`、`-v /usr/local/sbin/:/usr/local/sbin/:ro`以及`-v /root/ascend/:/root/ascend/:rw`：挂载华为工具链到容器内部，固定参数。 + +7. `--name`：指定容器的名字，此参数用户可根据需要变更。 + +8. `--hostname`：指定主机名，此参数用户可根据需要变更。 + +9. `xh-ascend-runtime:v1.0`：指定华为运行时镜像的名称和版本，固定参数。 + +10. `bash`：指定容器内需要运行的进程命令，固定参数。 + +通过`docker`命令进入容器进行操作。 + +```shell +sudo docker exec ${CONTAINER_ID} bash +``` + +## 3 示例 + +### 3.1 功能说明 + +以基于`resnet50`深度学习模型的图像识别算法为例（示例代码见：`resnet50_multibs_pinned.cpp`），介绍基于算力原生**芯合**平台的智算应用的开发、编译以及跨架构部署过程。 + +`resnet50`深度学习模型基于`ImageNet 2012`公开图像数据集进行训练，可实现对动物、交通工具、日常用品等`1000`种类别的识别，识别准确率高达`96.4%`。`ImageNet 2012`数据集地址为： + +​ + +基于resnet50的图像识别智算应用主要包括以下步骤： + +1. 从`ImageNet 2012`数据集中获取图像； + +2. 图像数据预处理； + +3. 使用`resnet50`模型进行图像识别； + +4. 打印模型的识别结果（`top 5`信息）。 + +mlri文件获取： + +链接: 提取码:9e4J + +### 3.2 如何编译 + +* 将工作目录切换到样例代码resnet50\_multibs\_pinned.cpp的同级目录下 + + ```shell + cd RESNET50_SAMPLE_DIR # SAMPLE_DIR为样例代码resnet50_multibs_pinned.cpp的目录 + ``` + +* 执行编译脚本 + + ```shell + ./compile.sh + + # 编译脚本compile.sh的内容如下所示： + #!/usr/bin/env bash + set -e + + OUTPUT_PATH=/opt/share/pub_share/project/demo_resnet50 # 指定编译好的目标二进制文件的保存位置 + BINARY_NAME=resnet50_multibs_pinned_binary # 目标二进制文件的文件名 + CMCC_HOME=${CMCC_INSTALL} # 原生编译器的安装目录 + + source ~/anaconda3/etc/profile.d/conda.sh + conda activate cmcc # 进入cmcc虚拟环境 + + # 编译源代码 + cmcc \ + -fsycl-device-only \ # 仅编译设备端代码 + -tensor_compute \ # tensor_compute为添加编译参数，指定后编译器将执行Tensor Compute编译过程；需要注意的是，指定该参数后，需指定对应的mlir路径，否则会编译报错。 + -I./model \ # resnet50模型（mlir格式）的存放路径 + -I${CMCC_HOME}/llvm/build/install/include \ # 包含llvm安装目录下的头文件 + -I${CMCC_HOME}/llvm/build/install/include/sycl \ # 包含llvm安装目录下的sycl相关头文件 + -I${CMCC_HOME}/opencv/include/opencv4 \ # 包含OpenCV安装目录下的头文件 + -L ${CMCC_HOME}/llvm/build/install/lib \ # 链接LLVM安装目录下的库文件 + -L ${CMCC_HOME}/opencv/lib/ \ # 链接OpenCV相关库 + -lopencv_calib3d \ # 链接opencv的calib3d库 + -lopencv_core \ # 链接opencv的core库 + -lopencv_dnn \ # 链接opencv的dnn库 + -lopencv_features2d \ # 链接opencv的features2d库 + -lopencv_flann \ # 链接opencv的flann库 + -lopencv_highgui \ # 链接opencv的highgui库 + -lopencv_imgcodecs \ # 链接opencv的imgcodecs库 + -lopencv_imgproc \ # 链接oepncv的imgproc库 + -lopencv_ml \ # 链接opencv的ml库 + -lopencv_objdetect \ # 链接opencv的objdetect库 + -lopencv_photo \ # opencv的photo库 + -lopencv_stitching \ # 链接opencv的stitching库 + resnet50_multibs_pinned.cpp \ # 要进行编译的源代码文件名 + -o ${OUTPUT_PATH}/${BINARY_NAME} # 指定编译好的目标二进制代码的保存目录和文件名 + + # 编译完成，退出conda虚拟环境 + conda deactivate + ``` + +### 3.3 如何部署及运行 + +1. 将编译好的目标二进制文件拷贝到特定硬件环境，且确保该环境上已配置好对应的docker运行时环境 + + ```shell + scp resnet50_multibs_pinned username@ip_address:path of current hardware environment + ``` + +2. 将工作目录切换到目标二进制文件所在的目录下，执行二进制文件运行脚本 + + ```shell + # 在命令行执行此命令运行二进制可执行文件 + ./resnet50_multibs_pinned + ``` + +注意：当前应用可在瀚博、`NVIDIA`和华为环境下运行，只需要将编译好的二进制文件拷贝到瀚博/`NVIDIA`/华为环境下执行即可。 diff --git a/sample/resnet50_video_pinned.cpp b/sample/resnet50_video_pinned.cpp new file mode 100644 index 0000000..11506a7 --- /dev/null +++ b/sample/resnet50_video_pinned.cpp @@ -0,0 +1,232 @@ +#include +#include +#include +#include +#include + +#include +#include + +#define BATCHSIZE 2 +using namespace cv; +using namespace std::chrono; +using namespace sycl; + +// Provide the compiler with information such as input model, target hardware, etc +#pragma tensor_compute model(ResNet) input(resnet50_relay_bs2.mlir) subkind(tc_spirv) target_spec(VASTAI_OAK :.*, NVIDIA_CUDA :.*, HUAWEI_ASCEND :.*, CAMBRICON_MLU :.*, nullptr) + +const size_t array_size = 1000; +size_t imageLen = 0; + +constexpr cl::sycl::access::mode sycl_read = cl::sycl::access::mode::read; +constexpr cl::sycl::access::mode sycl_write = cl::sycl::access::mode::write; + +typedef struct +{ + int index; + float value; +} sort_st; + +class nnet_resnet50 +{ +public: + nnet_resnet50():sharp_width(224), sharp_height(224) {} + ~nnet_resnet50() {} + + void get_file_names(std::string basePath, std::vector &files); + int resnet_tensor_compute(std::vector fileNames, size_t count); + cv::Mat nnet_preprocess(cv::Mat src, int sharp_width, int sharp_height, float *data); +private: + size_t sharp_width; + size_t sharp_height; + uint32_t batch = BATCHSIZE; +}; + +void nnet_resnet50::get_file_names(std::string basePath, std::vector &files) +{ + DIR *dir; + struct dirent *ptr; + + if ((dir = opendir(basePath.c_str())) == NULL) + { + std::cerr << "Open dir error input Image error....\n"; + exit(1); + } + + while ((ptr = readdir(dir)) != NULL) + { + if (strcmp(ptr->d_name, ".") == 0 || strcmp(ptr->d_name, "..") == 0) + { + continue; + } + else if (ptr->d_type == DT_REG) //DT_REG: This is a regular file. + { + std::string a = ptr->d_name; + int pe = a.find_last_of("."); + std::string pic_name = a.substr(pe + 1); + if (pic_name == "JPEG" || pic_name == "jpg") + { + std::string tmpname = basePath + "/" + ptr->d_name; + files.push_back(tmpname); + } + } + else if (ptr->d_type == DT_DIR) //DT_DIR: This is a directory. + { + std::string base = basePath + "/" + ptr->d_name; + get_file_names(base, files); + } + } + closedir(dir); + return; +} + +static bool compares(sort_st a, sort_st b) +{ + return a.value > b.value; +} + +cl::sycl::queue deviceQueue; + +template +void testApp(T *input, TT *output) +{ + // Define arrays of model inputs and outputs + cl::sycl::buffer bufferA(input, imageLen * BATCHSIZE, {sycl::property::buffer::use_host_ptr()}); + cl::sycl::buffer bufferC(output, array_size * BATCHSIZE, {sycl::property::buffer::use_host_ptr()}); + + deviceQueue.submit([&](cl::sycl::handler &cgh) + { + auto accessorA = bufferA.template get_access(cgh); + auto accessorC = bufferC.template get_access(cgh); + + std::vector inputs; + std::vector outputs; + + inputs.push_back(&accessorA); + outputs.push_back(&accessorC); + cgh.tensor_compute_task(inputs, outputs); + }).wait(); + + host_accessor h_c(bufferC, read_only); +} + +int nnet_resnet50::resnet_tensor_compute(std::vector fileNames, size_t count) +{ + auto alltime_start = system_clock::now(); + unsigned char *tempD; + + imageLen = this->sharp_width * this->sharp_height * 3; + cl::sycl::cl_float *imageData = (cl::sycl::cl_float *)malloc_pinned(imageLen * batch * sizeof(cl::sycl::cl_float), deviceQueue); + cl::sycl::cl_float *result = (cl::sycl::cl_float *)malloc_pinned(array_size * batch * sizeof(cl::sycl::cl_float), deviceQueue); + + for (int i = 0; i < count; i++) + { + std::string file_name = fileNames.back(); + fileNames.pop_back(); + std::cout << "\npic-name:" << file_name << std::endl; + cv::Mat img = cv::imread(file_name.c_str(), cv::IMREAD_COLOR); + cv::Mat dst; + { + float data[sharp_height * sharp_width * 3]; + std::cout << "sharp_height " << sharp_height << " , sharp_width " << sharp_width << std::endl; + dst = nnet_preprocess(img, this->sharp_width, this->sharp_height, data); + float *p = data; + int pos = 0; + for (int i = 0; i < imageLen * batch; ++i) + { + imageData[i] = *(p + pos); + pos = (pos + 1) % imageLen; + } + + auto hardprocess_starttime = system_clock::now(); + testApp(imageData, result); + } + + std::cout << std::endl; + std::vector result_vector(array_size * batch); + for (int m = 0; m < batch; m++) + { + int index = 0; + for (int n = m * (array_size); n < (m + 1) * (array_size); n++) + { + result_vector[index].value = result[n]; + result_vector[index].index = index; + index++; + } + + sort(result_vector.begin(), result_vector.end(), compares); + std::cout << "Top 5 is: "; + for (int i = 0; i < 5; ++i) + { + std::cout << result_vector[i].index << ": " << result_vector[i].value << " "; + } + std::cout << std::endl; + } + } + + free(imageData, deviceQueue); + free(result, deviceQueue); + duration alltime_end = system_clock::now() - alltime_start; + + std::cout << std::endl << "Finished!\n"; + return 0; +} + +// Set the image ROI +cv::Mat nnet_resnet50::nnet_preprocess(cv::Mat src, int sharp_width, int sharp_height, float *data) +{ + cv::Mat src_img; + cv::cvtColor(src, src, cv::COLOR_BGR2RGB); + + size_t resize = 256; + size_t resize_w = 0; + size_t resize_h = 0; + + if (src.cols > src.rows) + { + resize_h = resize; + resize_w = int(resize * float(src.cols) / float(src.rows)); + } + else + { + resize_w = resize; + resize_h = int(resize * float(src.rows) / float(src.cols)); + } + + size_t crop_x = int((resize_w - sharp_width) / 2); + size_t crop_y = int((resize_h - sharp_height) / 2); + + cv::resize(src, src_img, cv::Size(resize_w, resize_h)); + Rect roi(crop_x, crop_y, sharp_width, sharp_height); + + std::vector mean_value{0, 0, 0}; + std::vector std_value{1, 1, 1}; + + cv::Mat dst; + uint32_t count = 0; + src_img(roi).copyTo(dst); + for (int i = 0; i < sharp_height; i++) + { + uchar *uc_pixel = dst.data + i * dst.step; + for (int j = 0; j < sharp_width; j++) + { + data[count] = (uc_pixel[0] / 255. - mean_value[0]) / std_value[0]; + data[count + sharp_height * sharp_width] = (uc_pixel[1] / 255. - mean_value[1]) / std_value[1]; + data[count + 2 * sharp_height * sharp_width] = (uc_pixel[2] / 255. - mean_value[2]) / std_value[2]; + uc_pixel += 3; + count++; + } + } + return dst; +} + +int main(int argc, char **argv) +{ + nnet_resnet50 *resnet50 = new nnet_resnet50(); + + std::vector fileNames; + resnet50->get_file_names("/opt/share/pub_share/c_images/dataset/2012img", fileNames); + size_t count = fileNames.size(); + resnet50->resnet_tensor_compute(fileNames, count); + return 0; +} \ No newline at end of file