国产加速器海光DCU&GPGPU深算处理器程序开发常见问题

目录

一、硬件架构参数问题

二、DCU编程环境问题

三、 转码移植基本问题

四、 性能调优、MPI通信相关问题

五、编译及开发相关问题

1)编译类

2 )Fortran类

3 )核函数类


以计算服务平台昆山中心为例,罗列一些常见的DCU开发常见问题,请各位参考。

一、硬件架构参数问题

1、计算平台CPU、DCU对标NV哪个型号?

A:计算平台CPU#1,对标AMD Naples。DCU对标NV
Tesla V100、AMD Vega MI60。

2、DCU与AMD GPU的什么架构相似?

A:可参考AMD gfx906。

3、DCU(风冷及液冷)与NVIDIA V100在硬件架构与配置上有何差异?

A:DCU没有tensorcore,纹理单元和内存也没有,share memory机制类似,等基本是一样的,容量有些差别。

4、DCU支持tensorcore吗?

A:不支持。

5、计算平台上的cpu存在numa现象吗?与Intel至强CPU在硬件架构与配置上有何差异?

A:有的。每个cpu有4个numa节点,相当于4路cpu,每个cpu带一个加速卡,调优时注意绑定numa。Intel至强cpu 每个cpu有一个numa节点。

6、DCU系统软硬件架构

A:登录 - 光合开发者社区

7、DCU Z100产品参数

A:请联系秘书处获取产品手册,光合基金秘书处邮箱 ghfund@hieco.com.cn

二、DCU编程环境问题

1、DCU是否支持切换精度的宏?

A:支持。

2、是否支持boost库cpu端使用?

A:支持。

3、DCU是否支持thrust库?

A:支持,DCU中是rocThrust。

4、对于二级矩阵中的doublepoint**,CUDA里面是分配页锁定内存,DCU是否也支持这种用法?

A:支持锁页内存(pinned memory)。

5、DCU是否支持CUDACXX库?

A:目前还不支持,libcudacxx还没有做,load store接口是有的,但没有开放出来,可以联系开发工程师进行封装,用户订制开发。建议尽量绕过去CUDACXX。

6、ROCm对CUDA的支持列表有哪些?

A:CUDA4.0的大部分功能支持,如stream、内存管理、同步管理、kernel都差不多支,不支持的有:global调global(dinamic
parallel)、协作组。

7、目前昆山集群有哪些编译好的数学库可用呢?

A:rocBLAS、rocFFT、MIOpen等

8、能否使用gcc 10/11中的OpenMP/OpenACC做DCU的异构开发?

A:DTK-22.10已经添加了OpenMP和OpenACC C和C++语言支持,使用手册请见 https://cancon.hpccube.com:65024/1/main/DTK-22.10/Document

9、Unified Memory支持如何?

A:有api,更多是功能的封装,性能不好。

10、是否支持HPCG?

A:支持,性能还可以。

11、是否支持kokos?

A:后续发布的版本会支持,目前虽然编译通了,但性能不是太高。

12、是否支持PETSc?

A:目前PETSc对异构支持有HIP和Kokkos两种方式,支持的类型有限,支持vector相关的部分计算,matrix相关计算还在开发中,整体还不太可用,建议应用针对性的做异构移植和优化。

13、有无Clang环境?

A:直接用ROCm环境即可,ROCm环境就是基于LLVM+Clang环境构建的。

14、设备端数据结构是否可以用vector?

A:建议使用thrust cube等,不要用stl。

15、DCU芯片有没有CUDA的GPU direct和P2P类似的功能?

A:不支持,下一代有规划。

三、 转码移植基本问题

1、平台上如何进行cuda到hip的转码?

A:加载rocm环境可以直接使用hipify-perl命令,或者使用 hipconvertinplace-perl.sh 脚本对整个目录进行转码。

2、移植流程是怎样的?

A:1)profiling分析程序,找出耗时较多的部分——用时约60%-70%;2)移植,尽量减少cpu和dcu之间的数据搬迁,主要部分移植到dcu,较难移植的部分由cpu辅助去做。

3、核函数转换后字符串支持不太好,怎么办?

A:转换工具是辅助工具,转换后性能有差异,参考hip的一些开源项目。

四、 性能调优、MPI通信相关问题

1、DCU计算能力跟V100的区别?

A:50-70%。访存带宽差别不太大,影响约10-20%,根据数据访存特征,cache规格不一样。

2、使用pytorch写的程序需要做深度优化吗?

A:需要具体分析,如果是自定义算子需要自己深度优化,如果是通用算子如常用卷积计算我们则会提供相关优化。

3、DCU性能调试方法和工具有哪些?

A:dcuprof、roctracer、火焰图,具体参考开发者社区文档部分。

4、平台是否有dcuprof针对hip的性能检测工具?

A:有的,dtk22.04及以后版本有hipprof,具体使用方法请见开发者社区文档 登录 - 光合开发者社区

5、DCU性能分析,指令rocprof --hip-trace -o,为什么用chrome打开看不到具体时间?

A:添加-w扩展可以看到。

6、DCU并行计算有没有一些成功的性能调优经验?

A:可以复用V100的调优经验,注意硬件架构的区别。

7、计算平台CPU并行计算有没有一些成功的性能调优经验?

A:复用Intel调优经验。

8、为什么常量内存和纹理内存在hip api没有明显的性能优化?

A:DCU没有常量和纹理内存的硬件,是通过API兼容将这块空间映射到global上,所以和直接Global访问没啥差别

9、全局通信 all to all哪种比较好?hpcx吗?

A:建议使用Intel mpi,hpcx需要打开sharp功能(硬件),也看总进程数,规模不大时效果不明显,需要软硬件结合调参数。

10、提交任务失败,报错mpi问题怎么解决?

A:考虑可能原因是编译使用的mpi和运行mpi不一致,不兼容导致报错。

11、为什么运行程序时,单节点2个cpu慢于两个节点各一个cpu?

A:存在这样的可能性,单节点两个cpu的通信可能没有多节点通信快。

12、关于MPI并行IO,有没有相关包可以并行文件存储?

A:不建议直接使用MPI IO功能,IO问题需要具体问题具体分析,并行文件读写涉及到锁和对齐问题,是个相对比较复杂的问题。

13、怎么实现DCU节点负载均衡?

A:对于同化应用来说,每个进程上的任务是独立的,则可以单独设置主进程控制任务分发来做任务级并行。

14、如何配置I_MPI_FABRICS和I_MPI_PMI_LIBARARY?

A:前者是指定IntelMPI使用哪种底层通讯接口如OFI、DAPL还是UCX等;后者是大规模任务启动时IntelMPI和SLURM配合的库,程序能跑的话不用设置这个参数。

15、矩阵1000万元素个数的二维矩阵,在DCU上有多大的加速效果?

A:如果是矩阵GEMM或其他BLAS标准计算,直接使用BLAS库即可,加速比会非常高。

16、rocm是否支持支持动态并行?

A:目前不支持,rocm4.5之前不能支持,架构和nv不一样,占用很多卡资源,建议多线程并行。

17、BOOST SIGNALS2可否不拆分?

A:可以不拆分,需要检查参数类型是否正确,建议使用高版本rocm。

18、报错illegal instruction detested如何解决?

A:暂时屏蔽test模块。

19、运行程序时,DCU平均分配线程还是一块卡用满再用下一块?

A:自己可以通过程序控制。

20、使用DCU单卡并行2个以上的任务,怎么用stream?

A:计算之间没有关联,异步发送任务,stream多个数据,参考cuda c编程指南,控制十几个stream以内。

21、GPU的内建函数在DCU是否有相应的函数对应?是否同样将cuda改成hip?

A:是的,接口都是相似的,还有一些如push pop都是一样的。

22、是否可以提供简单的调blas的示例程序?(对比cuda的示例,以便快速上手)

A:跟CUDA类似,可先参考CUDA的编程手册。

23、遇到fft的相关问题,怎么解决呢?

A:ROCm发布的FFT会有正确性问题,我们发布的DTK解决了不少,可以试用,如果还有问题可反馈到开发者社区论坛。

24、DCU有没有向量化的操作?

A:有的,1个cu 64线程,包括4个simd,以wavefront为单位调用到线程

25、block分支多线程如何运行,需要原子操作吗?

A:block内分支执行流程,如果有if
else语句,若干个线程执行if,结束后再由若干线程去执行else。1个warp内按顺序执行,所有线程在执行同样的指令,只不过有些被mask掉不执行,if和else不会同时执行。如果线程数超过warp size 64,会进行重排序。

26、原子操作对性能影响多大?除了影响cache的有效利用外,对block内的线程运行是怎样影响的?

A:在当前硬件还不支持float的atomic操作,当前API是通过CAS的锁来实现的,性能不高。

27、显存中存在一个int变量,经过核函数的一系列执行操作,该变量的值会只是0或1吗?还是会产生不可预期的脏数据?

A:只会是0或1,不会有脏数据。

28、10万个int数组,不到M级别,数据传递开销占比很大,有什么优化技巧吗?

A:数据访问做合并,是否有数据和访存重叠,锁页内存、异步传输都可以用,没有数据预取,且数组较小,对优化没有影响。

29、share memory比较小,怎么使用比较合适?

A:看算法数据的重复性,数据共享时share memory效果较好,数据各不一样时不建议使用。

30、内存之间是连续的,并不是线程是连续的,对吧?

A:所有线程拼起来的地址是连续的效果最好。

31、线程提上去之后,需要很大的share memory怎么办?

A:share memory不是缓存,需要手动编程来控制这块存储空间,主要用来存放一个workgroup重复使用的数据。

32、并行计算优化两个重点

A:1)尽量将并行的合到一起,申请一个大的空间,一次传比多次传快;2)内核计算少的话也尽量合并在一起,多利用异步

33、DCU平台是否有性能计数器、profiling工具、调试工具

A:性能分析工具:hipprof,调试工具:dtk22.10版本会有rocgdb。

34、使用hipprof找到热点函数,但感觉粒度太粗,还需要找到函数里面的热点行,是否有合适的工具?

A:暂时没有,类似nvsight的工具短期内无法实现。

35、是否有Debug工具可以把原始的Binary定位到源代码?

A:没有,建议把大kernel拆成小一点的。减少数据拷贝(全局内存的读写)。按照NV的优化经验进行优化即可。

36、使用生成kernel,程序性能只到V100的五分之一,如何优化?

A:编译器处理能力有限,kernel越大的话处理越差。

37、程序中有多级嵌套和封装,大数组输入,约1G,如何优化?

A:1)建议将大结构体拆开,一个一个传参进去;2)也可以从数组的结构体和结构体的数组中优化一下,否则性能急剧下降;3)需要调整整个模块的参数空间,使得标量的、单double的、threshold放到一起,尽量避免一级矩阵、二级矩阵,地址类的单独传。

38、maple2C自动生成,代码可读性差,低效、每个线程多达几百个寄存器,编译器是否会优化?

A:dtk22.10增加了对local
memory的解析,可以在命令行里看,另外增加了可视化的方式,把程序下载下来后导入工具就可以看到。以上这些还没有汇总到hipprof中。但是local memory优化空间相对比较有限,建议从算法层面优化,比如数据分块、程序执行流程等过程。

39、跟其它老师沟通,写量化程序时,数据传输可以有不同方式,可以整块儿传进去,不一定要切分?

A:这种情况猜测程序之中还有一些计算,并不需要所有的数据,先将部分数据拷贝进去计算,计算的同时拷贝剩余数据。

40、之前培训视频中提到可以用指令延迟的方式提升性能,能够提升多少?

A:建议程序稳定、能够迭代多次后再使用指令延迟的方式,否则会导致后续修改程序时出现错误,涉及到非常细的语句排布。

41、数据传输耗时远大于计算耗时,如何优化?

A:可以考虑将相同区间内的数据进行合并,对齐合并后传输。更多的侧重算法层面的提升,比如数据分块、程序执行流程等过程。

42、一般share memory怎么使用?

A:一个CU里面或一个block里面,共用的数据放在shared memory。大数组放在local memory。

43、在DCU上统计时间时,用get time of day多次跑的时间差异较大。

A:get time of day是CPU端的,是异步的,计时前面一定要加device的同步,请注意三点:(1)kernel前后都要加同步;(2)同步有一定的开销;(3)第一遍kernel会比较慢,作为warmup。建议使用event时间,精度是高的。

44、调hipblas等库时时间会比较稳定,而自己写的kernel会有时间问题,如何处理?

A:确定kernel每次计算都是一样的,kernel直接循环运行,而不是分开提交作业运行。

45、1个节点多块卡都用上的话效率会非常低,1个节点1块卡就没问题。

A:完整独占节点,需确认不是所有进程都跑到一张卡了。用hipsetdevice设置卡,写在所有分配显存前面。

46、DCU负载为0%怎么办?

A:可以调一个库看下,如矩阵乘,使用watch rocm-smi监控。

47、传输时间对大体系计算有没有影响?

A:尽量减少cpu和dcu之间的数据传输,数据尽量全都放显存上,或用stream做掩盖。

48、报错GPU segmentation fault的原因?

A:数据访问地址越界,需定位到具体哪个内核函数,可以使用export HIP_KERNEL_PRINTF=1打印kernel函数,若不加此指令,可能会报错invalid address access:(nil),Error code 1.已中止(核心
倾倒))

49、lamps中的TIP4P水分子模型可以用DCU加速吗?

A:lammps 水分子模型里面长程力不能GPU加速,gromacs可以。Lammps主要面向材料的力场。

50、lamps中执行loop操作时使用clear命令清空模拟参数配置,引发访问权限错误问题。loop第一次正常、第二次报错,不调用DCU时loop正常运行。

A:建议使用新版dtk,如果还有报错,请反馈到社区论坛。

51、实践中网格是二维的,启动计算时,如何进行网格定义?

A:网格定义跟实际计算时的网格一样。

52、可以对4维张量加速吗?

A:网格划分只有三维。可以降维,最终都是一维。更多的维度可以用for循环。

53、对于购买的DCU Z100计算卡,在NVIDIA计算卡能够成功运行的一些程序放到DCU Z100上完成HIP转换再运行时发生错误,报错信息是内存分配失败。

A:intel服务器上插DCU卡可能会有这个问题。升级一下驱动或统一软件版本。建议购买配套服务器和加速卡。

54、新版HIP可以支持cuQuantum吗?

A:没有支持,cuQuantum做量子计算、张量网络的。

55、DCU是否支持rdc?当前做法是把核函数写在.cpp中,编译成object .o文件,最后链接时把.o文件链接成可执行文件。

A:rdc功能暂时不支持。当前建议您把printValue函数写在.h文件中,include进来。但是这样做有弊端:global变量不太能做了。根本解决方法:使用llvm相关的指令把cpp文件编译成bc的形式,再用mlink手动链接起来,需要对llvm比较熟。

56、hipify工具转码时,launchkernel如何转换?

A:hip和cuda的launchkernel基本是一样的,基本不用怎么转换。

57、有没有类似张量收缩TensorNet library的库

A:不建议使用。它是为rocblas服务的,比较底层,可以做高维和收缩,是生成其他库的工具,不是面向用户的库。

58、tensorflow框架训练,单DCU和双DCU时间接近,四DCU反而训练时间更久?

A:建议使用export HSA_FORCE_FINE_GRAIN_PCIE=1

五、编译及开发相关问题
1)编译类

1、作业方式提交编译与在编译窗口直接运行有什么异同?

A:登录节点只能用来编译程序,但由于没有加速卡,测试需要提交任务到计算节点进行,具体参考作业调度系统使用。

2、CMake中使用ENABLE_LANGUAGE命令报错

A:当前官方不支持这一命令,ENABLE_LANGUAGE是顶层语句,当前不支持开启,目前没办法写cmake扩展。

3、HIP和OpenCL可以混编吗?

A:可以混编,建议只使用HIP,比较高效。混编初始化、流管理比较麻烦,已验证功能性可以支持,但实际应用不能100%保证。

4、HIP和OpenMP可以混编吗?

A:HIP和OpenMP结合不是太好,OpenMP可以调用hip接口,但不建议写在一起,可以分开写,建议使用pthread做控制。

5、编译程序时导入其它机器环境的动态链接库,报错找不到函数名,如何解决?

A:分析动态链接库编译是否正确,函数导出是否跟现在的函数匹配,不同版本的编译器对函数命名稍有不同,建议当前节点再编译一遍动态链接库。

6、假设数据A在global函数里定义了值,若一直不调用该函数,后续它是一直存在于regester里面吗?

A:会被编译器优化掉,初始A存在于寄存器,但如果程序比较大,寄存器不够用,A可能会被移到globalspace。

7、数据先放到global,计算时必须送到寄存器上去,有些数据会用两三次,这些数据是一直存在于寄存器呢,还是每次用的时候再送到寄存器?

A:这是编译器的分配和优化工作,视程序/kernel大小而定,每线程最多使用256寄存器,空间不够时会使用global空间。

8、编译参数都有哪些?

A:新版本的dtk已内嵌一些参数,不需要显式指定。之前版本的dtk编译参数参考如下:

rocm2.9:
-L** (ROCM_PATH)/hip/lib -lstdc++ -L** (ROCM_PATH)/lib -lhsa-runtime64
-lhc_am

dtk21.04链接参数:-L** (ROCM_PATH)/hip/lib
-lamdhip64 -L** (ROCM_PATH)/lib -lhsa-runtime64,示例:cd
./TMP ;** (LD) ** (LINKFLAGS) -o ../** (目标) ** (对象) ** (库)
-L** (ROCM_PATH)/hip/lib -lamdhip64 -lstdc++ -L$(ROCM_PATH)/lib -lhsa-runtime64

9、编译时报错找不到文件怎么处理?

A:编译文件中指定路径。

10、hipSetDevice报错“hipErrorNoBinaryForGpu: Coudn't find binary for current devices!”

A:报错原因是编译时生成的目标机器二进制文件不匹配。需要添加-amdgpu-target=gfx906编译指令来指定gfx906。

11、Cmake选项太多了,HIP Language support功能没有加到搜索库中

A:enable-language在dtk22.04.2及以后的版本支持。

12、矢量化,AVX2或AVX512是否支持?

A:建议用SSE,指令比较标准。AVX在AMD和Intel之间的兼容性不太好,由于专利原因,指令集不一样。

13、HIP与OpenMP(CPU部分)共同编译有问题,如何处理?

A:dtk增加了OpenMP的原生支持(OpenMP和OpenMP target offload函数),建议使用hipcc,加参数-fopenmp。跟llvm的使用是一样的,把openmp库都内置进去了。建议使用OpenMP target offload——它会将这部分代码生成在DCU上执行的部分,直接使用DCU加速。多线程使用OpenMP的导语,类似OpenACC。跨节点使用OpenMPI,不建议使用DCU Direct(还不太行),RDMA不支持,现在有一个伪RDMA。DTK22.10版本会提供OpenACC的C和c++语言的支持。

14、自己台式机编译的openmp能否在DCU运行?

A:DCU支持openmp编程,hip5.0之后 dtk22.10支持openmp
offload。

15、CPU多线程转换到DCU openmp,有没有转换工具或使用方法?

A:加载dtk22.04,自带llvm编译器支持的openmp库,里面既有CPU多线程的,也有DCU offload的。加上openmp的编译参数。

2 )Fortran类

1、可否使用Flang编译器编译fortan程序遇到问题?

A:Flang编译器不成熟,不支持,CPU程序建议使用Intel的。

2、Fortran的OpenMP可以使用吗?

A:不建议使用,Flang还在快速开发中,到距离能用还有一段距离。

3、如何将基于Fortran纯CPU的应用调用DCU?

A:主程序使用原本框架,DCU加速部分使用动态库方式实现,链接至主程序。

4、怎样通过静态库或动态库实现Fortran调用c++?

A:建议 export导出函数接口,加下划线,调用时调这个接口,或者使用iso_binding功能。

3 )核函数类

1、核函数运行异常时整个程序都会退出,能否改进?

A:目前版本的设计理念是出现报错整个进程会退出,后续会优化为核函数报错时程序不退出。

2、关于核函数的返回值,是否可以有condifitonal
success的机制,能够在中间按条件退出?

A:目前dcu设计理念函数类型大体为void类型,不能从kernal返回值的,程序正常执行完成返回成功,错误时退出。可以申请一个global变量,在host端检查。

3、是否支持Fortran版的核函数?

A:不支持。虽然dcu核函数跟c程序写在一个文件里,但不是相互调用的,仅限于编译器本身的工作。Fortran只能调用cpu的程序,有cuda Fortran,但hip目前不能支持。

4、OpenMP可以制导到核函数吗?

A:可以,Fortran和c都可以,c最好,Fortran部分语法不支持。使用llvm的编译器,需要写OpenMP的offload
target。

5、kernel最大可以写到多大?

A:kernel大的话,计算资源会很大,启动的线程数受限。具体根据算法来定,核可以大,防止资源浪费,但不能过大,没有明确边界,需要平衡。

  • 13
    点赞
  • 10
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 0
    评论
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

打赏作者

技术瘾君子1573

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

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

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

打赏作者

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

抵扣说明:

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

余额充值