前言
在Nvidia的服务器GPU部署AI模型一定要使用TensorRT, 因为cuDNN远慢于TensorRT. 参考我最新的文章基于NvidiaGPU的AI模型结构优化:
最后说一点, 目前能够完美发挥出TensorCore性能的引擎, 据我所知只有TensorRT. 一些训练推理平台(比如pytorch)底层实际调用的是cuDNN来完成模型的推理. 但据我的经验来看, 相同参数的卷积cuDNN的实现要比TensorRT的实现慢3倍以上. 我也尝试过实现基于TensorCore的卷积, 但最终还是会比TensorRT慢20%多的样子.
————————————————
版权声明:本文为CSDN博主「Mr_L_Y」的原创文章,遵循CC 4.0 BY-SA版权协议,转载请附上原文出处链接及本声明。
原文链接:https://blog.csdn.net/luoyu510183/article/details/117385131
最近的项目是在帮助服务器落地一些AI算法, 用到的Nvidia的GPU来推理, 主要会用到CUDA, cuDNN和TensorRT. 另外我也在开发独立的CUDA引擎来替换cuDNN, 由于cuDNN目前不支持Direct这种直接卷积算法, 只能使用GEMM,FFT和WINOGRAD这三种算法. 我的目前测试感觉这三种并不会一定会比直接卷积快, 另外cuDNN的TRUE_HALF_CONFIG , 并不会比float计算快, 这个和我用直接卷积的结论是相反的.
先讲解下结构关系:
CUDA是Nvidia GPU开发的基础工具集, 不但包含了.cu编译器nvcc, 还有一大堆好用的工具库 比如npp, nvjpeg, nvblas等等.
cuDNN是基于CUDA的卷积算法库, 里面有关于神经网络卷积计算的很多算子, 比如Tensor的+ - * /, 激活函数, 卷积函数等等.
TensorRT是基于cuDNN的一个面向模型应用库, 到这里不需要了解具体卷积的实现, 只需要把模型交给TensorRT用对应parser解析就可以进行推理或者训练.
或者结构如下:
模型 => TensorRT Parser => TensorRT cuDNN Wrapper => cuDNN ops => CUDA Kernel => GPU driver
本篇主要是作为编译入门讲解, 搭建一个基础的双平台CUDA+cuDNN项目.
正文
准备工作
安装CUDA
首先需要用到的工具集主要是VS2019和它的CMake项目模板. 本篇属于进阶内容, 不讲CMake相关的基础, 可以去之前的文章看.
VS版本没有具体需要, CUDA我使用的是10.1版本, Driver 版本大于r418就行. CUDA10.2也是可以的, 11.0以后我目前还未测试.
注意这里是最新版本,一般我也喜欢用最新版本, 但是由于服务器部署要求, 所以只能下载10.1版本, 需要点击下面这个链接去下载历史版本:
网址:https://developer.nvidia.com/cuda-toolkit-archive
Windows安装没啥好说的, 就是让你选择需要的组件, 这里提醒下, 同时会安装GPU Driver, 一般10.1的driver都比较老, 所以你可以选择不安装driver只安装cuda toolkit.
Linux安装参照https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html, 一定要按照官网的步骤安装, 建议使用deb local的方式安装.
安装cuDNN
地址:https://developer.nvidia.com/rdp/cudnn-download
这里我还是下载老版本7.6.5
Windows的安装我有点忘了, cuDNN好像是解压成一个文件夹需要配置环境变量. 把解压的几个文件夹直接复制到cuda的文件夹下面最直接, 除非你需要多个版本的cudnn对应同一个cuda版本.
解压出来应该有include lib等几个文件夹,直接复制到对应的cuda目录下就好:
Linux这边要简单一点, 直接安装两个deb文件就行了. 这里有顺序关系, 一定不能乱序安装.
先是运行库, 然后是开发库, 注意文件名有没有dev. 最后有一个sample和documentation我一般不装.
CMake配置
基础的比如项目怎么创建, Linux或者WSL目标怎么添加请看我之前的文章. 这里主要讲怎么在Win和Linux下设置CUDA的链接位置和头文件位置.
请看下面的CMakeLists.txt:
cmake_minimum_required(VERSION 3.18)
project(cudnnBug)
if (UNIX)
set(CUDA_TOOLKIT_ROOT_DIR "/usr/local/cuda")
set(CMAKE_CUDA_COMPILER "${CUDA_TOOLKIT_ROOT_DIR}/bin/nvcc") #要在15行之前先设置Nvcc的路径
set(CUDA_LIB_DIR "${CUDA_TOOLKIT_ROOT_DIR}/lib64")
elseif (WIN32)
find_package(CUDA)
set(CUDA_LIB_DIR "${CUDA_TOOLKIT_ROOT_DIR}/lib/x64")
endif()
set(CUDA_INCLUDE "${CUDA_TOOLKIT_ROOT_DIR}/include")
project(${PROJECT_NAME} LANGUAGES CXX CUDA) #使能CUDA, 在这里会检查nvcc路径是否正确
FILE(GLOB SRCS "*.cpp" "*.cc" "*.cu") #将当前路径下所有的源文件写入SRCS , 如果需要搜索文件夹内的源文件则把GLOB_RECURSE替换GLOB
FILE(GLOB INCS "*.h" "*.hpp" "*.cuh") #同上,这里是头文件
add_executable(${PROJECT_NAME} ${INCS} ${SRCS})
set_property(TARGET ${PROJECT_NAME} PROPERTY CUDA_ARCHITECTURES 61 70 75)#这个很重要, 这是编译的gpu代码支持的平台, 比如需要在1080TI运行就需要加上61
#另外CUDA_ARCHITECTURES好像是cmake 3.18的新特性,老的cmake 需要用 set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -gencode=arch=compute_70,code=sm_70")
target_include_directories(${PROJECT_NAME} PUBLIC ${CUDA_INCLUDE})
target_link_directories(${PROJECT_NAME} PUBLIC ${CUDA_LIB_DIR})
target_link_libraries(${PROJECT_NAME} PUBLIC cudart cudnn)
if(WIN32)
elseif(UNIX)
target_link_libraries(${PROJECT_NAME} PUBLIC pthread)
endif()
上面注释基本都解释了核心知识点, 还需要补充下怎么查询显卡对应的CUDA_ARCHITECTURES.
网址在这里:https://developer.nvidia.com/cuda-GPUs
到这里就完成了Windows和Linux上的CMake配置.
一些测试代码
下面的代码是一个cudnn的例子, 是我曾经给Nvidia反馈的一个bug.
代码主要示范了三种卷积形式:
- 先用cudnnConvolutionForward卷积, 然后用cudnnAddTensor加上bias
- 直接用cudnnConvolutionBiasActivationForward完成卷积和bias叠加, 使用CUDNN_ACTIVATION_IDENTITY作为激活
- 直接用cudnnConvolutionBiasActivationForward完成卷积和bias叠加, 使用CUDNN_ACTIVATION_RELU作为激活
其中正常情况下1和2的结果应该一致, 3是有Relu的效果. 但是bug的情况下, 2和3都有Relu的效果. Bug发生在dilation rate大于3的时候. 经过英伟达的反馈这个bug在cuDNN8下面被修复了, cuDNN7应该还存在.
本项目只包含一个cudnn_convbug.cu, 代码如下:
#include "cuda_runtime.h"
#include "cudnn.h"
#include <stdio.h>
#include <vector>
#define CheckCUDNN(ret) \
{\
auto tmp = ret;\
if(tmp!=CUDNN_STATUS_SUCCESS)\
{\
printf("CuDNN Error %d: %d",__LINE__,tmp);\
return -1;\
}\
}
#define CheckCUDA(ret) \
{\
auto tmp=ret;\
if(tmp!=cudaSuccess)\
{\
printf("CUDA Error %d: %d",__LINE__,tmp);\
return -1;\
}\
}
static int cudnnBugTest(int dialation)
{
constexpr int width = 28, height = 16;
constexpr int bytesize = width * height * sizeof(float);
int logsize = 100;
cudnnHandle_t hcudnn = NULL;
CheckCUDNN(cudnnCreate(&hcudnn));
cudnnTensorDescriptor_t hinputtensor, houtputtensor, bias;
cudnnCreateTensorDescriptor(&hinputtensor);
cudnnCreateTensorDescriptor(&houtputtensor);
cudnnCreateTensorDescriptor(&bias);
cudnnSetTensor4dDescriptor(hinputtensor, cudnnTensorFormat_t::CUDNN_TENSOR_NCHW, cudnnDataType_t::CUDNN_DATA_FLOAT, 1, 1, height, width);
cudnnSetTensor4dDescriptor(houtputtensor, cudnnTensorFormat_t::CUDNN_TENSOR_NCHW, cudnnDataType_t::CUDNN_DATA_FLOAT, 1, 3, height, width);
cudnnSetTensor4dDescriptor(bias, cudnnTensorFormat_t::CUDNN_TENSOR_NCHW, cudnnDataType_t::CUDNN_DATA_FLOAT, 1, 3, 1, 1);
float* src_dev, * tar_dev;
float* src_h, * tar_h;
float* bias_dev;
float bias_h[] = { -100,-90,-80 };
cudaMalloc(&src_dev, bytesize);
cudaMalloc(&tar_dev, bytesize * 3);
cudaMallocHost(&src_h, bytesize);
cudaMallocHost(&tar_h, bytesize * 3);
cudaMalloc(&bias_dev, 3 * sizeof(float));
cudaMemcpy(bias_dev, bias_h, 3 * sizeof(float), cudaMemcpyKind::cudaMemcpyHostToDevice);
for (int i = 0; i < height; i++)
{
for (int j = 0; j < width; j++)
{
*(src_h + i * width + j) = (float)(j);
}
}
cudaMemcpy(src_dev, src_h, bytesize, cudaMemcpyKind::cudaMemcpyHostToDevice);
float alpha(1), beta(0);
cudnnConvolutionDescriptor_t hconv0;
cudnnCreateConvolutionDescriptor(&hconv0);
CheckCUDNN(cudnnSetConvolution2dDescriptor(hconv0, dialation, dialation, 1, 1, dialation, dialation, cudnnConvolutionMode_t::CUDNN_CROSS_CORRELATION, cudnnDataType_t::CUDNN_DATA_FLOAT));
float* filter_h, * filter_d;
constexpr int filtersizeb = 3 * 3 * 3 * sizeof(float);
cudaMalloc(&filter_d, filtersizeb);
cudaMallocHost(&filter_h, filtersizeb);
for (int i = 0; i < 3 * 3 * 3; i++)
{
*(filter_h + i) = 1;
}
cudaMemcpy(filter_d, filter_h, filtersizeb, cudaMemcpyKind::cudaMemcpyHostToDevice);
cudnnFilterDescriptor_t hfilter0;
cudnnCreateFilterDescriptor(&hfilter0);
CheckCUDNN(cudnnSetFilter4dDescriptor(hfilter0, cudnnDataType_t::CUDNN_DATA_FLOAT, cudnnTensorFormat_t::CUDNN_TENSOR_NCHW, 3, 1, 3, 3));
size_t sizeInBytes = 0;
void* workSpace = NULL;
cudnnConvolutionFwdAlgo_t convalgo = cudnnConvolutionFwdAlgo_t::CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM;
CheckCUDNN(cudnnGetConvolutionForwardWorkspaceSize(hcudnn, hinputtensor, hfilter0, hconv0, houtputtensor, convalgo, &sizeInBytes));
if (sizeInBytes)
{
cudaMalloc(&workSpace, sizeInBytes);
}
CheckCUDNN(cudnnConvolutionForward(hcudnn, &alpha, hinputtensor, src_dev, hfilter0, filter_d, hconv0
, convalgo
, workSpace, sizeInBytes, &beta, houtputtensor, tar_dev));
cudnnAddTensor(hcudnn, &alpha, bias, bias_dev, &alpha, houtputtensor, tar_dev);
std::vector<float> out0, out1, out2;
out0.resize(width * height * 3);
out1.resize(width * height * 3);
out2.resize(width * height * 3);
cudaMemcpy(out0.data(), tar_dev, bytesize * 3, cudaMemcpyKind::cudaMemcpyDeviceToHost);//这里tar_h是正确的计算结果
cudnnActivationDescriptor_t act0;
cudnnCreateActivationDescriptor(&act0);
CheckCUDNN(cudnnSetActivationDescriptor(act0, cudnnActivationMode_t::CUDNN_ACTIVATION_IDENTITY, cudnnNanPropagation_t::CUDNN_NOT_PROPAGATE_NAN, 0));
CheckCUDNN(cudnnConvolutionBiasActivationForward(hcudnn, &alpha, hinputtensor, src_dev, hfilter0, filter_d, hconv0
, convalgo
, workSpace, sizeInBytes, &beta, houtputtensor, tar_dev, bias, bias_dev, act0, houtputtensor, tar_dev));
cudaMemcpy(out1.data(), tar_dev, bytesize * 3, cudaMemcpyKind::cudaMemcpyDeviceToHost);//这里tar_h是被RELU
CheckCUDNN(cudnnSetActivationDescriptor(act0, cudnnActivationMode_t::CUDNN_ACTIVATION_RELU, cudnnNanPropagation_t::CUDNN_NOT_PROPAGATE_NAN, 10));
CheckCUDNN(cudnnConvolutionBiasActivationForward(hcudnn, &alpha, hinputtensor, src_dev, hfilter0, filter_d, hconv0
, convalgo
, workSpace, sizeInBytes, &beta, houtputtensor, tar_dev, bias, bias_dev, act0, houtputtensor, tar_dev));
cudaMemcpy(out2.data(), tar_dev, bytesize * 3, cudaMemcpyKind::cudaMemcpyDeviceToHost);//这里tar_h和CUDNN_ACTIVATION_IDENTITY的结果一样
printf("Method0\tMethod1\tMethod2\n");
for (size_t i = 0; i < logsize; i++)
{
printf("%f\t%f\t%f\n", out0[i], out1[i], out2[i]);
}
//TODO 清理cudaMalloc和cudaMallocHost的内存,由于只运行一次这里就暂时没加.
CheckCUDNN(cudnnDestroy(hcudnn));
return 1;
}
int main(int ac,const char**as)
{
int dilation = 5;
if (ac>1)
{
dilation = atoi(as[1]);
}
cudnnBugTest(dilation);
CheckCUDA(cudaDeviceReset());
return 0;
}
上面的程序主要是测试在更改dilation rate的情况下, 对三个卷积结果的改变. 运行如下:
三列数据分别对应上面说的三种卷积方式, 当某个dilation rate导致第一列和第二列数据不相等的时候, bug就产生了.
后文
本篇没有对CUDA语法和API进行讲解, 也没有示范Launch kernel, 这部分内容我建议直接去看CUDA Samples.
下面是我推荐的几个sample: