使用Tensor Core

最近在使用NNFusion的时候发现Codegen出来的FP16的网络在V100上的性能打不过FP32(甚至要慢一倍以上),但是理论上FP16应该要比FP32有两倍的性能收益才对(V100 Cuda Core的half precision的最大吞吐量是single的两倍,在s9234的slides中看到直接使用half的情况下peak performance其实和single差不多,都是15Tops,但是Cuda Core提供了half2类型,可一次做两个half类型的运算,这是half在CUDA Core上的收益来源;V100卡上的Tensor Core只支持FP16,利用好Tensor Core可以获得非常强的加速,A100卡上的Tensor Core增加更多的精度支持)。

建议阅读的文章:

聊聊 GPU 峰值计算能力

A100 Tensor Float 32 性能实测

拿nvprof测试了一下发现主要的性能瓶颈是:half卷积算子的实现速度要比single慢一倍,而这部分运算又占了总体运行时间的绝大部分。

CUDNN的卷积性能

NNFusion直接codegen出来的卷积算子默认使用CUDNN的实现,是CUDNN本身对卷积算子的时间不行,还是没有用好CUDNN呢?

注意到一篇2019年的论文,作者在V100上做了CUDNN_V7,CUDA_v9的卷积性能分析:Performance Evaluation of cuDNN Convolution Algorithms on NVIDIA Volta GPUs

这篇论文的结论是:

the filter size and the number of inputs are the most significant parameters when selecting a GPU convolution algorithm for 32-bit FP data. For 16-bit FP, leveraging specialized arithmetic units (NVIDIA Tensor Cores) is key to obtain the best performance.

在为32位FP数据选择GPU的卷积算法时,filter的大小和输入个数是最重要的参数。对于FP16运算,利用专门的算术单元(NVIDIA Tensor Cores)是获得最佳性能的关键。

CUDNN_V7的卷积使用Tensor Core的方法官方有给blog,必须得对运算做专门的限制,比如需要指定特定的算法,kernel和input的通道数需要跟8对齐等,我在自己的实验环境里用V8,也有相似的限制(但是没翻到V8用Tensor Core的文档)。

测试你的应用程序是否使用了Tensor Core

猜测是CUDNN没有使用到Tensor Core,于是开始分析应用程序的Tensor Core使用率,这里我用了两种分析方法,一种是用nvprof,另一种是用dlprof。

使用nvprof

nvprof的tensor_precision_fu_utilization这个metrics可以看到tensorcore的使用信息,可以参考Using Nsight Compute or Nvprof to Show Mixed Precision Use in Deep Learning Models,他会吐出不同kernel对tensorcore的利用率,分为0-10十一个等级,如果是0就代表没有使用tensorcore。

宿主机上使用:

sudo nvprof  --metrics tensor_precision_fu_utilization ./main_test

sudo是因为想要获取tensor_precision_fu_utilization的metrics需要用到底层驱动的API,普通用户没有权限获取,会报如下的错误:

==616== Warning: The user does not have permission to profile on the target device. See the following link for instructions to enable permissions and get more information: https://developer.nvidia.com/NVSOLN1000
==616== Warning: Some profiling data are not recorded. Make sure cudaProfilerStop() or cuProfilerStop() is called before application exit to flush profile data.

Nvidia官方提供的解决办法在:https://developer.nvidia.com/NVSOLN1000 ,需要先卸载驱动再在挂载驱动的过程中加个flag,但是我测试下来发现直接用sudo就可以。

容器内使用:

需要在起docker容器的时候给容器权限,否则无法使用nvprof拿到tensor core metrics: https://codeyarns.com/tech/2019-05-08-nvprof-in-nvidia-docker-permissions-warning.html

nvprof  --metrics tensor_precision_fu_utilization ./main_test

测试一个没有使用tensor core的cudnn单层卷积:

Invocations                               Metric Name                           Metric Description         Min         Max         Avg
Device "Tesla V100-PCIE-16GB (0)"
    Kernel: void implicit_convolve_sgemm<float, float, int=128, int=5, int=5, int=3, int=3, int=3, int=1, bool=0, bool=1, bool=1>(int, int, int, float const *, int, float*, float const *, kernel_conv_params, __int64, int, float, float, int, float const *, float const *, bool, int, int)
        100           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
    Kernel: void implicit_convolve_hhgemm<__half, int=0, int=6, int=6, int=5, int=4, int=4, bool=0, int=1, bool=1>(int, int, int, __half const *, int, __half*, __half const *, kernel_conv_params, __int64, int, float, float, int, __half const *, __half const *, bool, int, int)
        108           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)

这里的Idle为0代表没有使用tensor core。

测试一个使用了tensor core的cudnn单层卷积:

Invocations                               Metric Name                           Metric Description         Min         Max         Avg
Device "Tesla V100-PCIE-16GB (0)"
    Kernel: cask_cudnn::computeOffsetsKernel(cask_cudnn::ComputeOffsetsParams)
        100           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
    Kernel: volta_scudnn_128x32_relu_medium_nn_v1
        100           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
    Kernel: void implicit_convolve_hhgemm<__half, int=0, int=6, int=6, int=5, int=4, int=4, bool=0, int=1, bool=1>(int, int, int, __half const *, int, __half*, __half const *, kernel_conv_params, __int64, int, float, float, int, __half const *, __half const *, bool, int, int)
          3           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
    Kernel: void cutlass::Kernel<cutlass_tensorop_h884fprop_analytic_64x64_32x2>(cutlass_tensorop_h884fprop_analytic_64x64_32x2Params)
        106           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization     Low (1)     Low (2)     Low (1)
    Kernel: void nchwToNhwcKernel<__half, __half, float, bool=1, bool=0, cudnnKernelDataType_t=0>(int, int, int, int, __half const *, __half*, float, float)
        106           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)
    Kernel: void nhwcToNchwKernel<__half, __half, float, bool=1, bool=0, cudnnKernelDataType_t=0>(int, int, int, int, __half const *, __half*, float, float)
        106           tensor_precision_fu_utilization   Tensor-Precision Function Unit Utilization    Idle (0)    Idle (0)    Idle (0)

关于如何在cudnn的卷积中使用tensor core,可以参考这一篇blog

分析nvprof的日志可以看到cutlass_tensorop_h884fprop_analytic_64x64_32x2这个kernel使用了tensor core,但是使用率也很低,是Low(1)。

查看tensorrt过程中

使用dlprof

实际测试的过程中发现使用nvprof测tensor core metrics的时候会影响程序运行的时间,例如原来的一次推理只需要运行6ms,加上了tensor_precision_fu_utilization 这个metrics之后,就变成了200ms左右才能完成一次推理,而为了测benchmark,我使用了5次warm_up和100次loop,时间还是很长的,而且nvprof拿到的信息也比较少,相比之下dlprof拿到的信息更加全面,不会影响程序的运行时间,输出也是的也是csv格式,更胜一步的是可以用tensorboard打开,这样方便在没有图形界面的情况下也能用网页端可视化查看资源:

dlprof --force=true --reports=all --mode=simple ./testCudNNConv

执行该命令之后,会在当下前目录下生成若干文件:

-rw-r--r-- 1 root root       622 May 31 15:48 dlprof_detailed.csv
-rw-r--r-- 1 root root        47 May 31 15:48 dlprof_expert_systems.csv
-rw-r--r-- 1 root root       525 May 31 15:48 dlprof_group_node.csv
-rw-r--r-- 1 root root     43709 May 31 15:48 dlprof_iteration.csv
-rw-r--r-- 1 root root      2234 May 31 15:48 dlprof_kernel.csv
-rw-r--r-- 1 root root       501 May 31 15:48 dlprof_op_type.csv
-rw-r--r-- 1 root root       940 May 31 15:48 dlprof_summary.csv
-rw-r--r-- 1 root root       217 May 31 15:48 dlprof_tensor_core.csv
-rw-r--r-- 1 root root    262357 May 31 15:48 nsys_profile.qdrep
-rw-r--r-- 1 root root    950272 May 31 15:48 nsys_profile.sqlite

其中summary可以看到总情况,tensor_core可以看到使用和没有使用tensor_core的kernel数量。

同样来看在没有使用tensorcore的单层卷积上,tensorcore的使用情况:

cat dlprof_tensor_core.csv
Tensor Core Report
Op ID,Op Name,Op Type,Uses TC,Total GPU Time (ns),TC GPU Time (ns),Non-TC GPU Time (ns),TC Utilization (%),Total Kernel Count,TC Kernel Count,TC Kernel Names,Non-TC Kernel Count,Non-TC Kernel Names

再看使用tensorcore的单层卷积上,tensorcore的使用情况:

cat dlprof_tensor_core.csv
Tensor Core Report
Op ID,Op Name,Op Type,Uses TC,Total GPU Time (ns),TC GPU Time (ns),Non-TC GPU Time (ns),TC Utilization (%),Total Kernel Count,TC Kernel Count,TC Kernel Names,Non-TC Kernel Count,Non-TC Kernel Names
"SIMPLE_MODE_GLOBAL_NODE_OP_TYPE_1","simple_mode_global_node_name",simple_mode_global_node_op_type,yes,6583335,12063,6571272,0.2,17,1,"Kernel",16,"computeOffsetsKernel, explicit_convolve_sgemm, fft1d_c2r_32, fft1d_r2c_32, fft2d_c2r_64x64, fft2d_r2c_64x64, flip_filter, gemv2T_kernel_val, im2col4d_kernel, implicit_convolve_hhgemm, implicit_convolve_sgemm, nchwToNhwcKernel, nhwcToNchwKernel, transpose_readWrite_alignment_kernel, volta_gcgemm_32x32_nt, volta_scudnn_128x32_relu_medium_nn_v1"

只有一个kernel使用了tensor core,和nvprof的报告一致,但是kernel name却变成了kernel,没有nvprof的清晰..

但是这个工具测试tensorrt等一些框架就能得到很好的结果,例如测试一下tensorrt的resnet50,tensor core的利用率都打满了:

image-20220601000732050

未完待续..

Windows Server AMD 玄学驱动安装 踩坑nvidia driver

Comments

Your browser is out-of-date!

Update your browser to view this website correctly. Update my browser now

×