学网站开发培训机构_360浏览器最新版本下载安装_企业文化理念_深圳推广网络

【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