【CUDA】cuDNN:加速深度学习的核心库
1. 什么是 cuDNN?
cuDNN(CUDA Deep Neural Network library)是 NVIDIA 提供的一个高性能 GPU 加速库,专为深度学习任务中常见的操作进行了高度优化。它不仅提供了单操作的高效实现,还支持 多操作融合(fusion),旨在最大化地利用 NVIDIA GPU 的计算能力。
cuDNN 能做什么?
cuDNN 支持以下常见深度学习操作:
- 卷积操作(Convolution forward/backward,包括交叉相关)。
- GEMM(通用矩阵乘法,General Matrix Multiply)。
- 池化操作(Pooling forward/backward)。
- 激活函数(如 ReLU、Tanh、Sigmoid、ELU、GELU、Softplus、Swish)。
- Softmax(forward/backward)。
- 点操作(Pointwise operations:算术、逻辑、关系操作)。
- 张量变换(如 reshape、transpose、concat)。
- 归一化操作:Batch Normalization、Instance Normalization、Layer Normalization。
- 运行时融合:动态融合多个操作(如卷积 + 激活函数),减少内存访问。
特点:cuDNN 提供了高度优化的单操作引擎,并在新版本中引入了 Graph API,允许用户定义操作图,实现更灵活的内核融合。
2. 卷积操作:从理论到实践
2.1 卷积的两种实现方式
卷积在深度学习中广泛用于图像分类、检测等任务。cuDNN 支持高效实现卷积操作,主要依赖于以下两种方法:
- 直接卷积(Slow Convolution):基于数学定义逐元素计算卷积,计算复杂度较高。
- 快速卷积(Fast Convolution):通过 FFT(快速傅里叶变换)或者将卷积转化为矩阵乘法(GEMM)来加速计算。
在 cuDNN 中,快速卷积通过 GEMM 的实现更为常见,因为现代 GPU 对矩阵乘法的优化非常强大。
2.2 cuDNN 卷积 API 的使用流程
cuDNN 中实现卷积操作的主要步骤如下:
1. 创建 cuDNN 句柄
所有 cuDNN 操作都需要一个上下文句柄 cudnnHandle_t
,用于初始化库环境。
cudnnHandle_t cudnn;
cudnnCreate(&cudnn);
2. 定义输入和输出张量描述符
使用 cudnnTensorDescriptor_t
来描述输入、输出张量的形状和数据格式。例如:
cudnnTensorDescriptor_t inputDesc, outputDesc;
cudnnCreateTensorDescriptor(&inputDesc);
cudnnSetTensor4dDescriptor(inputDesc, CUDNN_TENSOR_NCHW, // 数据格式:批量、通道、高度、宽度CUDNN_DATA_FLOAT, // 数据类型:floatbatch_size, channels, height, width);cudnnCreateTensorDescriptor(&outputDesc);
cudnnSetTensor4dDescriptor(outputDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, batch_size, output_channels, output_height, output_width);
3. 定义卷积操作描述符
使用 cudnnConvolutionDescriptor_t
来描述卷积核的参数,比如步幅(stride)、填充(padding)等:
cudnnConvolutionDescriptor_t convDesc;
cudnnCreateConvolutionDescriptor(&convDesc);
cudnnSetConvolution2dDescriptor(convDesc, pad_h, pad_w, // 填充stride_h, stride_w, // 步幅dilation_h, dilation_w, // 扩张CUDNN_CROSS_CORRELATION, // 交叉相关CUDNN_DATA_FLOAT);
4. 定义卷积核(Filter)描述符
通过 cudnnFilterDescriptor_t
来设置卷积核的形状和数据格式:
cudnnFilterDescriptor_t filterDesc;
cudnnCreateFilterDescriptor(&filterDesc);
cudnnSetFilter4dDescriptor(filterDesc, CUDNN_DATA_FLOAT, // 数据类型CUDNN_TENSOR_NCHW, // 数据格式output_channels, input_channels, kernel_h, kernel_w);
5. 选择卷积前向算法
cuDNN 提供了多种卷积前向算法(如 CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM
),可以通过性能测试选择最优算法:
cudnnConvolutionFwdAlgo_t algo;
cudnnGetConvolutionForwardAlgorithm(cudnn, inputDesc, filterDesc, convDesc, outputDesc, CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo);
6. 分配工作空间(Workspace)
某些卷积算法需要额外的 GPU 内存工作空间:
size_t workspaceSize;
cudnnGetConvolutionForwardWorkspaceSize(cudnn, inputDesc, filterDesc, convDesc, outputDesc, algo, &workspaceSize);void *workspace;
cudaMalloc(&workspace, workspaceSize);
7. 执行卷积前向操作
使用 cudnnConvolutionForward
完成卷积计算:
float alpha = 1.0f, beta = 0.0f;
cudnnConvolutionForward(cudnn, &alpha, inputDesc, d_input, filterDesc, d_kernel, convDesc, algo, workspace, workspaceSize, &beta, outputDesc, d_output);
8. 释放资源
执行完毕后,释放分配的内存和描述符:
cudaFree(workspace);
cudnnDestroyTensorDescriptor(inputDesc);
cudnnDestroyTensorDescriptor(outputDesc);
cudnnDestroyFilterDescriptor(filterDesc);
cudnnDestroyConvolutionDescriptor(convDesc);
cudnnDestroy(cudnn);
3. cuDNN 内核融合:高效执行多操作
3.1 什么是内核融合?
内核融合(Kernel Fusion)是指将多个操作组合成一个 CUDA 内核执行,从而减少 GPU 的内存读写次数,提升计算性能。例如:
output = torch.sigmoid(tensor1 + tensor2 * tensor3)
传统执行:每个操作(加法、乘法、激活)会触发一个独立的 CUDA 内核。 融合执行:所有操作合并为一个内核,避免冗余的内存访问。
3.2 cuDNN 的内核融合引擎
cuDNN 提供以下几种融合引擎:
- 通用运行时融合引擎(Generic Runtime Fusion Engines):支持灵活组合多个操作。
- 特定运行时融合引擎(Specialized Runtime Fusion Engines):针对特定操作序列进行了优化(如卷积 + 激活)。
- 预编译融合引擎(Pre-compiled Fusion Engines):对特定操作序列进行预编译,性能极高但缺乏灵活性。
3.3 Graph API:灵活定义操作图
cuDNN 在 v8 版本引入了 Graph API,允许用户以操作图的形式定义计算。操作节点代表计算(如卷积、激活),边代表张量。
- 优势:提供更大的灵活性,支持动态融合和运行时编译。
- 应用:特别适用于需要高度优化的复杂操作序列。
4. 性能优化与实践
4.1 性能基准测试
对于卷积操作,cuDNN 提供多种前向算法。可以测试不同算法的性能,选择最快的实现:
CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM
CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM
CUDNN_CONVOLUTION_FWD_ALGO_FFT
CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD
4.2 自定义内核
对于特殊场景(如非批处理任务),可以编写自定义 CUDA 内核,结合 cuDNN 进行优化。
5. cuDNN Graph API:灵活定义和执行计算图
5.1 什么是 Graph API?
Graph API 是 cuDNN v8 引入的一个新特性,它允许用户将一系列深度学习操作以 计算图(computation graph)的形式定义,并通过一次性执行整个图来提高性能。
在传统的计算模式中,每个操作(例如卷积、激活、归一化)都是独立的 CUDA 内核,执行时需要多次进行 GPU 内存读写,导致性能瓶颈。
Graph API 将多个操作融合成一个计算图,优势包括:
- 减少内存读写:数据在 GPU 上的中间结果不会频繁写回内存,而是直接在图中流动。
- 动态编译优化:cuDNN 可以自动编译并优化整个计算图。
- 减少调度开销:CUDA 内核调度的次数减少,整体执行更快。
5.2 Graph API 的操作流程
使用 cuDNN 的 Graph API 可以分为以下几个步骤:
1. 创建 Graph 句柄
使用 cudnnBackendDescriptor_t
创建一个计算图的描述符。
cudnnHandle_t cudnn;
cudnnCreate(&cudnn);cudnnBackendDescriptor_t graph;
cudnnBackendCreateDescriptor(CUDNN_BACKEND_EXECUTION_PLAN_DESCRIPTOR, &graph);
2. 定义操作节点
在计算图中,每个操作(如卷积、激活、池化)都会成为一个 节点,这些节点通过张量(tensor)进行连接。
定义输入和输出张量
cudnnBackendTensorDescriptor_t inputTensor, outputTensor;
// 输入张量
cudnnBackendCreateDescriptor(CUDNN_BACKEND_TENSOR_DESCRIPTOR, &inputTensor);
cudnnBackendSetAttribute(inputTensor, CUDNN_ATTR_TENSOR_DATA_TYPE, CUDNN_TYPE_DATA_TYPE, 1, &dataType);
cudnnBackendSetAttribute(inputTensor, CUDNN_ATTR_TENSOR_DIMENSIONS, CUDNN_TYPE_INT64, 4, dims);
添加卷积操作
cudnnBackendDescriptor_t convNode;
cudnnBackendCreateDescriptor(CUDNN_BACKEND_OPERATION_CONVOLUTION_DESCRIPTOR, &convNode);
cudnnBackendSetAttribute(convNode, CUDNN_ATTR_OPERATION_CONVOLUTION_CONV_DESC, CUDNN_TYPE_CONVOLUTION_DESC, 1, &convDesc);
添加激活操作
cudnnBackendDescriptor_t reluNode;
cudnnBackendCreateDescriptor(CUDNN_BACKEND_OPERATION_POINTWISE_DESCRIPTOR, &reluNode);
cudnnBackendSetAttribute(reluNode, CUDNN_ATTR_OPERATION_POINTWISE_MODE, CUDNN_TYPE_POINTWISE_MODE, 1, &reluMode);
3. 将节点连接成计算图
通过设置张量的输入输出,来连接各个操作节点,形成完整的计算图。
cudnnBackendSetAttribute(convNode, CUDNN_ATTR_OPERATION_CONVOLUTION_X, CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &inputTensor);
cudnnBackendSetAttribute(convNode, CUDNN_ATTR_OPERATION_CONVOLUTION_Y, CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &outputTensor);// 将激活操作的输入设为卷积的输出
cudnnBackendSetAttribute(reluNode, CUDNN_ATTR_OPERATION_POINTWISE_X, CUDNN_TYPE_BACKEND_DESCRIPTOR, 1, &outputTensor);
4. 执行计算图
构建好计算图后,使用 cuDNN 的 cudnnBackendFinalize
函数对图进行编译并执行。
cudnnBackendFinalize(graph);
cudnnBackendExecute(graph, executionPlan);
5. Graph API 的性能优化
Graph API 可以根据实际的计算图进行多种优化:
- 内核融合:自动将多个操作融合成一个高效的 CUDA 内核。
- 调度优化:减少 GPU 的调度开销。
- 内存优化:避免不必要的内存复制,数据流在 GPU 内高效传输。
6. cuDNN 内核融合 (Kernel Fusion)
6.1 内核融合的原理
内核融合是 cuDNN 提高性能的重要手段,目标是减少 GPU 内核之间的内存读写开销,将多个操作合并为一个内核执行。例如:
- 卷积 + 激活函数(ReLU)
- 卷积 + 批量归一化(BatchNorm)+ 激活函数
6.2 内核融合的两种模式
- 静态融合(Static Fusion):
- 预定义常用操作的融合模式,比如卷积 + ReLU。
- 性能最佳,但缺乏灵活性。
- 动态融合(Dynamic Fusion):
- 在运行时动态组合用户定义的操作。
- 使用 Graph API 实现,灵活性更高,但需要一定的编译开销。
6.3 使用内核融合的最佳实践
在 cuDNN 中,用户可以选择直接使用 Pointwise 操作 和 Graph API 来实现内核融合:
Pointwise 操作示例
Pointwise 操作可以执行逐元素的运算,例如 Add
、Multiply
和 ReLU
等:
cudnnBackendDescriptor_t pointwiseDesc;
cudnnBackendCreateDescriptor(CUDNN_BACKEND_OPERATION_POINTWISE_DESCRIPTOR, &pointwiseDesc);
cudnnBackendSetAttribute(pointwiseDesc, CUDNN_ATTR_OPERATION_POINTWISE_MODE, CUDNN_TYPE_POINTWISE_MODE, 1, &pointwiseMode);
Graph API 实现复杂融合
通过 Graph API 将多个点操作与卷积结合,形成更复杂的内核融合计算图。
7. cuDNN 优化技巧总结
- 选择最优卷积算法: 使用
cudnnGetConvolutionForwardAlgorithm
动态选择性能最优的卷积前向算法。 - 最小化内存工作空间: 对于 GPU 内存有限的场景,可以通过指定工作空间大小来选择算法。
- 使用 Graph API 进行内核融合: 将多个操作合并成一个计算图,减少内存读写和调度开销。
- 预热 GPU 内核: 在实际训练之前,先运行一遍前向和反向计算,让 GPU 完成内核编译和优化。
代码示例
Tanh.cu
这段代码是一个完整的CUDA和cuDNN示例程序,用于比较使用朴素CUDA核函数和cuDNN库实现tanh
激活函数的性能和正确性。
#include <cuda_runtime.h>
#include <cudnn.h>
#include <math.h>
#include <stdio.h>
#include <stdlib.h>#define CHECK_CUDA(call) \{ \cudaError_t err = call; \if (err != cudaSuccess) { \fprintf(stderr, "CUDA error in file '%s' in line %i : %s.\n", __FILE__, __LINE__, \cudaGetErrorString(err)); \exit(EXIT_FAILURE); \} \}#define CHECK_CUDNN(call) \{ \cudnnStatus_t err = call; \if (err != CUDNN_STATUS_SUCCESS) { \fprintf(stderr, "cuDNN error in file '%s' in line %i : %s.\n", __FILE__, __LINE__, \cudnnGetErrorString(err)); \exit(EXIT_FAILURE); \} \}__global__ void NaiveTankKernel(float* input, float* output, int size) {int idx = blockIdx.x * blockDim.x + threadIdx.x;if (idx < size) {output[idx] = tanhf(input[idx]);}
}float CpuTanh(float x) { return tanhf(x); }void InitializeData(float* data, int size) {for (int i = 0; i < size; ++i) {// Random values between -1 and 1data[i] = (float)rand() / RAND_MAX * 2.0f - 1.0f;}
}bool VerifyResults(float* cpu_output, float* gpu_output, int size, float tolerance = 1e-5) {for (int i = 0; i < size; ++i) {if (fabs(cpu_output[i] - gpu_output[i]) > tolerance) {printf("Mismatch at index %d: CPU = %f, GPU = %f\n", i, cpu_output[i], gpu_output[i]);return false;}}return true;
}int main() { // Set up tensor dimensions for a scenario where cuDNN is likely to outperformconst int batch_size = 256; // NCHW formatconst int channels = 32;const int height = 224;const int width = 224;const int tensor_size = batch_size * channels * height * width;// Allocate host memoryfloat *h_input, *h_output_naive, *h_output_cudnn, *h_output_cpu;h_input = (float*)malloc(tensor_size * sizeof(float));h_output_naive = (float*)malloc(tensor_size * sizeof(float));h_output_cudnn = (float*)malloc(tensor_size * sizeof(float));h_output_cpu = (float*)malloc(tensor_size * sizeof(float));InitializeData(h_input, tensor_size);// Allocate device memoryfloat *d_input, *d_output_naive, *d_output_cudnn;CHECK_CUDA(cudaMalloc(&d_input, tensor_size * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_output_naive, tensor_size * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_output_cudnn, tensor_size * sizeof(float)));// Copy input data to deviceCHECK_CUDA(cudaMemcpy(d_input, h_input, tensor_size * sizeof(float), cudaMemcpyHostToDevice));// Create CUDA events for timingcudaEvent_t start, stop;CHECK_CUDA(cudaEventCreate(&start));CHECK_CUDA(cudaEventCreate(&stop));// Warmup and benchmark parametersconst int num_warmup = 10;const int num_benchmark = 100;float naive_times[num_benchmark];float cudnn_times[num_benchmark];// Naive CUDA kerneldim3 block(256);dim3 grid((tensor_size + block.x - 1) / block.x);// Warmup runs for naive kernelfor (int i = 0; i < num_warmup; ++i) {NaiveTankKernel<<<grid, block>>>(d_input, d_output_naive, tensor_size);}CHECK_CUDA(cudaDeviceSynchronize());for (int i = 0; i < num_benchmark; ++i) {// cudaEventRecord(start) 将当前时间记录在 start 事件中CHECK_CUDA(cudaEventRecord(start));NaiveTankKernel<<<grid, block>>>(d_input, d_output_naive, tensor_size);// cudaEventRecord(stop) 将当前时间记录在 stop 事件中CHECK_CUDA(cudaEventRecord(stop));// cudaEventSynchronize(stop) 等待 stop 事件完成。CHECK_CUDA(cudaEventSynchronize(stop));// cudaEventElapsedTime(&naive_times[i], start, stop) 计算从 start 事件到 stop 事件之间的时间差CHECK_CUDA(cudaEventElapsedTime(&naive_times[i], start, stop));}// cuDNN setup// cudnnHandle_t 是 cuDNN 的句柄,用于管理 cuDNN 库的上下文。cudnnHandle_t cudnn;CHECK_CUDNN(cudnnCreate(&cudnn));cudnnTensorDescriptor_t input_descriptor;/*cudnnSetTensor4dDescriptor 用于设置 4D 张量的描述信息:CUDNN_TENSOR_NCHW:指定张量的布局为 NCHW(Batch, Channels, Height, Width)。CUDNN_DATA_FLOAT:指定张量的数据类型为 float。batch_size:批量大小(即一次处理的样本数量)。channels:通道数(例如 RGB 图像的通道数为 3)。height:张量的高度。width:张量的宽度。*/CHECK_CUDNN(cudnnCreateTensorDescriptor(&input_descriptor));CHECK_CUDNN(cudnnSetTensor4dDescriptor(input_descriptor, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, batch_size, channels,height, width));// cudnnActivationDescriptor_t 是用于描述激活函数的结构。cudnnActivationDescriptor_t activation_descriptor;CHECK_CUDNN(cudnnCreateActivationDescriptor(&activation_descriptor));/*cudnnSetActivationDescriptor 用于设置激活函数的参数:CUDNN_ACTIVATION_TANH:指定激活函数为 tanh。CUDNN_PROPAGATE_NAN:指定在计算过程中如何处理 NaN 值(这里选择传播 NaN)。0.0:对于 tanh 激活函数,不需要额外的参数,因此设置为 0.0。*/CHECK_CUDNN(cudnnSetActivationDescriptor(activation_descriptor, CUDNN_ACTIVATION_TANH, CUDNN_PROPAGATE_NAN, 0.0));float alpha = 1.0f, beta = 0.0f;// Warmup runs for cuDNNfor (int i = 0; i < num_warmup; ++i) {/*cudnnActivationForward 是 cuDNN 提供的函数,用于执行激活函数的前向传播:cudnn:cuDNN 句柄。activation_descriptor:激活函数描述符。&alpha 和 &beta:缩放因子。input_descriptor 和 d_input:输入张量的描述符和设备指针。input_descriptor 和 d_output_cudnn:输出张量的描述符和设备指针。*/CHECK_CUDNN(cudnnActivationForward(cudnn, activation_descriptor, &alpha, input_descriptor, d_input, &beta,input_descriptor, d_output_cudnn));}CHECK_CUDA(cudaDeviceSynchronize());// Benchmark runs for cuDNNfor (int i = 0; i < num_benchmark; ++i) {CHECK_CUDA(cudaEventRecord(start));CHECK_CUDNN(cudnnActivationForward(cudnn, activation_descriptor, &alpha, input_descriptor, d_input, &beta,input_descriptor, d_output_cudnn));CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));CHECK_CUDA(cudaEventElapsedTime(&cudnn_times[i], start, stop));}// Calculate average timesfloat avg_naive_time = 0.0f, avg_cudnn_time = 0.0f;for (int i = 0; i < num_benchmark; ++i) {avg_naive_time += naive_times[i];avg_cudnn_time += cudnn_times[i];}avg_naive_time /= num_benchmark;avg_cudnn_time /= num_benchmark;// Copy results back to hostCHECK_CUDA(cudaMemcpy(h_output_naive, d_output_naive, tensor_size * sizeof(float), cudaMemcpyDeviceToHost));CHECK_CUDA(cudaMemcpy(h_output_cudnn, d_output_cudnn, tensor_size * sizeof(float), cudaMemcpyDeviceToHost));// CPU verificationfor (int i = 0; i < tensor_size; ++i) {h_output_cpu[i] = CpuTanh(h_input[i]);}// Verify resultsbool naive_correct = VerifyResults(h_output_cpu, h_output_naive, tensor_size);bool cudnn_correct = VerifyResults(h_output_cpu, h_output_cudnn, tensor_size);// Print resultsprintf("Tensor size: %d x %d x %d x %d\n", batch_size, channels, height, width);printf("Average Naive CUDA kernel time: %.3f ms\n", avg_naive_time);printf("Average cuDNN activation time: %.3f ms\n", avg_cudnn_time);printf("Speedup: %.2fx\n", avg_naive_time / avg_cudnn_time);printf("Naive kernel results correct: %s\n", naive_correct ? "Yes" : "No");printf("cuDNN results correct: %s\n", cudnn_correct ? "Yes" : "No");// Clean upCHECK_CUDA(cudaFree(d_input));CHECK_CUDA(cudaFree(d_output_naive));CHECK_CUDA(cudaFree(d_output_cudnn));CHECK_CUDA(cudaEventDestroy(start));CHECK_CUDA(cudaEventDestroy(stop));CHECK_CUDNN(cudnnDestroyTensorDescriptor(input_descriptor));CHECK_CUDNN(cudnnDestroyActivationDescriptor(activation_descriptor));CHECK_CUDNN(cudnnDestroy(cudnn));free(h_input);free(h_output_naive);free(h_output_cudnn);free(h_output_cpu);return 0;
}
结果:
Tensor size: 256 x 32 x 224 x 224
Average Naive CUDA kernel time: 18.201 ms
Average cuDNN activation time: 18.377 ms
Speedup: 0.99x
Naive kernel results correct: Yes
cuDNN results correct: Yes
使用 cuDNN 的性能与朴素 CUDA 核函数几乎相同,甚至略慢一点点,可能是因为激活函数tanh本身已经足够简单,同时cuDNN有一些额外的计算(alpha和beta),所以使用cuDNN不一定会比自定义CUDA内核快。
但如果你不使用CUDA 内核来实现tanh的话,会慢很多,代码见https://github.com/Infatoshi/cuda-course/blob/master/06_CUDA_APIs/02%20CUDNN/00%20torch-compare.py。所以使用CUDA重写确实会快很多。
Conv2d_HCHW.cu
这段代码实现了一个基于CUDA和cuDNN的二维卷积操作的性能对比。它首先定义了一个简单的CUDA核函数 NaiveConv2d
,用于执行朴素的二维卷积操作。然后,代码使用cuDNN库来执行相同的卷积操作,并选择性能最佳的卷积算法。通过对比cuDNN和朴素卷积核的输出结果,代码验证了两者的计算结果是否一致,并测量了它们的执行时间。最终,代码输出卷积结果以及两者的最大差异,并打印了平均执行时间。
#include <cuda_runtime.h>
#include <cudnn.h>
#include <stdio.h>
#include <stdlib.h>#include <iostream>
#include <limits>#define CHECK_CUDA(call) \{ \cudaError_t err = call; \if (err != cudaSuccess) { \printf("CUDA error: %s\n", cudaGetErrorString(err)); \exit(1); \} \}
#define CHECK_CUDNN(call) \{ \cudnnStatus_t err = call; \if (err != CUDNN_STATUS_SUCCESS) { \printf("cuDNN error: %s\n", cudnnGetErrorString(err)); \exit(1); \} \}// Complex multi-channel 2D convolution kernel
__global__ void NaiveConv2d(float* input, float* kernel, float* output, int width, int height, int in_channels,int out_channels, int kernel_size, int batch_size) {int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;int out_channel = blockIdx.z % out_channels;int batch_idx = blockIdx.z / out_channels;// 因为卷积后宽高不变,所以按理是要padding的,但是这里认为padding填充的是0,所以实际上要padding的区域跳过计算,体现在"-half_kernel"开始if (x < width && y < height && out_channel < out_channels && batch_idx < batch_size) {float sum = 0;int half_kernel = kernel_size / 2;for (int in_channel = 0; in_channel < in_channels; ++in_channel) {for (int ky = -half_kernel; ky <= half_kernel; ++ky) {for (int kx = -half_kernel; kx <= half_kernel; ++kx) {int ix = x + kx;int iy = y + ky;if (ix >= 0 && ix < width && iy >= 0 && iy < height) {int input_idx = ((batch_idx * in_channels + in_channel) * height + iy) * width + ix;int kernel_idx = ((out_channel * in_channels + in_channel) * kernel_size + (ky + half_kernel)) *kernel_size +(kx + half_kernel);sum += input[input_idx] * kernel[kernel_idx];}}}}int output_idx = ((batch_idx * out_channels + out_channel) * height + y) * width + x;output[output_idx] = sum;}
}int main() {// Smaller, predefined sizes for human-readable outputconst int width = 4;const int height = 4;const int kernel_size = 3;const int in_channels = 1;const int out_channels = 1;const int batch_size = 1;const int input_size = width * height * in_channels * batch_size;const int output_size = width * height * out_channels * batch_size;const int kernel_elements = kernel_size * kernel_size * in_channels * out_channels;std::cout << "Image size: " << width << "x" << height << "x" << in_channels << std::endl;std::cout << "Kernel size: " << kernel_size << "x" << kernel_size << "x" << in_channels << "x" << out_channels<< std::endl;std::cout << "Batch size: " << batch_size << std::endl;// Allocate host memoryfloat* h_input = (float*)malloc(input_size * sizeof(float));float* h_kernel = (float*)malloc(kernel_elements * sizeof(float));float* h_output_cudnn = (float*)malloc(output_size * sizeof(float));float* h_output_naive = (float*)malloc(output_size * sizeof(float));// Initialize input and kernel with predefined valuesfloat input_values[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16,};float kernel_values[] = {1, 2, 3, 4, 5, 6, 7, 8, 9,};memcpy(h_input, input_values, input_size * sizeof(float));memcpy(h_kernel, kernel_values, kernel_elements * sizeof(float));// Allocate device memoryfloat *d_input, *d_kernel, *d_output_cudnn, *d_output_naive;CHECK_CUDA(cudaMalloc(&d_input, input_size * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_kernel, kernel_elements * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_output_cudnn, output_size * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_output_naive, output_size * sizeof(float)));// Copy data to deviceCHECK_CUDA(cudaMemcpy(d_input, h_input, input_size * sizeof(float), cudaMemcpyHostToDevice));CHECK_CUDA(cudaMemcpy(d_kernel, h_kernel, kernel_elements * sizeof(float), cudaMemcpyHostToDevice));// cuDNN setupcudnnHandle_t cudnn;CHECK_CUDNN(cudnnCreate(&cudnn));cudnnTensorDescriptor_t input_desc, output_desc;cudnnFilterDescriptor_t kernel_desc;cudnnConvolutionDescriptor_t conv_desc;CHECK_CUDNN(cudnnCreateTensorDescriptor(&input_desc));CHECK_CUDNN(cudnnCreateTensorDescriptor(&output_desc));CHECK_CUDNN(cudnnCreateFilterDescriptor(&kernel_desc));CHECK_CUDNN(cudnnCreateConvolutionDescriptor(&conv_desc));CHECK_CUDNN(cudnnSetTensor4dDescriptor(input_desc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, batch_size, in_channels,height, width));CHECK_CUDNN(cudnnSetTensor4dDescriptor(output_desc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, batch_size, out_channels,height, width));/*cudnnStatus_t cudnnSetFilter4dDescriptor(cudnnFilterDescriptor_t filterDesc, // 卷积核的描述符对象cudnnDataType_t dataType, // 卷积核的数据类型(如 CUDNN_DATA_FLOAT 或 CUDNN_DATA_DOUBLE)cudnnTensorFormat_t format, // 卷积核的存储格式(如 CUDNN_TENSOR_NCHW 或 CUDNN_TENSOR_NHWC)int k, // 卷积核的数量(输出通道数)int c, // 卷积核的输入通道数(输入特征图的通道数)int h, // 卷积核的高度int w // 卷积核的宽度)*/CHECK_CUDNN(cudnnSetFilter4dDescriptor(kernel_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, out_channels, in_channels,kernel_size, kernel_size));/*cudnnStatus_t cudnnSetConvolution2dDescriptor(cudnnConvolutionDescriptor_t convDesc, // 卷积操作的描述符对象int pad_h, // 输入特征图在高度方向上的填充大小(padding)int pad_w, // 输入特征图在宽度方向上的填充大小(padding)int u, // 卷积核在高度方向上的步幅(stride)int v, // 卷积核在宽度方向上的步幅(stride)int dilation_h, // 卷积核在高度方向上的膨胀率(dilation)int dilation_w, // 卷积核在宽度方向上的膨胀率(dilation)cudnnConvolutionMode_t mode, // 卷积模式(如 CUDNN_CONVOLUTION 或 CUDNN_CROSS_CORRELATION)cudnnDataType_t computeType // 卷积计算的数据类型(如 CUDNN_DATA_FLOAT 或 CUDNN_DATA_DOUBLE))这里因为卷积后的宽高尺寸不变,所以特征图四周分别填充kernel_size / 2*/CHECK_CUDNN(cudnnSetConvolution2dDescriptor(conv_desc, kernel_size / 2, kernel_size / 2, 1, 1, 1, 1,CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT));// Find the fastest cuDNN aogorithm// CUDNN_CONVOLUTION_FWD_ALGO_COUNT 是 cuDNN 支持的卷积前向传播算法的总数。int requested_algo_count = CUDNN_CONVOLUTION_FWD_ALGO_COUNT;int returned_algo_count;cudnnConvolutionFwdAlgoPerf_t perf_results[CUDNN_CONVOLUTION_FWD_ALGO_COUNT];/*获取所有可用的卷积前向传播算法,并返回它们的性能信息cudnnStatus_t cudnnGetConvolutionForwardAlgorithm_v7(cudnnHandle_t handle, // cuDNN 句柄cudnnTensorDescriptor_t srcDesc, // 输入张量的描述符cudnnFilterDescriptor_t filterDesc, // 卷积核的描述符cudnnConvolutionDescriptor_t convDesc, // 卷积操作的描述符cudnnTensorDescriptor_t destDesc, // 输出张量的描述符int requestedAlgoCount, // 请求的算法数量int *returnedAlgoCount, // 实际返回的算法数量cudnnConvolutionFwdAlgoPerf_t *perfResults // 算法性能结果数组)具体来说,它可以返回以下几种卷积前向传播算法(cudnnConvolutionFwdAlgo_t枚举类型):1. CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM描述: 隐式 GEMM 算法。通过将卷积操作转换为矩阵乘法(GEMM)来实现。特点: 实现简单,但性能可能不如其他算法。2. CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM描述: 隐式预计算 GEMM 算法。在 GEMM 之前进行一些预计算以提高性能。特点: 性能优于 IMPLICIT_GEMM,但仍然可能不如其他算法。3. CUDNN_CONVOLUTION_FWD_ALGO_GEMM描述: 显式 GEMM 算法。直接使用矩阵乘法来实现卷积。特点: 适用于某些特定场景,但通常不如其他算法高效。4. CUDNN_CONVOLUTION_FWD_ALGO_DIRECT描述: 直接卷积算法。直接在空间域中执行卷积操作。特点: 性能较好,适用于大多数常见场景。5. CUDNN_CONVOLUTION_FWD_ALGO_FFT描述: 快速傅里叶变换(FFT)算法。通过将卷积转换为频域中的乘法来实现。特点: 适用于大卷积核或大输入尺寸,但计算复杂度较高。6. CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING描述: FFT 分块算法。通过将输入数据分块并在频域中执行卷积来实现。特点: 适用于中等大小的卷积核和输入尺寸。7. CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD描述: Winograd 算法。通过数学变换减少乘法操作的数量。特点: 性能优异,尤其适用于小卷积核(如 3x3)。8. CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED描述: 非融合 Winograd 算法。与 WINOGRAD 类似,但避免了某些融合操作。特点: 性能略低于 WINOGRAD,但可能更稳定。9. CUDNN_CONVOLUTION_FWD_ALGO_COUNT描述: 算法数量的计数器。用于表示所有可用算法的总数。返回的性能信息cudnnGetConvolutionForwardAlgorithm_v7 返回的 perfResults 数组中,每个元素包含以下性能信息:algo: 算法类型(cudnnConvolutionFwdAlgo_t)。status: 算法的状态(cudnnStatus_t)。time: 算法的执行时间(以毫秒为单位)。memory: 算法所需的工作区内存大小(以字节为单位)。determinism: 算法是否是确定性的(cudnnDeterminism_t)。mathType: 算法的数学类型(cudnnMathType_t)。*/CHECK_CUDNN(cudnnGetConvolutionForwardAlgorithm_v7(cudnn, input_desc, kernel_desc, conv_desc, output_desc,requested_algo_count, &returned_algo_count, perf_results));cudnnConvolutionFwdAlgo_t algo = perf_results[0].algo;for (int i = 1; i < returned_algo_count; ++i) {std::cout << "Algorithm: " << perf_results[i].algo << " Time: " << perf_results[i].time << std::endl;if (perf_results[i].status == CUDNN_STATUS_SUCCESS && perf_results[i].time < perf_results[0].time) {algo = perf_results[i].algo;}}std::cout << "Selected algorithm: " << algo << std::endl;size_t workspace_size;/*cudnnGetConvolutionForwardWorkspaceSize用于返回指定卷积前向传播算法所需的最小工作区大小。工作区是GPU内存的一部分,用于存储卷积操作中的中间结果。通过调用此函数,用户可以为卷积操作分配足够的内存空间,从而确保卷积操作能够顺利执行。cudnnStatus_t cudnnGetConvolutionForwardWorkspaceSize(cudnnHandle_t handle, // cuDNN 句柄cudnnTensorDescriptor_t xDesc, // 输入张量的描述符cudnnFilterDescriptor_t wDesc, // 卷积核的描述符cudnnConvolutionDescriptor_t convDesc, // 卷积操作的描述符cudnnTensorDescriptor_t yDesc, // 输出张量的描述符cudnnConvolutionFwdAlgo_t algo, // 卷积前向传播算法size_t *sizeInBytes // 返回的工作区大小(以字节为单位))*/CHECK_CUDNN(cudnnGetConvolutionForwardWorkspaceSize(cudnn, input_desc, kernel_desc, conv_desc, output_desc, algo,&workspace_size));void* d_workspace;CHECK_CUDA(cudaMalloc(&d_workspace, workspace_size));// Define grid and block sizes for the naive kerneldim3 block_size(16, 16);dim3 grid_size((width + block_size.x - 1) / block_size.x, (height + block_size.y - 1) / block_size.y,out_channels * batch_size);// Warmup and benckmark runsconst int warmup_runs = 5;const int benchmark_runs = 20;float total_time_cudnn = 0;float total_time_naive = 0;float alpha = 1.0f, beta = 0;// Warmup runsfor (int i = 0; i < warmup_runs; ++i) {/*cudnnStatus_t cudnnConvolutionForward(cudnnHandle_t handle, // cuDNN 句柄const void *alpha, // 输入张量的缩放因子cudnnTensorDescriptor_t xDesc, // 输入张量的描述符const void *x, // 输入张量的数据指针cudnnFilterDescriptor_t wDesc, // 卷积核的描述符const void *w, // 卷积核的数据指针cudnnConvolutionDescriptor_t convDesc, // 卷积操作的描述符cudnnConvolutionFwdAlgo_t algo, // 卷积前向传播算法void *workSpace, // 工作区指针size_t workSpaceSizeInBytes, // 工作区大小(以字节为单位)const void *beta, // 输出张量的缩放因子cudnnTensorDescriptor_t yDesc, // 输出张量的描述符void *y // 输出张量的数据指针)*/CHECK_CUDNN(cudnnConvolutionForward(cudnn, &alpha, input_desc, d_input, kernel_desc, d_kernel, conv_desc, algo,d_workspace, workspace_size, &beta, output_desc, d_output_cudnn));NaiveConv2d<<<grid_size, block_size>>>(d_input, d_kernel, d_output_naive, width, height, in_channels,out_channels, kernel_size, batch_size);CHECK_CUDA(cudaDeviceSynchronize());}cudaEvent_t start, stop;CHECK_CUDA(cudaEventCreate(&start));CHECK_CUDA(cudaEventCreate(&stop));for (int i = 0; i < benchmark_runs; ++i) {// cuDNN benchmark/*cudaEventRecord是一个用于记录CUDA事件的函数,其作用是在GPU上异步标记一个时间点,以便后续测量事件之间的时间差。cudaEventRecord实际上并不是执行到该点,然后把时间给start,虽然看起来像,但并没有传入指针不是吗所以本质是一个记录CUDA事件的函数,事件的标记由CUDA内部完成*/CHECK_CUDA(cudaEventRecord(start));CHECK_CUDNN(cudnnConvolutionForward(cudnn, &alpha, input_desc, d_input, kernel_desc, d_kernel, conv_desc, algo,d_workspace, workspace_size, &beta, output_desc, d_output_cudnn));CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));float milliseconds = 0;CHECK_CUDA(cudaEventElapsedTime(&milliseconds, start, stop));total_time_cudnn += milliseconds;// Naive kernel benchmarkCHECK_CUDA(cudaEventRecord(start));NaiveConv2d<<<grid_size, block_size>>>(d_input, d_kernel, d_output_naive, width, height, in_channels,out_channels, kernel_size, batch_size);CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));CHECK_CUDA(cudaEventElapsedTime(&milliseconds, start, stop));total_time_naive += milliseconds;}float avg_time_cudnn = total_time_cudnn / benchmark_runs;float avg_time_naive = total_time_naive / benchmark_runs;printf("cuDNN average time: %f ms\n", avg_time_cudnn);printf("Naive kernel average time: %f ms\n", avg_time_naive);// Copy results back to hostCHECK_CUDA(cudaMemcpy(h_output_cudnn, d_output_cudnn, output_size * sizeof(float), cudaMemcpyDeviceToHost));CHECK_CUDA(cudaMemcpy(h_output_naive, d_output_naive, output_size * sizeof(float), cudaMemcpyDeviceToHost));// Compare resultsfloat max_diff = 0;for (int i = 0; i < output_size; ++i) {float diff = fabs(h_output_cudnn[i] - h_output_naive[i]);if (diff > max_diff) max_diff = diff;}// %e:科学计数法输出浮点数printf("Max difference between cuDNN and naive kernel: %e\n", max_diff);// Print the outputfor (int b = 0; b < batch_size; ++b) {for (int c = 0; c < out_channels; ++c) {printf("Channel %d:\n", c);for (int h = 0; h < height; ++h) {for (int w = 0; w < width; ++w) {int idx = ((b * out_channels + c) * height + h) * width + w;printf("%f ", h_output_cudnn[idx]);}printf("\n");}printf("\n");}}printf("\nNaive Kernel Output:\n");for (int b = 0; b < batch_size; b++) {for (int c = 0; c < out_channels; c++) {printf("Channel %d:\n", c);for (int h = 0; h < height; h++) {for (int w = 0; w < width; w++) {int idx = ((b * out_channels + c) * height + h) * width + w;printf("%f ", h_output_naive[idx]);}printf("\n");}printf("\n");}}// Print flattened output for easier comparison with PyTorchprintf("\nFlattened cuDNN Output:\n");for (int i = 0; i < output_size; i++) {printf("%f", h_output_cudnn[i]);if (i < output_size - 1) printf(", ");}printf("\n");// Clean upCHECK_CUDNN(cudnnDestroyTensorDescriptor(input_desc));CHECK_CUDNN(cudnnDestroyTensorDescriptor(output_desc));CHECK_CUDNN(cudnnDestroyFilterDescriptor(kernel_desc));CHECK_CUDNN(cudnnDestroyConvolutionDescriptor(conv_desc));CHECK_CUDNN(cudnnDestroy(cudnn));CHECK_CUDA(cudaFree(d_input));CHECK_CUDA(cudaFree(d_kernel));CHECK_CUDA(cudaFree(d_output_cudnn));CHECK_CUDA(cudaFree(d_output_naive));CHECK_CUDA(cudaFree(d_workspace));CHECK_CUDA(cudaEventDestroy(start));CHECK_CUDA(cudaEventDestroy(stop));free(h_input);free(h_kernel);free(h_output_cudnn);free(h_output_naive);return 0;
}
输出:
Image size: 4x4x1
Kernel size: 3x3x1x1
Batch size: 1
Algorithm: 0 Time: -1
Algorithm: 2 Time: -1
Algorithm: 6 Time: -1
Algorithm: 4 Time: -1
Algorithm: 5 Time: -1
Algorithm: 7 Time: -1
Algorithm: 3 Time: -1
Selected algorithm: 1
cuDNN average time: 0.031240 ms
Naive kernel average time: 0.006974 ms
Max difference between cuDNN and naive kernel: 0.000000e+00
Channel 0:
111.000000 178.000000 217.000000 145.000000
231.000000 348.000000 393.000000 252.000000
363.000000 528.000000 573.000000 360.000000
197.000000 274.000000 295.000000 175.000000 Naive Kernel Output:
Channel 0:
111.000000 178.000000 217.000000 145.000000
231.000000 348.000000 393.000000 252.000000
363.000000 528.000000 573.000000 360.000000
197.000000 274.000000 295.000000 175.000000 Flattened cuDNN Output:
111.000000, 178.000000, 217.000000, 145.000000, 231.000000, 348.000000, 393.000000, 252.000000, 363.000000, 528.000000, 573.000000, 360.000000, 197.000000, 274.000000, 295.000000, 175.000000
虽然在这里你看到Naive kernel要比cuDNN快,但实际上是数据量太小,在下面的示例中你会看到cuDNN的真正实力。
Compaer_Conv.cu
这段程序实现了一个使用CUDA和cuDNN进行二维卷积操作的示例,比较了基于cuDNN优化的卷积与简单的CUDA卷积实现("naive"实现)在性能上的差异。
#include <cuda_runtime.h>
#include <cudnn.h>
#include <stdio.h>
#include <stdlib.h>#include <iostream>
#include <limits>#define CHECK_CUDA(call) \{ \cudaError_t err = call; \if (err != cudaSuccess) { \printf("CUDA error: %s\n", cudaGetErrorString(err)); \exit(1); \} \}
#define CHECK_CUDNN(call) \{ \cudnnStatus_t err = call; \if (err != CUDNN_STATUS_SUCCESS) { \printf("cuDNN error: %s\n", cudnnGetErrorString(err)); \exit(1); \} \}// Complex multi-channel 2D convolution kernel
__global__ void naiveConv2d(float* input, float* kernel, float* output, int width, int height, int inChannels,int outChannels, int kernelSize, int batchSize) {int x = blockIdx.x * blockDim.x + threadIdx.x;int y = blockIdx.y * blockDim.y + threadIdx.y;int outChannel = blockIdx.z % outChannels;int batchIdx = blockIdx.z / outChannels;if (x < width && y < height && outChannel < outChannels && batchIdx < batchSize) {float sum = 0.0f;int halfKernel = kernelSize / 2;for (int inChannel = 0; inChannel < inChannels; inChannel++) {for (int ky = -halfKernel; ky <= halfKernel; ky++) {for (int kx = -halfKernel; kx <= halfKernel; kx++) {int ix = x + kx;int iy = y + ky;if (ix >= 0 && ix < width && iy >= 0 && iy < height) {int inputIdx = ((batchIdx * inChannels + inChannel) * height + iy) * width + ix;int kernelIdx =((outChannel * inChannels + inChannel) * kernelSize + (ky + halfKernel)) * kernelSize +(kx + halfKernel);sum += input[inputIdx] * kernel[kernelIdx];}}}}int outputIdx = ((batchIdx * outChannels + outChannel) * height + y) * width + x;output[outputIdx] = sum;}
}int main() {// Smaller, predefined sizes for human-readable outputconst int width = 224;const int height = 224;const int kernelSize = 11;const int inChannels = 32;const int outChannels = 64;const int batchSize = 4;const int inputSize = width * height * inChannels * batchSize;const int outputSize = width * height * outChannels * batchSize;const int kernelElements = kernelSize * kernelSize * inChannels * outChannels;std::cout << "Image size: " << width << "x" << height << "x" << inChannels << std::endl;std::cout << "Kernel size: " << kernelSize << "x" << kernelSize << "x" << inChannels << "x" << outChannels<< std::endl;std::cout << "Batch size: " << batchSize << std::endl;// Allocate host memoryfloat* h_input = (float*)malloc(inputSize * sizeof(float));float* h_kernel = (float*)malloc(kernelElements * sizeof(float));float* h_output_cudnn = (float*)malloc(outputSize * sizeof(float));float* h_output_naive = (float*)malloc(outputSize * sizeof(float));// Initialize input and kernel with random valuessrand(time(NULL));for (int i = 0; i < inputSize; i++) {h_input[i] = static_cast<float>(rand()) / RAND_MAX;}for (int i = 0; i < kernelElements; i++) {h_kernel[i] = static_cast<float>(rand()) / RAND_MAX;}// Allocate device memoryfloat *d_input, *d_kernel, *d_output_cudnn, *d_output_naive;CHECK_CUDA(cudaMalloc(&d_input, inputSize * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_kernel, kernelElements * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_output_cudnn, outputSize * sizeof(float)));CHECK_CUDA(cudaMalloc(&d_output_naive, outputSize * sizeof(float)));// Copy data to deviceCHECK_CUDA(cudaMemcpy(d_input, h_input, inputSize * sizeof(float), cudaMemcpyHostToDevice));CHECK_CUDA(cudaMemcpy(d_kernel, h_kernel, kernelElements * sizeof(float), cudaMemcpyHostToDevice));// cuDNN setupcudnnHandle_t cudnn;CHECK_CUDNN(cudnnCreate(&cudnn));cudnnTensorDescriptor_t inputDesc, outputDesc;cudnnFilterDescriptor_t kernelDesc;cudnnConvolutionDescriptor_t convDesc;CHECK_CUDNN(cudnnCreateTensorDescriptor(&inputDesc));CHECK_CUDNN(cudnnCreateTensorDescriptor(&outputDesc));CHECK_CUDNN(cudnnCreateFilterDescriptor(&kernelDesc));CHECK_CUDNN(cudnnCreateConvolutionDescriptor(&convDesc));CHECK_CUDNN(cudnnSetTensor4dDescriptor(inputDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, batchSize, inChannels,height, width));CHECK_CUDNN(cudnnSetTensor4dDescriptor(outputDesc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, batchSize, outChannels,height, width));CHECK_CUDNN(cudnnSetFilter4dDescriptor(kernelDesc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, outChannels, inChannels,kernelSize, kernelSize));CHECK_CUDNN(cudnnSetConvolution2dDescriptor(convDesc, kernelSize / 2, kernelSize / 2, 1, 1, 1, 1,CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT));// Find the fastest cuDNN algorithmint requestedAlgoCount = CUDNN_CONVOLUTION_FWD_ALGO_COUNT;int returnedAlgoCount;cudnnConvolutionFwdAlgoPerf_t perfResults[CUDNN_CONVOLUTION_FWD_ALGO_COUNT];CHECK_CUDNN(cudnnGetConvolutionForwardAlgorithm_v7(cudnn, inputDesc, kernelDesc, convDesc, outputDesc,requestedAlgoCount, &returnedAlgoCount, perfResults));cudnnConvolutionFwdAlgo_t algo = CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM; // Default algorithmsize_t workspaceSize;CHECK_CUDNN(cudnnGetConvolutionForwardWorkspaceSize(cudnn, inputDesc, kernelDesc, convDesc, outputDesc, algo,&workspaceSize));void* d_workspace;CHECK_CUDA(cudaMalloc(&d_workspace, workspaceSize));// Define grid and block sizes for the naive kerneldim3 blockSize(16, 16);dim3 gridSize((width + blockSize.x - 1) / blockSize.x, (height + blockSize.y - 1) / blockSize.y,outChannels * batchSize);// Warmup and benchmark runsconst int warmupRuns = 5;const int benchmarkRuns = 20;float totalTime_cudnn = 0.0f;float totalTime_naive = 0.0f;float alpha = 1.0f, beta = 0.0f;// Warmup runsfor (int i = 0; i < warmupRuns; i++) {CHECK_CUDNN(cudnnConvolutionForward(cudnn, &alpha, inputDesc, d_input, kernelDesc, d_kernel, convDesc, algo,d_workspace, workspaceSize, &beta, outputDesc, d_output_cudnn));naiveConv2d<<<gridSize, blockSize>>>(d_input, d_kernel, d_output_naive, width, height, inChannels, outChannels,kernelSize, batchSize);CHECK_CUDA(cudaDeviceSynchronize());}// Benchmark runscudaEvent_t start, stop;CHECK_CUDA(cudaEventCreate(&start));CHECK_CUDA(cudaEventCreate(&stop));for (int i = 0; i < benchmarkRuns; i++) {// cuDNN benchmarkCHECK_CUDA(cudaEventRecord(start));CHECK_CUDNN(cudnnConvolutionForward(cudnn, &alpha, inputDesc, d_input, kernelDesc, d_kernel, convDesc, algo,d_workspace, workspaceSize, &beta, outputDesc, d_output_cudnn));CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));float milliseconds = 0;CHECK_CUDA(cudaEventElapsedTime(&milliseconds, start, stop));totalTime_cudnn += milliseconds;// Naive kernel benchmarkCHECK_CUDA(cudaEventRecord(start));naiveConv2d<<<gridSize, blockSize>>>(d_input, d_kernel, d_output_naive, width, height, inChannels, outChannels,kernelSize, batchSize);CHECK_CUDA(cudaEventRecord(stop));CHECK_CUDA(cudaEventSynchronize(stop));CHECK_CUDA(cudaEventElapsedTime(&milliseconds, start, stop));totalTime_naive += milliseconds;}// Calculate average timesfloat avgTime_cudnn = totalTime_cudnn / benchmarkRuns;float avgTime_naive = totalTime_naive / benchmarkRuns;printf("cuDNN average time: %f ms\n", avgTime_cudnn);printf("Naive kernel average time: %f ms\n", avgTime_naive);// Copy results back to hostCHECK_CUDA(cudaMemcpy(h_output_cudnn, d_output_cudnn, outputSize * sizeof(float), cudaMemcpyDeviceToHost));CHECK_CUDA(cudaMemcpy(h_output_naive, d_output_naive, outputSize * sizeof(float), cudaMemcpyDeviceToHost));// Compare resultsfloat maxDiff = 0.0f;for (int i = 0; i < outputSize; i++) {float diff = fabs(h_output_cudnn[i] - h_output_naive[i]);if (diff > maxDiff) maxDiff = diff;}printf("Max difference between cuDNN and naive kernel: %e\n", maxDiff);// Clean upCHECK_CUDNN(cudnnDestroyTensorDescriptor(inputDesc));CHECK_CUDNN(cudnnDestroyTensorDescriptor(outputDesc));CHECK_CUDNN(cudnnDestroyFilterDescriptor(kernelDesc));CHECK_CUDNN(cudnnDestroyConvolutionDescriptor(convDesc));CHECK_CUDNN(cudnnDestroy(cudnn));CHECK_CUDA(cudaFree(d_input));CHECK_CUDA(cudaFree(d_kernel));CHECK_CUDA(cudaFree(d_output_cudnn));CHECK_CUDA(cudaFree(d_output_naive));CHECK_CUDA(cudaFree(d_workspace));CHECK_CUDA(cudaEventDestroy(start));CHECK_CUDA(cudaEventDestroy(stop));free(h_input);free(h_kernel);free(h_output_cudnn);free(h_output_naive);return 0;
}
输出:
Image size: 224x224x32
Kernel size: 11x11x32x64
Batch size: 4
cuDNN average time: 19.572138 ms
Naive kernel average time: 107.169754 ms
Max difference between cuDNN and naive kernel: 0.000000e+00
可以看到在我的机器上,cuDNN实现的大型卷积操作的速度是简单的CUDA卷积的5倍左右。
Larger Rigs or Datacenters(大型工作站 vs 数据中心)
这里简单补充一下在大型工作站和数据中心上CUDA的一些相关知识
cuBLAS-mp vs NCCL vs MIG(multi instance GPU):关键区别和使用场景
这三种技术各自有不同的应用场景,但它们都在分布式计算和高性能计算中优化GPU性能方面扮演重要角色。以下是对每种技术的详细分析:
1. cuBLAS-mp (多进程cuBLAS)
定义:
cuBLAS-mp(多进程cuBLAS)是NVIDIA提供的一个高性能、GPU加速的线性代数库,专为在单节点(一个物理机器)中进行多GPU计算而设计。
使用场景:
- 单节点、多GPU计算:当一个模型过大,无法适应单个GPU时,cuBLAS-mp可以将工作负载分配到同一台机器上的多个GPU。这种情况通常发生在深度学习模型的大小超过单个GPU的显存时。
- 矩阵乘法(Matmul):cuBLAS-mp优化了矩阵乘法操作,这是训练深度学习模型中的关键操作,适用于将多个GPU上的计算任务分配并同步。
关键特点:
- 高性能线性代数计算:优化了矩阵运算(例如矩阵-矩阵乘法,
GEMM
)的GPU操作。 - 多进程支持:允许多个进程共享单节点上的GPU资源。
- 单节点扩展:适用于在单台机器上进行大规模的张量计算,尤其是在模型无法完全放入单个GPU时。
使用案例:
- 大模型训练:例如训练像GPT-5这样的大型模型时,由于模型过大无法放入单个GPU的显存,cuBLAS-mp可以将计算任务分配到多个GPU上。
2. NCCL (NVIDIA Collective Communications Library)
定义:
NCCL(“nickel”)是NVIDIA提供的一个用于分布式集群通信的库,主要用于在多个机器或节点之间进行GPU之间的高效通信。
使用场景:
- 分布式训练:NCCL对于大规模的分布式深度学习训练至关重要,尤其是涉及多个节点(每个节点上有多个GPU)的场景。它负责GPU和节点之间的数据通信。
- 集体通信:包括操作如All-Reduce、Broadcast、Gather和Scatter,这些操作是并行化训练和模型权重同步的基础。
关键特点:
- 集群级通信:NCCL负责处理分布式训练中的通信部分,而cuBLAS-mp负责GPU端的计算任务。
- 高效的集体操作:优化了在多个节点或GPU之间共享和同步数据的操作。
- 与PyTorch的集成:在PyTorch中,分布式数据并行(DDP) 是基于NCCL的,它支持跨多个GPU和节点的高效模型并行训练。
使用案例:
- 多节点集群训练:如果你在多个节点上训练一个大模型(例如GPT-5),NCCL会负责在不同机器上的GPU之间进行梯度和模型更新的通信,确保分布式训练的高效进行。
一些对你可能有用的链接:
https://pytorch.org/tutorials/intermediate/ddp_tutorial.html
https://www.youtube.com/watch?v=T22e3fgit-A&ab_channel=CUDAMODE
https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-gpu-memory
https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/overview.html
https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/api.html
3. MIG (Multi-Instance GPU)
定义:
MIG是一种将大GPU划分为多个小型独立GPU实例的技术。每个实例都有自己的内存和计算资源,相互隔离,适用于将一个大的GPU资源分配给多个不同的用户或应用。
使用场景:
- 数据中心优化:MIG特别适合在数据中心的环境中使用,尤其是在多用户或多个应用共享同一个GPU时,能够提高GPU资源的利用率。
- 提高资源利用率:通过将单个GPU划分为多个独立的小GPU实例,MIG能确保每个工作负载都能获得足够的GPU资源。
关键特点:
- GPU分割:MIG将一个大的GPU(如NVIDIA A100或H100)划分为多个小的实例,每个实例都具有独立的内存、计算能力和内存带宽。
- 资源隔离:每个MIG实例都是独立的,具有完全的资源隔离,避免了不同任务之间的干扰。
- 数据中心扩展性:MIG在云环境或数据中心中非常有用,可以让多个任务共享同一台物理GPU,提高GPU的资源利用率。
使用案例:
- 多个独立任务:例如,训练多个较小的模型,而不是将整个GPU资源分配给单一任务。通过MIG,可以将一个大GPU划分为多个独立的小GPU实例,从而最大化资源使用。
比较总结:
特性 | cuBLAS-mp | NCCL | MIG |
---|---|---|---|
范围 | 单节点,多GPU张量操作 | 分布式集群中多节点间的通信 | 将单个GPU分割成多个独立的小GPU实例 |
主要用途 | 适用于无法在单个GPU上运行的大模型 | 在多个节点的GPU之间进行同步和数据分发 | 在数据中心环境中提高GPU资源的利用率 |
通信 | 无(专注于计算) | 集体操作(All-reduce, Broadcast, Gather, Scatter) | 无(专注于GPU资源分配) |
关键操作 | 矩阵乘法、张量计算 | 集体通信、模型并行训练 | 独立GPU实例化 |
最适用场景 | 单节点内多GPU训练 | 多节点分布式训练 | 数据中心环境中资源优化 |
总结:
- cuBLAS-mp 适用于单节点的多GPU计算,特别是在模型无法完全放入单个GPU时,适合大规模的张量计算任务。
- NCCL 主要用于分布式训练中的多节点集群通信,负责GPU和节点间的数据同步和集体操作。
- MIG 则是将单个GPU划分为多个小的独立实例,可以提高GPU资源的利用率,尤其是在数据中心或云环境中,适合同时处理多个较小的任务。
这三者各自针对不同的应用场景和需求,优化了GPU在大规模分布式计算中的性能。
参考:https://github.com/Infatoshi/cuda-course/tree/master