版本须知¶
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信息