[嵌入式AI从0开始到入土]17_Ascend C算子开发

已于 2024-05-04 17:57:30 修改
阅读量1.7k 收藏 14
点赞数 14
分类专栏： 嵌入式ai从0到入土 文章标签： 人工智能 c语言 开发语言 ascend 华为
于 2024-05-04 17:47:50 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/weixin_44354614/article/details/138424102
版权

[嵌入式AI从0开始到入土]嵌入式AI系列教程

注：等我摸完鱼再把链接补上
可以关注我的B站号工具人呵呵的个人空间，后期会考虑出视频教程，务必催更，以防我变身鸽王。

第1期 昇腾Altas 200 DK上手
第2期 下载昇腾案例并运行
第3期 官方模型适配工具使用
第4期 炼丹炉的搭建（基于Ubuntu23.04 Desktop）
第5期 炼丹炉的搭建（基于wsl2_Ubuntu22.04）
第6期 Ubuntu远程桌面配置
第7期 下载yolo源码及样例运行验证
第8期 在线Gpu环境训练（基于启智ai协作平台）
第9期 转化为昇腾支持的om离线模型
第10期 jupyter lab的使用
第11期 yolov5在昇腾上推理
第12期 yolov5在昇腾上应用
第13期_orangepi aipro开箱测评
第14期 orangepi_aipro小修补含yolov7多线程案例
第15期 orangepi_aipro欢迎界面、ATC bug修复、镜像导出备份
第16期 ffmpeg_ascend编译安装及性能测试
第17期 Ascend C算子开发
未完待续…

文章目录

前言

我在24年3月和我的小伙伴一起参加了第一届昇腾AI原生创新精英挑战赛，在这里做一下总结。这里以orangepi Ai Pro为例。
注：我们的代码仓最早将于24.05.10开放，大家可以直接看op_kernel内的compute，kernelLaunch内可能有错，实在来不及改了
代码仓地址：https://gitee.com/toolsmanhehe/acl_ops

一、环境配置

我们基于正常能够使用的镜像作为基础镜像。这里我推荐使用minimal镜像。这样就不用先卸载cann了，甚至你可以直接删除/opt/compress目录，反正咱后面直接远程连接敲代码了，也用不上。

1、CANN包安装

wget https://ascend-repo.obs.cn-east-2.myhuaweicloud.com/Milan-ASL/Milan-ASL%20V100R001C17SPC702/Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run
chmod +x Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run

#卸载旧的CANN
./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run --uninstall
sudo rm -rf /usr/local/Ascend/ascend-toolkit/*

#安装指定版本的CANN
./Ascend-cann-toolkit_8.0.RC1.alpha002_linux-aarch64.run --install
#安装依赖
pip install protobuf==3.20.0
#添加环境变量
echo “source /usr/local/Ascend/ascend-toolkit/set_env.sh” >> /home/HwHiAiUser/.bashrc
source /home/HwHiAiUser/.bashrc

2、配置ssh密钥（可选）

主要是vscode等ide连接时都需要输入密码，比较麻烦。
这里可以参考我之前的文章来实现免密登录，在七、问题 的第5点

3、配置git（可选）

因为我们三个人在三个城市，因此为了方便讨论和开发，我们建立了代码仓库，但是每次推送和拉取都需要账号密码（在完赛前是不可能公开的），这不符合本懒人的性格啊。
这里我们需要在开发环境上执行

cd ~
touch .git-credentials
vim .git-credentials
#输入以下内容，请自行替换username和password
https://username:password@gitee.com
git config --global credential.helper store

二、获取sample样例

cd 
git clone https://gitee.com/ascend/samples.git

在不修改算子名称，输入输出的时候，我们只需要关注图中框出来的文件即可。

1、add算子

打开目录operator/AddCustomSample

1、KernelLaunch

在这里插入图片描述
我们的调用顺序是main.cpp->add_custom_do->add_custom->op.Init->op.Process。
在这里插入图片描述
因为我们要实现的算子的Z=X+Y，因此我们需要将这三个变量传入计算过程。
虽然这里只有2个输入，但是输出也需要申请内存，因此是3个输入参数

然后我们需要申明相关的变量和常量（这里使用静态shape）。
在这里插入图片描述
在这里插入图片描述
接着就是初始化，为各个张量申请内存
在这里插入图片描述
接着就是计算过程，这里因为使用的是静态shape，因此循环次数是定值（芯片内存空间有限，不可能一次性全部计算完成）
在这里插入图片描述
在copyin的时候从xGM和yGM分别取出TILE_LENGTH个数据，存入xLocal和yLocal以供compute使用。
在compute结束以后，我们需要先使用outQueueZ.EnQue来表示计算完成，但是此时不能释放zLocal的内存，因为我们还没有保存到zGM。
在copyout环节，将输出结果存入zGM。

接着我们看生成测试数据的程序，这里我们生成了2条16384个1~100随机half格式的数据。我们最后可以直接对比output/golden.bin和output/output_z.bin的md5值来判断算子正确与否。或者修改scripts/verify_result.py直接打印误差数量。
在这里插入图片描述
最后来到KernelLaunch目录执行以下命令，测试核函数正确性。
务必先进行cpu测试，通过后执行npu测试，在npu下有些报错不显示

su			   #使用root用户执行，否则可能报错
bash run.sh -r cpu -v ascend310B1   #cpu测试
bash run.sh -r npu -v ascend310B1   #npu测试

以下为cpu测试结果
在这里插入图片描述
以下为npu测试结果
在这里插入图片描述
测试均通过的情况下，我们就可以进行下一步的framework的编写了

2、Framework

我们先看AddCustomSample/FrameworkLaunch/AddCustom.json这个文件，上面为输入变量，下面为输出变量。我们需要使用这个配置文件来生成framework工程。此处的变量应该和工程内的一致。
在这里插入图片描述
接着我们看工程。
在这里插入图片描述
op_host没什么可说的，可以去看本文下一个案例Addcdiv。
op_kernel基本上就是把上面在kernelLaunch中测试通过的代码cv过来。
注意图中的地方就可以了，这个tiling是从host侧传入的。然后在开头将静态shape删除了，因为这里我们是通过op_host实现的动态shape的切分，然后传入kernel侧的。
在这里插入图片描述
接下来修改CMakePresets.json，将框出来的地方改成你的CANN路径。
在这里插入图片描述
最后，我们进入framework目录，编译算子并安装

bash build.sh
./build_out/custom_opp_ubuntu_aarch64.run

在这里插入图片描述

3、AclNN

在算子大赛的时候，这个是由官方发布的（就是可能有错误），我们直接使用即可，一般测试能通过，就会有4-8分（10分满分）。
在这里插入图片描述
这里的gen_data和kernelLaunch里是一样的，我们执行以下命令，验证算子正确与否。

bash run.sh

测试通过会有如下提示
在这里插入图片描述

2、Addcdiv算子

打开目录operator/AddcdivCustomSample
大部分与add算子相似，因此我们这里只看op_host和op_kernel部分。
在头文件中你会发现多了许多东西,所有的东西我们都需要传入kernel侧。具体实现过程就去阅读代码吧，就是这个案例也是赶出来的，可能里面的切分策略不是最好的，但是确实是能用的。

#ifndef ADDCDIV_CUSTOM_TILING_H
#define ADDCDIV_CUSTOM_TILING_H
#include "register/tilingdata_base.h"

namespace optiling {
BEGIN_TILING_DATA_DEF(AddcdivCustomTilingData)
  TILING_DATA_FIELD_DEF(float, value); 	//参与计算的标量
  TILING_DATA_FIELD_DEF(uint32_t, blockLength);
  TILING_DATA_FIELD_DEF(uint32_t, tileNum);
  TILING_DATA_FIELD_DEF(uint32_t, tileLength);
  TILING_DATA_FIELD_DEF(uint32_t, lasttileLength);
  TILING_DATA_FIELD_DEF(uint32_t, formerNum);
  TILING_DATA_FIELD_DEF(uint32_t, formerLength);
  TILING_DATA_FIELD_DEF(uint32_t, formertileNum);
  TILING_DATA_FIELD_DEF(uint32_t, formertileLength);
  TILING_DATA_FIELD_DEF(uint32_t, formerlasttileLength);
  TILING_DATA_FIELD_DEF(uint32_t, tailNum); 
  TILING_DATA_FIELD_DEF(uint32_t, tailLength);
  TILING_DATA_FIELD_DEF(uint32_t, tailtileNum);
  TILING_DATA_FIELD_DEF(uint32_t, tailtileLength);
  TILING_DATA_FIELD_DEF(uint32_t, taillasttileLength);    
END_TILING_DATA_DEF;

REGISTER_TILING_DATA_CLASS(AddcdivCustom, AddcdivCustomTilingData)
}
#endif // ADDCDIV_CUSTOM_TILING_H

以下为op_kernel内的部分代码

 private:
  TPipe pipe;
  // TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueY, inQueueZ;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueIN;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueOUT;
  GlobalTensor<half> xGm;
  GlobalTensor<half> yGm;
  GlobalTensor<half> zGm;
  GlobalTensor<half> outGm;
  half value;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileLength;
  uint32_t lasttileLength;
  uint32_t formerNum;
  uint32_t formerLength;
  uint32_t formertileNum;
  uint32_t formertileLength;
  uint32_t formerlasttileLength;
  uint32_t tailNum;
  uint32_t tailLength;
  uint32_t tailtileNum;
  uint32_t tailtileLength;
  uint32_t taillasttileLength;
};

extern "C" __global__ __aicore__ void addcdiv_custom(GM_ADDR x, GM_ADDR y,
                                                     GM_ADDR z, GM_ADDR out,
                                                     GM_ADDR workspace,
                                                     GM_ADDR tiling) {
  GET_TILING_DATA(tiling_data, tiling);
  // TODO: user kernel impl
  KernelAddcdiv op;

  uint32_t tilingKey = 1;
  if (TILING_KEY_IS(1)) {
    tilingKey = 1;
  } else if (TILING_KEY_IS(2)) {
    tilingKey = 2;
  } else {
    tilingKey = 1;
  }

  op.Init(x, y, z, out, tiling_data.value, tiling_data.blockLength,
          tiling_data.tileNum, tiling_data.tileLength,
          tiling_data.lasttileLength, tiling_data.formerNum,
          tiling_data.formerLength, tiling_data.formertileNum,
          tiling_data.formertileLength, tiling_data.formerlasttileLength,
          tiling_data.tailNum, tiling_data.tailLength, tiling_data.tailtileNum,
          tiling_data.tailtileLength, tiling_data.taillasttileLength,
          tilingKey);
  op.Process();
}

#ifndef __CCE_KT_TEST__
// call of kernel function
void addcdiv_custom_do(uint32_t blockDim, void* l2ctrl, void* stream,
                       uint8_t* x, uint8_t* y, uint8_t* z, uint8_t* out,
                       uint8_t* workspace, uint8_t* tiling) {
  addcdiv_custom<<<blockDim, l2ctrl, stream>>>(x, y, z, out, workspace, tiling);
}
#endif

三、编写自己的算子

1、搭建框架

我们可以使用参考add算子搭建以下目录结构。以下文件夹内的文件没有特别说明就直接从add算子工程内复制。

myCustom
├── AclNNInvocation
│   ├── inc
│   ├── scripts
│   └── src
│   ├── run.sh
├── myCustom 		<-由msopgen工具生成
├── KernelLaunch
│   ├── myCustom.cpp
│   ├── cmake
│   ├── CMakeLists.txt
│   ├── data_utils.h
│   ├── run.sh
│   └── scripts
└── myCustom.json

2、 KernelLaunch编写

1、myCustom.cpp

我们直接cv add算子的，对输入做下修改，然后修改compute就行了。

2、main.cpp

这里主要是将算子名称以及传入的参数修改下
在这里插入图片描述

3、scripts/gen_data.py

这里根据你要实现的代码编写生成数据和真值的程序就行了，在比赛时，我们可以直接从官方给出的AclNN中取。

3、 framework编写

在kernelLaunch测试通过后我们直接修改myCustom.json。如果是多个数据类型，如下所示。

[
    {
        "op": "myCustom",
        "language": "cpp",
        "input_desc": [
            {
                "name": "x",
                "param_type": "required",
                "format": [
                    "ND","ND"
                ],
                "type": [
                    "fp16","fp32"
                ]
            }
        ],
        "output_desc": [
            {
                "name": "y",
                "param_type": "required",
                "format": [
                    "ND","ND"
                ],
                "type": [
                    "fp16","fp32"
                ]
            }
        ]
    }
]

然后生成工程（具体目录请自行修改）

/usr/local/Ascend/ascend-toolkit/latest/python/site-packages/bin/msopgen gen -i /home/HwHiAiUser/myCustom/myCustom.json -c ai_core-ascend310B1 -lan cpp -out /home/HwHiAiUser/myCustom/myCustom

接着就是参考add和addcdiv算子在op_host中实现tiling策略，将kernelLaunch中测试通过的代码加上tiling相关的代码后搬运到op_kernel。编译安装算子。

4、 Aclnn测试

这里因为我做的是比赛里给出的题目，因此直接使用官方给的案例进行测试。对于自定义算子，除修改gen_data外，我们还需要修改op_runner以及main.cpp。

四、torch_npu重新编译（可选）

参考仓库说明：https://gitee.com/ascend/op-plugin

五、常用api

为了简化使用，以下仅列出常用的2级接口，如需高性能实现，请使用0级接口。310b系列似乎不支持高级api，因此也不列出了。详细内容请直接看api文档

名称功能表达式二级接口样例
Exp按元素取自然指数在这里插入图片描述Exp(dstLocal, srcLocal, 512);
Abs按元素取绝对值在这里插入图片描述Abs(dstLocal, srcLocal, 512);
Reciprocal按元素取倒数在这里插入图片描述Reciprocal(dstLocal, srcLocal, 512);
Sqrt按元素做开方在这里插入图片描述Sqrt(dstLocal, srcLocal, 512);
Ln按元素取自然对数在这里插入图片描述Ln(dstLocal, srcLocal, 512);
Add按元素求和在这里插入图片描述Add(dstLocal, src0Local, src1Local, 512);
Mul按元素求积在这里插入图片描述Mul(dstLocal, src0Local, src1Local, 512);
Adds/Muls矢量内每个element与标量求和/积同上Adds(dstLocal, srcLocal, half(2), 512);
Sub按元素求差在这里插入图片描述Sub(dstLocal, src0Local, src1Local, 512);
Div按element求商在这里插入图片描述Div(dstLocal, src0Local, src1Local, 512);
Max按element求最大值在这里插入图片描述Max(dstLocal, src0Local, src1Local, 512);
Min按element求最小值在这里插入图片描述Min(dstLocal, src0Local, src1Local, 512);
Duplicate将一个变量或一个立即数，复制多次并填充到向量在这里插入图片描述Duplicate(dstLocal, half(18.0), 512);

注意：标量双目指令中没有减法和除法，基础api没有log只有ln。

问题

一句话，多看文档，有问题就先去社区搜一下。160001，error code 0这种就直接查代码吧，没有具体原因。

1、fatal error: register/tilingdata_base.h: No such file or directory

在这里插入图片描述
在这里插入图片描述
检查一下CANN路径

其他能稳定复现的bug等我后面遇到了再补充解决办法吧。

总结

也许，有时歪门邪道比正道更简单。不要被文档和案例限制了，不要问能不能，跑下试试最快。
就像adds直接乘标量不好使，那就直接把这个标量填满整个local，直接用张量去计算嘛。而且这样能用的api还更多呢。

工具人呵呵
关注
  • 14
    点赞
  • 踩
  • 14
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 9
    评论
专栏目录
「嵌入式 AI」这个时髦的名词，究竟是一个新理念，还是一个旧噱头？
机器之心
11-06 2123
撰文 | 高静宜 编辑 | 宇多田 在人工智能迅猛发展的当下，几乎每一家企业都在探索行业的切入点，试图在人工智能大潮中分一杯羹。 而在追寻 AI 商业化的道路上，人们逐渐发现，要实现技术的落地，不仅需要性能优越的算法模型以及可靠的硬件支持，还需要把 AI 技术和硬件环境进行有机结合，再应用到具体的实际场景中，进而满足用户的需求。 而「嵌入式 AI」就是当前最热门
[嵌入式AI从0开始到入土]12_yolov5在昇腾上应用
工具人呵呵的博客
01-07 1648
本文参考案例程序，完善了yolov5在Atlas200DK上的推理应用代码
新风口下：嵌入式AI学习中较好的练手项目（附代码资料/学习视频/学习规划）...
嵌入式Linux
10-27 3160
有粉丝问我：“当前乃至未来5-10年，嵌入式开发者还有哪些风口？”画外音：风口的本质，其实就是一段时间的人才供需不平衡。说白了就是由于行业突变，敏锐的资本快速进入，导致短时间内行业大量扩张...
嵌入式系统开发教程 - 创新课程设计- “AI人脸识别+扩展系统”
weixin_46771530的博客
04-25 2525
文章目录前言一、人脸识别功能如何实现二、快递柜的软件系统三、人脸测温总结 前言 这里只是一些网上资料的合集，具体的设计方案需要各位结合项目情况考虑。 一、人脸识别功能如何实现 简单可以总结为用一个树莓派小系统作为一个本地终端，用来采集摄像头数据上传到网上，在云终端进行人脸识别并返回人脸识别信息，本地终端接收到信息后进行对应的操作：控制舵机、电机或其他器件。 具体可以参考B站视频: 树莓派人脸识别门禁系统快速上手系列教程. 本系统可作为你的毕设中扩展的一个小功能，配合Arduino控制舵机的树莓派人脸识
嵌入式AI每日学习之路
bucan804228552的博客
02-08 1644
【2.7】虚拟机+LInux学习 vmware中打开浏览器不能上网问题： 1.点击VMware Workstation右下角的 “网络适配器” 小图标 确认是NAT模式 2.ctrl+alt+t打开命令窗口，输入 $ sudo service network-manager stop $ sudo rm /var/lib/NetworkManager/NetworkManager.state $ sudo service network-manager start 3.重新打开 Firefox浏览器 Li
嵌入式AI基础知识介绍
Jocelin47的博客
08-16 2545
一、嵌入式AI概述 大部分神经网络计算都在卷积层，基本上99%都在卷积层， 二、模型加速 2-1.权值量化 tengine量化后的准确率之下降了1.5%。 2-2.网络剪枝 2-3.知识蒸馏 大网络是一个老师，小网络是一个学生 让小网络逐渐学习到大网络的能力 三、边缘侧的训练方法 传统的机器学习训练方法是把数据在服务器上训练好，然后使用 联邦学习是用户不把数据上传，而是把模型下发到对应的用户那里 用户训练以后用户会把训练好的模型上传到服务器，服务器再把训练的模型进行合成，
Embed_C.rar_C语言程序设计_嵌入式C语言_嵌入式C语言编程_嵌入式c
09-19
"嵌入式C语言程序设计"是一个深入研究如何利用C语言进行嵌入式系统开发的主题，它涵盖了从基本概念到高级技巧的广泛知识领域。 首先，"嵌入式C语言程序设计基础"文档可能包含了对C语言基础知识的讲解，包括变量、...
ARM--PDAtest.rar_44b0_arm pda_嵌入式_嵌入式系统设计与开发实例详解
09-14
标题中的"ARM--PDAtest.rar_44b0_arm pda_嵌入式_嵌入式系统设计与开发实例详解"揭示了本压缩包主要包含与ARM处理器、PDA（个人数字助手）以及嵌入式系统设计和开发相关的实例教程。其中，S3C44B0是一款基于ARM7TDMI...
嵌入式系统开发的好文章.rar_嵌入式_嵌入式 linux_嵌入式linux_嵌入式系统_嵌入式系统开发
09-14
嵌入式系统无处不在，从智能手机到家用电器，再到汽车电子系统，它们都是现代生活的重要组成部分。这篇文章集合将深入探讨嵌入式系统开发的关键概念和技术，特别关注嵌入式Linux的使用。 首先，我们要理解嵌入式...
C语言主从机开发_C-C++_嵌入式_主从机开发_
10-01
在嵌入式系统设计中，主从机模式是一种常见的通信架构，它允许多个设备或处理器协同工作，其中一个作为主设备（主机），控制整个系统的工作流程，而其他设备作为从设备（从机）响应主机的指令。这个压缩包文件包含两...
从零开始学习ARM嵌入式系列——Linux驱动开发简介-综合文档
05-24
从零开始学习ARM嵌入式系列——Linux驱动开发简介
图像识别VPU——易用的嵌入式AI支持深度学习平台介绍
weixin_34396103的博客
07-21 2648
公司玩了大半年的嵌入式AI平台，现在产品进入量产模式，也接触了很多嵌入式方案，有了一些心得体会，本人不才，在这里介绍一下一款简单易用的嵌入式AI方案——Movidius Myriad 2 VPU(MA2450) 和 Myriad X VPU(MA2485)。这里本人重点提示：简单易用的嵌入式AI。现在好多家半导体厂商已经推出嵌入式AI平台，比如华为海思今年4月份发布的Hi3559A，这个样品超过1...
【嵌入式AI入门日记】将 AI 模型移植到 RT-Thread 上（1）
FILLMOREad的博客
07-07 1202
本期我们分享主题是如何将 AI 模型部署到嵌入式系统中，下一期将介绍如何在 RT-Thread 操作系统上运行 Mnist Demo（手写数字识别）。 嵌入式关联 AIAI落地一直是一...
课程：《嵌入式Linux工程师实战进阶人工智能》
Linux阅码场
07-24 2217
课程背景人工智能似乎一直上在云端的高高在上的概念，但是随着5G、IoT、边缘计算等概念和技术广泛深入地推进，人工智能与嵌入式Linux的关系逐步拉近。时下火遍全宇宙的概念...
【自然语言处理】概论（一）：自然语言处理概要
古月居
07-25 598
语言：语言就广义而言，是采用一套具有共同处理规则来进行表达的沟通指令，指令会以视觉、声音或者触觉等方式来传递。可以分为:自然语言、动物语言、计算机语言。。。自然语言人类在发展过程中形成的一种信息交流方式，包括口语和书面语，反映人类的思想。它自然地随文化演化，是人类用来沟通的指令，可以通过视觉、声音或触觉等方式传递。定义：自然语言处理是计算机科学和人工智能的一个分支，它旨在开发能够理解和生成人类语言的技术。通过建立形式化计算模型来分析、理解和生成自然语言的学科。
2024钉钉杯A题思路详解
最新发布
全栈川川
07-26 1260
2024钉钉杯思路
随机数种子的作用
帆的博客
07-23 1284
设置随机数种子（random seed）的目的是为了确保随机数生成器在每次运行时产生相同的随机数序列，从而保证实验结果的一致性。随机数种子通过初始化随机数生成器的内部状态，使得在相同的种子值下，随机数生成器每次调用时生成的序列是相同的。
常见的文心一言的指令
知码客
07-25 966
文心一言，作为百度研发的预训练语言模型“ERNIE 3.0”的一项功能，能够与人对话互动，回答问题，协助创作，高效便捷地帮助人们获取信息、知识和灵感。
Linux入门与嵌入式开发：从安装到C编程
第三章着重于Linux下的C编程基础，从C语言简介、编程环境配置（如Vi和Emacs编辑器）、Gcc编译器、Gdb调试器以及Make工程管理器的使用等方面，全面介绍了编程和调试工具的使用方法，这对于嵌入式Linux应用开发至关...
评论 9
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

工具人呵呵

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值