【CUDA】cuDNN：加速深度学习的核心库
【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, // 数据类型：float
batch_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_GEMMCUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMMCUDNN_CONVOLUTION_FWD_ALGO_FFTCUDNN_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 1
data[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 outperform
const int batch_size = 256; // NCHW format
const int channels = 32;
const int height = 224;
const int width = 224;
const int tensor_size = batch_size * channels * height * width;
// Allocate host memory
float *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 memory
float *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 device
CHECK_CUDA(cudaMemcpy(d_input, h_input, tensor_size * sizeof(float), cudaMemcpyHostToDevice));
// Create CUDA events for timing
cudaEvent_t start, stop;
CHECK_CUDA(cudaEventCreate(&start));
CHECK_CUDA(cudaEventCreate(&stop));
// Warmup and benchmark parameters
const int num_warmup = 10;
const int num_benchmark = 100;
float naive_times[num_benchmark];
float cudnn_times[num_benchmark];
// Naive CUDA kernel
dim3 block(256);
dim3 grid((tensor_size + block.x - 1) / block.x);
// Warmup runs for naive kernel
for (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 cuDNN
for (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 cuDNN
for (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 times
float 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 host
CHECK_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 verification
for (int i = 0; i < tensor_size; ++i) {
h_output_cpu[i] = CpuTanh(h_input[i]);
}
// Verify results
bool naive_correct = VerifyResults(h_output_cpu, h_output_naive, tensor_size);
bool cudnn_correct = VerifyResults(h_output_cpu, h_output_cudnn, tensor_size);
// Print results
printf("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 up
CHECK_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 output
const 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 memory
float* 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 values
float 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 memory
float *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 device
CHECK_CUDA(cudaMemcpy(d_input, h_input, input_size * sizeof(float), cudaMemcpyHostToDevice));
CHECK_CUDA(cudaMemcpy(d_kernel, h_kernel, kernel_elements * sizeof(float), cudaMemcpyHostToDevice));
// cuDNN setup
cudnnHandle_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 kernel
dim3 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 runs
const 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 runs
for (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 benchmark
CHECK_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 host
CHECK_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 results
float 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 output
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_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 PyTorch
printf("\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 up
CHECK_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 output
const 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 memory
float* 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 values
srand(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 memory
float *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 device
CHECK_CUDA(cudaMemcpy(d_input, h_input, inputSize * sizeof(float), cudaMemcpyHostToDevice));
CHECK_CUDA(cudaMemcpy(d_kernel, h_kernel, kernelElements * sizeof(float), cudaMemcpyHostToDevice));
// cuDNN setup
cudnnHandle_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 algorithm
int 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 algorithm
size_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 kernel
dim3 blockSize(16, 16);
dim3 gridSize((width + blockSize.x - 1) / blockSize.x, (height + blockSize.y - 1) / blockSize.y,
outChannels * batchSize);
// Warmup and benchmark runs
const 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 runs
for (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 runs
cudaEvent_t start, stop;
CHECK_CUDA(cudaEventCreate(&start));
CHECK_CUDA(cudaEventCreate(&stop));
for (int i = 0; i < benchmarkRuns; i++) {
// cuDNN benchmark
CHECK_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 benchmark
CHECK_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 times
float 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 host
CHECK_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 results
float 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 up
CHECK_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