The performance of ResNet50 using CUDNN and CUBLAS is significantly slower compared to Cutlass

umiswing · May 12, 2023, 11:10am

Hello! I am currently testing the ResNet50 model on NVIDIA’s 40G A100 platform using https://github.com/masahi/tvm-cutlass-eval. After utilizing Cutlass, the compiled ResNet50 is 2.9 times faster compared to compiling it with CUDNN and CUBLAS. However, I find it strange that Cutlass can achieve such a significant speed-up. Is there anything I am overlooking? Below are the specifics of my testing process:

In order to run the resnet50/run.py using TVM’s main branch API, I made some modifications to the testing scripts:Update to run on main branch of tvm. · umiswing/tvm-cutlass-eval@3b4bd37 · GitHub

Link to reproduce the testing (python run.py):https://github.com/umiswing/tvm-cutlass-eval/blob/master/resnet50/run.py

Performance using cutlass:

Execution time summary:
  mean (ms)   median (ms)    max (ms)     min (ms)     std (ms)
    1.7029       1.5734       1.8596       1.5677       0.1418

Performance using cudnn+cublas:

Execution time summary:
  mean (ms)   median (ms)    max (ms)     min (ms)     std (ms)
    4.7330       4.6858       5.0248       4.6459       0.0859

cuda version and hardware

11.7

NVIDIA A100 40G

tvm version

commit 1d145f112115ca20a0cd2e37a726b1d1519cac4b

config.cmake

@@ -46,7 +46,7 @@
 # - ON: enable CUDA with cmake's auto search
 # - OFF: disable CUDA
 # - /path/to/cuda: use specific path to cuda toolkit
-set(USE_CUDA OFF)
+set(USE_CUDA ON)

 # Whether enable ROCM runtime
 #
@@ -142,7 +142,7 @@ set(USE_MICRO_STANDALONE_RUNTIME OFF)
 # - OFF: disable llvm, note this will disable CPU codegen
 #        which is needed for most cases
 # - /path/to/llvm-config: enable specific LLVM when multiple llvm-dev is available.
-set(USE_LLVM OFF)
+set(USE_LLVM /usr/bin/llvm-config-11)

 #---------------------------------------------
 # Contrib libraries
@@ -217,10 +217,10 @@ set(USE_EDGETPU OFF)
 # - ON: enable cuDNN with cmake's auto search in CUDA directory
 # - OFF: disable cuDNN
 # - /path/to/cudnn: use specific path to cuDNN path
-set(USE_CUDNN OFF)
+set(USE_CUDNN ON)

 # Whether use cuBLAS
-set(USE_CUBLAS OFF)
+set(USE_CUBLAS ON)

 # Whether use MIOpen
 set(USE_MIOPEN OFF)
@@ -416,7 +416,7 @@ set(USE_GTEST AUTO)

 # Enable using CUTLASS as a BYOC backend
 # Need to have USE_CUDA=ON
-set(USE_CUTLASS OFF)
+set(USE_CUTLASS ON)

masahi · May 12, 2023, 11:20am

Indeed something seems off. Even on RTX 3070 I had a better result according to https://github.com/masahi/tvm-cutlass-eval. Have you tried nvprof?

umiswing · May 12, 2023, 11:56am

Hi @masahi . I just ran a profiling on an A100, and I noticed that the GPU kernel times were similar. However, when compiling the model with CuDNN and CuBLAS, I discovered that the compiled model is taking a lot of time on cudaFree() and cudaStreamCreateWithFlag(), which are synchronous operations. This might be the reason why ResNet50 compiled with CuDNN and CuBLAS is running slow.

Links to profiling results:

github.com

umiswing/tvm-cutlass-eval/blob/master/resnet50/cutlass_nsys.txt

Generating SQLite file report1.sqlite from report1.nsys-rep
Processing [report1.sqlite] with [/opt/nvidia/nsight-systems/2023.2.1/host-linux-x64/reports/nvtx_sum.py]... 

 ** NVTX Range Summary (nvtx_sum):

 Time (%)  Total Time (ns)  Instances    Avg (ns)      Med (ns)     Min (ns)    Max (ns)   StdDev (ns)   Style                Range              
 --------  ---------------  ---------  ------------  ------------  ----------  ----------  -----------  -------  --------------------------------
     92.5       10,537,391          1  10,537,391.0  10,537,391.0  10,537,391  10,537,391          0.0  PushPop  nvFuser::Manager::CudaFuseGraph 
      7.3          827,126        584       1,416.3         400.5         176      20,645      1,997.4  PushPop  isFusibleCudaFusionGroup        
      0.1           11,388          2       5,694.0       5,694.0       5,154       6,234        763.7  PushPop  compileFusionRecursive          
      0.1           10,311          1      10,311.0      10,311.0      10,311      10,311          0.0  PushPop  PeepholeOptimizeShapeExpressions

Processing [report1.sqlite] with [/opt/nvidia/nsight-systems/2023.2.1/host-linux-x64/reports/osrt_sum.py]... 

 ** OS Runtime Summary (osrt_sum):

 Time (%)  Total Time (ns)  Num Calls    Avg (ns)      Med (ns)     Min (ns)   Max (ns)    StdDev (ns)           Name         
 --------  ---------------  ---------  ------------  -------------  --------  -----------  ------------  ---------------------
     43.3   11,513,996,879        127  90,661,392.7  100,119,499.0     7,573  100,176,355  26,846,413.3  poll                 
     38.6   10,255,595,219      2,086   4,916,392.7       82,120.5       257  339,303,908  27,605,735.4  ioctl

This file has been truncated. show original

github.com

umiswing/tvm-cutlass-eval/blob/master/resnet50/cudnn_cublas_nsys.txt

Generating SQLite file report2.sqlite from report2.nsys-rep
Processing [report2.sqlite] with [/opt/nvidia/nsight-systems/2023.2.1/host-linux-x64/reports/nvtx_sum.py]... 

 ** NVTX Range Summary (nvtx_sum):

 Time (%)  Total Time (ns)  Instances    Avg (ns)      Med (ns)     Min (ns)    Max (ns)   StdDev (ns)   Style                Range              
 --------  ---------------  ---------  ------------  ------------  ----------  ----------  -----------  -------  --------------------------------
     92.0       10,549,026          1  10,549,026.0  10,549,026.0  10,549,026  10,549,026          0.0  PushPop  nvFuser::Manager::CudaFuseGraph 
      7.7          886,869        584       1,518.6         514.0         202      18,756      1,955.0  PushPop  isFusibleCudaFusionGroup        
      0.1           17,191          2       8,595.5       8,595.5       5,857      11,334      3,872.8  PushPop  compileFusionRecursive          
      0.1           10,186          1      10,186.0      10,186.0      10,186      10,186          0.0  PushPop  PeepholeOptimizeShapeExpressions

Processing [report2.sqlite] with [/opt/nvidia/nsight-systems/2023.2.1/host-linux-x64/reports/osrt_sum.py]... 

 ** OS Runtime Summary (osrt_sum):

 Time (%)  Total Time (ns)  Num Calls    Avg (ns)      Med (ns)     Min (ns)   Max (ns)    StdDev (ns)           Name         
 --------  ---------------  ---------  ------------  -------------  --------  -----------  ------------  ---------------------
     55.7   30,035,002,283        312  96,266,033.0  100,113,772.5    36,878  100,155,652  17,725,295.1  poll                 
     43.8   23,620,449,524      3,968   5,952,734.3       86,821.5       249  485,291,305  32,856,503.4  ioctl

This file has been truncated. show original

umiswing · May 14, 2023, 3:15pm

Hi @masahi. Is this a performance issue with TVM or an expected acceleration since cutlass provides more fusion patterns, thereby avoiding the allocation and freeing of intermediate memory?

masahi · May 14, 2023, 10:18pm

Not sure, but in general cuDNN shouldn’t be allocating a lot of intermediate memory. CUTLASS doesn’t allocate any temp storage for resnet at all.

umiswing · May 18, 2023, 6:51am

Hi @masahi. Some update for this issue:

According to the timeline, when TVM compiles ResNet50 with cuDNN, sum of kernel’s duration is similar with ResNet50 compiled with cutlass, but ResNet50 compiled with cuDNN seems spends a lot of time on waiting something when executing the kernel, while model ResNet50 compiled with cutlass does not. The waiting behavior is strange but it could be another reason why ResNet50 compiled with cuDNN is so slow.

timeline of ResNet50 compiled with cuDNN

timeline of ResNet50 compiled with cutlass

Does TVM plan to debug with this issue recently?

masahi · May 18, 2023, 7:16am

No, at least I don’t have a plan to work on cuDNN related issues. We are happy with cutlass.