版本须知¶
TopsPlatform 1.2.0¶
TopsPlatform_v1.2.0版本KMD提升了底层驱动的稳定性，Runtime优化launch kernel和多流下发的性能，提升模型效率。 编译器优化dte api codesize，减少icache miss；优化编译器产生的指令效率，提升算子性能。TopsProfiler、TopsGDB、Efsmi等工具也做了改进和优化，方便用户调试使用。
新增¶
TopsGraph mode支持factor算子和topscc算子混跑
新增runtime接口topsExecutableGetBinaryPtr()
新增runtime接口topsStreamCreateWithLaunchLimit API
默认使能L2C/LLC，提升模型性能
新增查询函数，查询buffer是L2还是L3，可通过topsPointerGetAttribute接口查询，topsPointerGetAttributes接口未支持
新增FW热更新
- 新增efsmi接口 efsmi -mcm single/dual -i x，用于配置卡的hash模式，该功能有如下limitation:
仅支持Asic，Asic Passthrough切换hash mode，SRIOV虚拟化环境下仅支持读取hash mode
驱动重新安装或机器重启后，mcm变回默认状态
docker里面切换mcm，需要docker有privileged权限
不支持多个docker并行切换（即使docker内是不同的卡）
新增TopsVisualProfiler选择视图导出csv时支持平铺list导出
新增TopsProfiler支持显示kernel隐式参数信息
新增TopsProfiler支持使能SIP的部分profiling事件而非全部事件
变更¶
efsmi reset移除-F参数
libefdrv.so和libefdrv_static.a以及相关软连接文件从topsruntime package中移除，头文件不变，一些必要的符号从libefdrv中移动到libefrt。如果之前没有直接依赖libefdrv，但是编译脚本中有类似-lefdrv这样的链接语句，需要删除，对于已经编译好的，依赖libefdrv的动态库或者可执行文件，需要重新编译
stream packet 默认使用单线程下发
Hwsync替换gsync，更新TopsPlatform_v1.2.0.4版本后，需拿1.2.0.4的TopsPlatform版本重新编译算子，否则会出现不兼容问题
修复及优化¶
解决kernel param size大且并发度高场景下，导致kernel param分配失败的问题
打开2M device memory pool功能，提升性能
graph support topsLaunchCooperativeKernel
launch kernel优化
多stream性能优化
fix graph capture下申请大内存OOM
efsmi安全增强，在SRIOV模式下，如果检查到当前存在VF，则禁止切换ECC模式或者主动触发reset，并且给用户删除VF的提示
fix 整卡透传VM 触发ras 偶现 ap timeout
fix efsmi -mcm single/dual未检查gcu是否有任务占用，执行该切换命令会导致任务直接中断
fix 容器非特权模式下，efsmi不支持读取加速卡sleep状态
- PF FLR后对VF处理优化：
支持host直接使用VF的场景下，PF FLR，VF销毁后自动重建VF。重新创建的VF，会进行重新编号，设备的编号和发生reset前可能不同，测试脚本需要重新check设备编号再使用；
Docker下使用VF的场景，PF FLR后，VF销毁后也自动创建VF，但需运维人员手动重启docker；
VF透传到VM的场景下，PF RAS后，KMD会halt住SP，终止VF中的业务继续下发，然后主动做PF FLR。后续需运维人员介入进行如下操作：关闭虚拟机、VF重新绑定回KMD、手动删除VF、再重新创建VF
fix 设备在reset后，可用device memory减少的问题
fix S60在多卡环境下profiler，如指定其中一张卡会报错 “topsDeviceGetPCIBusId failed error”，功能无影响
fix 多用户多进程profiling场景时的非法内存访问和机器重启
fix 1vf场景下profiling时，Die1 sip错误得显示在Die0
fix Profiling时，如全开所有的engine则会出现丢profiler数据的情况。建议减少profiling的时间或通过单算子场景做cdte/sdte/odte抓取vpd文件做性能分析
fix TopsProfiler不只是显示tops50的kernel，需要打印所有kernel信息
fix TopsGDB 除0号卡外无法debug的问题
fix虚拟化环境下，开启cdte sdte extra sip&sp等大数据量profile场景下，数据分发处理异常，导致服务器hang
TopsProfiler –enable-activities选项改为optional，不指定时默认使能operator与memcpy
fix 2vf环境下，开启cdte sdte extra sip&sp等大数据量进行profiling会存在时间戳为负数（丢失部分数据，使得数据没有对齐）
TopsProfiler按照stream对GCU Kernels进行分组
TopsProfiler支持区分不同的ODTE搬运方向(DToH, HToD, DToD)
TopsVisualProfiler优化事件视图中的属性显示
编译器codesize优化
编译器实现c slot替换优化，提升算子性能
已知问题¶
L2下沉管理默认开启后，会导致__topsMallocSharedMem接口在多线程多进程情况下概率性crash，后续在DRS方案中修复
1vf环境下，开启cdte sdte extra sip&sp等大数据量进行profiling，kmd dmesg信息会有call trace
TopsPlatform 1.1.0¶
TopsPlatform_v1.1.0版本上TopsCC及TopsRuntime进行了性能优化，增强了资源隔离，提供运行时库Graph功能将整个计算流定义为一个图而不是单个操作的列表，以优化计算流性能。同时，TopsProfiler、TopsGDB、Efsmi等工具也做了改进和优化，方便用户调试使用。
新增¶
新增TopsCC的DTE增量配置优化
新增多进程支持（L2下沉管理）
在topsDeviceSetLimit接口中新增topsLimitMultiProcessorCount和topsLimitMaxThreadsPerBlock选项，用于设置线程上下文中可用的最高processor count和thread per block
launch kernel开销优化：12个sip的启动斜率开销加boot code时间开销优化至3.5us
支持基于stream capture方式构建runtime graph，仅支持如下接口（__topsPushKernelNodeName、__topsPopKernelNodeName、topsStreamBeginCapture、topsStreamEndCapture、topsStreamGetCaptureInfo、topsStreamGetCaptureInfo_v2、topsStreamIsCapturing、topsStreamUpdateCaptureDependencies、topsThreadExchangeStreamCaptureMode、topsGraphGetNodes、topsGraphNodeGetType、topsGraphKernelNodeGetAttribute、topsGraphDebugDotPrint、topsGraphDestroy）
支持基于runtime graph查询算子使用block shared memory大小
支持用户配置算子使用的block shared memory的开始地址
支持命令行安装TopsPlatform.run包增加–no-kernel-modules参数，不自动编译和加载kmd ko
新增TopsGDB支持分析程序生成的coredump功能
TopsProfiler支持按Stream显示timeline
TopsProfiler支持提示“DTE partial write”
在S60上支持多媒体Video进程强杀
S60G默认只支持2VF驱动安装方式
默认支持Sleep功能
提高GCU boot成功率
- Efsmi的改进和优化：
vGCU状态改为GCU virt，并显示为disabled、vgcu、mdev、sriov
efsmi -r默认的reset行为改为FLR
hot reset流程，由kmd ioctl改为unbind+ sysfs reset + bind
efsmi -ecc on/off + efsmi -r hot 即可完成L3容量更新
ECC在VF上不可修改，仅只读
变更¶
TopsPlatform.run不再包含TopsVisualProfiler的windows和mac安装包，TopsVisualProfiler的windows和mac安装包单独提供
修复¶
循环中更新leaptr的offset， 出现tar spill happen
ctrl+c或者SIP超时后，这之前的printf信息不能输出到控制台
当使用-O0对topscc kernel进行编译时，执行该kernel的操作会默认为同步操作
S60 4VF下，虽然property中查询Shared Memory上限为32MB，但运行kernel时使用动态分配32MB Shared Memory时（如kernel<<<1,1,SIZE_32MB>>>），会发生crash。（1VF/2VF/asic运行模式下，Shared Memory上限为64MB，kernel使用动态分配64MB Shared Memory时，如kernel<<<1,1,SIZE_64MB>>>，可以正常运行）
修复pyefml的权限问题
TopsProfiler修改launch kernel与算子的统计方式，默认一个launch kernel对应一个算子
多线程场景下进行profiling，launch kernel数量偶现存在错误
已知问题¶
L2下沉管理默认开启后，会导致__topsMallocSharedMem接口在多线程多进程情况下概率性crash，后续在SVG/PVG方案中修复
S60 4vf打开所有profiler引擎会引起call trace，功能无影响
S60在多卡环境下profiler，如指定其中一张卡会报错 “topsDeviceGetPCIBusId failed error”，功能无影响
TopsPlatform 1.0.2¶
TopsPlatform_v1.0.2版本上TopsCC及TopsRuntime实现了部分的内部优化，TopsProfiler在S60/S60G上支持单机多卡profiling，TopsGDB对GDB版本作了升级
新增¶
TopsCC优化整型vector除法性能
TopsCC增强bank conflict算法，提升程序性能
支持Event底层使用hardware timer计算时间
默认打开kmd odte ifb和vmsys prefetch
支持S60/S60G单机多卡profiling
加载Driver时默认打开kernel profiler
Profiling时支持RAS enable
更新GDB版本，从12.1更新到13.2 GDB 13.2 Changes
变更¶
NA
修复¶
S60/S60G不支持SIP上的kernel程序运行时的强杀
topsEventRecord会造成host memory的泄露
修复DTE操作的size >= 4GB时报错返回
kernel越界访问或L2L3多sip搬运等错误情况就会产生DTE、PTE问题
S60/S60G上使用一个event来同步多个stream时无效，可能造成硬件错误
缺少 acosh asinh atanh cbrt copysign logb 的函数接口
修复需要用户显式添加reserve_r29_for_scavenge属性
如果算子编译的时候不添加fgcu参数，并尝试去收集dma操作时，pti会抛出 the activities missed meta and not find meta for this sip launch debug message
尝试profile sdte/cdte时如果未增加-fgcu参数，会有assert sdma failed call trace
i20上批量执行op+cpu profile用例会抛出unable to handle kernel paging request at 0000000000002559
gcu-utilzation 在kernel_profile=1 的情况下，收集不到数据
TopsGDB attach/ -p pid 形式调试hang的程序，无法查看到gcu线程
已知问题¶
TopsCC kernel函数不支持纯虚类子类的实例化
Launch TopsCC kernel时，线程总数（grid dim * block dim）不能大于1024
topscc使用-O0编译报错
循环中更新leaptr的offset， 出现tar spill happen
kernel使用两个private DTE比只使用一个带来性能下降
cmake不能够使用topscc作为cc（c语言编译器）
S60/S60G不支持local关键字修饰的动态数组
ctrl+c或者SIP超时后，这之前的printf信息不能输出到控制台
当使用-O0对topscc kernel进行编译时，执行该kernel的操作会默认为同步操作
打开staging memcpy功能时，如果EFRT_STAGING_BUFFER_SIZE设置很小（如4KB），在直接使用系统接口malloc()分配的内存进行H2D或D2H操作时可能看到比较大的性能下降。topsHostMalloc()分配的内存，或进行过topsHostRegister()注册的内存则不会受到影响
attach 方式做kernel profile会引起rpc链接失败，从而导致attach无法进行数据收集
topsdnn 进行legacy profiler 收集不到 gcu op 的数据
S60先进行 legacy 只收集op 的profiler，再进行收集全部数据的kernel profiler ,kernel profielr 收集不到sdte的数据
TopsPlatform 0.9.0¶
TopsPlatform_v0.9.0版本首次在S6/S60上使能TopsGDB和TopsProfiler，TopsCC及TopsRuntime提供了更丰富的功能
新增¶
支持Staging Memcpy功能
TopsCC支持软流水
TopsCC支持Barrier同步方式
TopsCC在S6/S60上支持Block超发
TopsRuntime增加打印启动的kernel信息
TopsRuntime增加TopsMemGetInfoExt API以提供每个MC上的内存使用信息
TopsGDB在S6/S60设备上支持C/C++源码级别的调试，包括断点、单步执行、backtrace、查看栈变量和device变量、exception/abort信号捕获
TopsGDB支持多卡环境
TopsProfiler在S6/S60上支持单卡单进程
IDE提供configuration的TopsGDB选项，供用户选择使用
efsmi增加强制reset某张卡
变更¶
移除TopsCC的elemwise、reduction、nn计算接口
修复¶
将设备端函数代码直接写在调用者（caller）的内部和采用函数内联（forceinline）这两种方式生成的代码大小不一致
printf以格式转换说明符结尾时，会多输出(null)
printf(“%%”)多打印了一个%
kernel 函数中最多定义 8 个 constexpr，超过会出现运行时错误
某些intrinsic函数参数需要强制转换成addrespace(5)才能编译
i20上启用TOPS_ENABLE_DTE_CHECK宏之后，DTE一些错误行为会导致程序hang
使用char*类型的指令报错
内层循环中使用printf+set_address之后，qa的targ地址发生变化，导致store出现错误
在同一块memory中，使用循环load和store qacc， 出现core dump，dacc store 地址更新出错
LLVM ERROR: unsupported swap.smr spill
Kernel输入参数为nullptr，应用程序异常退出
topscc中set address store结果错误
Fix 无法从const section进行memcpy h2d的问题
trace-api 在 kernel_profile=1 attach 模式下，收集不到数据
收集不到ODMA event的数据
Profiler算子kernel气泡大
Factor profiler收集OP meta data时没有shape信息
已知问题¶
TopsCC kernel函数不支持纯虚类子类的实例化
Launch TopsCC kernel时，线程总数（grid dim * block dim）不能大于1024
Topscc使用-O0编译报错
kernel使用两个private DTE比只使用一个带来性能下降
cmake不能够使用topscc作为cc（c语言编译器）
循环中更新leaptr的offset， 出现tar spill happen
ctrl+c或者SIP超时后，这之前的printf信息不能输出到控制台
S6/S60不支持local关键字修饰的动态数组
当使用-O0对topscc kernel进行编译时，执行该kernel的操作会默认为同步操作
S6/S60上不支持SIP上的kernel程序运行时的强杀
topsEventRecord会造成host memory的泄露
S6/S60上使用一个event来同步多个stream时无效，可能造成硬件错误
对4G以上的host memory进行topsMemset、topsMemcpy等操作时，可能造成DTE Error
打开staging memcpy功能时，如果EFRT_STAGING_BUFFER_SIZE设置很小（如4KB），在直接使用系统接口malloc()分配的内存进行H2D或D2H操作时可能看到比较大的性能下降。topsHostMalloc()分配的内存，或进行过topsHostRegister()注册的内存则不会受到影响
如果算子编译的时候不添加fgcu参数，并尝试去收集dma操作时，pti会抛出 the activities missed meta and not find meta for this sip launch debug message
尝试profile sdte/cdte时如果未增加-fgcu参数，会有assert sdma failed call trace
i20上批量执行op+cpu profile用例会抛出unable to handle kernel paging request at 0000000000002559
KMD安装时需增加环境变量kernel_profile=1，才能使能topscc profiler
gcu-utilzation 在kernel_profile=1 的情况下，收集不到数据
TopsGDB attach/ -p pid 形式调试hang的程序，无法查看到gcu线程
TopsPlatform 0.8.1¶
TopsPlatform v0.8.1版本包含了如下新功能：TopsCC在S6、S60设备上使能、TopsGDB在i20上支持C/C++源码级别调试、TopsProfiler实现TopsCC-Profiler和Factor-Profiler的整合
新增¶
TopsCC在S6、S60设备上使能
支持配置dma的offset/layout参数的数组初始化，如： config_deslice(hbm_output, csb_input, {0, 0, 0, 0})
新增Runtime API：
topsExtSetProfileMeta for setting profile meta data
topsExecutableGetConstManagedData for getting weight to be refit
topsExecutableGetConstManagedDataV2 for getting weight to be refit
topsExecutableUpdateRuntimeResource for weight in-place replacement
topsExecutableLoadConstData for sharing constant in different executable
topsExecutableLoadConstDataV2 for sharing constant in different executable
topsExecutableQueryInputName for get executable input name
topsExecutableQueryOutputName for get executable output name
topsExtMallocWithBankV2 for expand topsExtMallocWithBank with flags parameter
topsExtMallocWithAffinity for malloc dev memory with logical memory bank
topsScatterInplace for replace scatter old sub mem with new sub mem descriptor
topsDeviceEnablePeerAccessRegion for setup p2p access mapping with peer device’s specified address region
topsDeviceDisablePeerAccessRegion for destroy p2p access mapping with peer device’s specified address region
TopsGDB在i20上支持C/C++源码级别的调试，包括断点、单步执行、backtrace、查看栈变量和device变量
TopsProfiler支持GCU指定事件和活动的采集
TopsProfiler支持CPU堆栈和内存等信息采集
TopsProfiler输出各种采集数据信息
TopsProfiler控制采集程序的启停
变更¶
TopsCC v0.8.1存在默认安装路径及库名字的改变，默认安装路径从/opt/topscc变为/opt/tops，请在使用时做相应修改
TopsGDB默认安装路径从/opt/topscc变为/opt/tops，请在使用时做对应修改
Runtime3.0库的默认安装路径从/usr变为/opt/tops，请在使用时做相应修改
Runtime2.0中头文件路径含dtu字样改成gcu，为兼容原用法，目前保留dtu软链接。软链接将在下次发布中去除，请在该段时间内完成相应修改
修复¶
在__device__ __forceinline__函数内L1数值错乱
编译报错：Incomplete scavenging after 2nd pass
matmul精度问题
topscc编译出来的二进制size较大，比factor编译的code size大3~4倍
二维数组使用vstore编译报错如下：Incomplete scavenging after 2nd pass
使用4个cdma，6个sdma，然后调用printf，hang
1c12s同时d2c，只启动了2个cdma，预期启动4个cdma
支持使用数组初始化列表来配置dmaoffset/layout参数
TopsCC 在 scorpio 上运行，cfunc kernel intrinsic 指令行为与预期不符
TopsCC反汇编intrinsic指令与预期不符，没有打包。需要加attribute
在T20、T21上TopsGDB通过Ctrl+D强制结束后再使用存在异常
已知问题¶
kernel c++不能支持纯虚成员函数
TopsCC在S60上只支持单Block
TopsCC在S6、S60设备上不支持libtopscxx提供的计算接口
kernel使用两个private DTE比只使用一个带来性能下降
cmake不能够使用topscc作为cc（c语言编译器）
ctrl+c或者SIP超时后，这之前的printf信息不能输出到控制台
将设备端函数代码直接写在调用者（caller）的内部和采用函数内联（forceinline）这两种方式生成的代码大小不一致
printf以格式转换说明符结尾时，会多输出(null)
printf(“%%”)多打印了一个%
kernel 函数中最多定义 8 个 constexpr，超过会出现运行时错误
某些intrinsic函数参数需要强制转换成addrespace(5)才能编译
[gcu300]不支持local关键字修饰的动态数组
[gcu200]启用TOPS_ENABLE_DTE_CHECK宏之后，DTE一些错误行为会导致程序hang
Build topscc kernel with -O0 flag and the kernel launch will be synchronized by default
TopsGDB仅支持单卡环境，暂不支持多卡环境
TopsGDB在i20设备上支持C/C++源码调试；在T20、T21设备上支持汇编级别的调试，对于S6、S60等设备的支持以及T20、T21设备上源码调试功能将在后续版本支持
必须增加环境变量kernel_profile=1，才能使能topscc profiler
gcu-utilzation 在kernel_profile=1 的情况下，收集不到数据
trace-api 在 kernel_profile=1 attach 模式下，收集不到数据
收集不到ODMA event的数据
factor profiler收集OP meta data时没有shape信息，Topscc profiler不支持获取OP meta data的shape信息