当前位置：首页 > news >正文

CUDA从入门到精通（六）——CUDA编程模型（二）

news 来源：原创 2025/6/10 0:06:54

1. 核函数类型限定符

CUDA 核函数的常用函数类型限定符及其相关信息的表格：

限定符	执行端	调用方式	备注
`__global__`	设备端（GPU）	从主机代码使用 `<<<...>>>` 调用核函数	用于声明核函数，在 GPU 上执行。只能从主机代码调用。通常没有返回值。
`__device__`	设备端（GPU）	只能从设备代码（核函数或其他设备函数）调用	用于声明设备函数，只能在 GPU 上执行，不能从主机代码调用。
`__host__`	主机端（CPU）	只能从主机代码调用	用于声明主机函数，必须在 CPU 上执行，不能从设备代码调用。
`__host__ __device__`	主机端（CPU）和设备端（GPU）	可以从主机或设备代码调用	该函数可以在主机和设备上执行，适用于需要兼容主机和设备的通用函数。
`__launch_bounds__`	设备端（GPU）	用于核函数声明	用于提示编译器优化线程块的大小和寄存器的使用。
`__restrict__`	设备端（GPU）	用于指针参数声明	用于声明指针，告诉编译器该指针所指向的内存不会被其他指针修改，有助于性能优化。

详细说明：

__global__：
- 核函数限定符，表示该函数是由 GPU 上的线程执行的。
- 从主机代码中调用，使用 <<<...>>> 语法进行配置。
__device__：
- 用于声明设备函数，函数仅在 GPU 代码中执行。
- 只能被核函数或其他设备函数调用，无法从主机代码直接调用。
__host__：
- 用于声明主机函数，表示该函数只能在 CPU 上执行。
- 只能从主机代码中调用，不能从设备代码中调用。
__host__ __device__：
- 允许函数在主机和设备上都执行，兼容两端的调用。
- 适用于那些通用的函数，它们可以同时在主机和设备上执行。
__launch_bounds__：
- 用于优化核函数的执行，提供线程块大小和寄存器使用的提示。
- 提示编译器对核函数的线程调度进行优化。
__restrict__：
- 用于指针类型，告知编译器该指针所指向的内存不会被其他指针修改。
- 允许编译器进行更有效的优化，减少内存访问冲突。

在 CUDA 编程中，核函数（kernel functions）是由 GPU 上的线程执行的函数。尽管 CUDA 提供了强大的并行计算能力，但在使用核函数时也存在一些限制。以下是一些主要的限制：

2. 核函数限制

1. 返回值限制

核函数不能返回值：核函数的返回类型必须是 void，因为它们不能直接返回值。所有的结果必须通过指针或引用传递回主机。

2. 线程和块的限制

最大线程数：每个线程块的最大线程数通常为 1024（具体取决于 GPU 架构）。这意味着在一个线程块中，您不能创建超过这个数量的线程。
最大线程块数：每个网格的最大线程块数也有限制，具体取决于 GPU 的计算能力。
线程块维度：线程块的维度（即线程的数量）通常限制为 1D、2D 或 3D，且每个维度的大小也有上限。

3. 内存限制

共享内存限制：每个线程块可以使用的共享内存量是有限的，通常为 48KB（具体取决于 GPU 架构）。如果需要更多的共享内存，可能需要调整线程块的大小。
全局内存访问延迟：虽然全局内存可以存储大量数据，但访问全局内存的延迟相对较高。频繁的全局内存访问可能会导致性能下降。

4. 设备函数限制

设备函数不能被主机代码调用：设备函数（使用 __device__ 限定符声明的函数）只能在设备代码中调用，不能从主机代码直接调用。

5. 递归限制

不支持递归：CUDA 核函数不支持递归调用。所有的函数调用必须是非递归的。

6. 线程同步限制

线程同步：在同一个线程块内，可以使用 __syncthreads() 进行线程同步，但不能跨线程块进行同步。跨块的同步需要其他机制，如原子操作或全局内存的协调。

7. 设备属性限制

设备属性：不同的 GPU 设备具有不同的计算能力和资源限制。开发者需要根据目标设备的属性进行优化。

8. 设备内存分配限制

动态内存分配：在核函数中使用动态内存分配（如 malloc）是有限制的，可能会导致性能下降。动态分配的内存也可能会导致内存碎片。

9. 计算能力限制

计算能力：不同的 GPU 具有不同的计算能力（如 CUDA 计算能力 2.0、3.0、5.0 等），某些功能和特性可能在较低的计算能力下不可用。

10. 设备和主机之间的数据传输

数据传输开销：在主机和设备之间传输数据（如从主机到设备的内存拷贝）会引入开销，频繁的数据传输会影响性能。

3.核函数计时

在 CUDA 编程中，计时核函数的执行时间是评估性能的重要步骤。可以使用 CUDA 提供的事件（events）来精确测量核函数的执行时间。以下是实现核函数计时的步骤和示例代码。

1. 使用 CUDA 事件计时

CUDA 事件是用于测量时间的高精度工具。通过创建事件并在核函数执行前后记录时间，可以计算出核函数的执行时间。

创建事件：使用 cudaEventCreate() 创建事件。
记录事件：在核函数调用前后使用 cudaEventRecord() 记录事件。
计算时间：使用 cudaEventElapsedTime() 计算两个事件之间的时间差。
清理事件：使用 cudaEventDestroy() 清理事件。

#include <iostream>
#include <cuda_runtime.h>__global__ void kernel_function() {// 核函数代码int idx = threadIdx.x + blockIdx.x * blockDim.x;// 进行一些计算 if (idx < 1000) {// 示例计算float value = idx * 2.0f;}
}int main() {// 创建 CUDA 事件cudaEvent_t start, stop;cudaEventCreate(&start);cudaEventCreate(&stop);// 设置线程块和网格大小int blockSize = 256;int numBlocks = (1000 + blockSize - 1) / blockSize;// 记录开始事件cudaEventRecord(start);// 调用核函数 kernel_function<<<numBlocks, blockSize>>>();// 记录结束事件 cudaEventRecord(stop);// 等待事件完成cudaEventSynchronize(stop);// 计算时间 float milliseconds = 0;cudaEventElapsedTime(&milliseconds, start, stop);// 输出执行时间std::cout << "Kernel execution time: " << milliseconds << " ms" << std::endl;// 清理事件cudaEventDestroy(start);cudaEventDestroy(stop);return 0;
}

核函数：kernel_function 是一个简单的核函数，执行一些计算。
事件创建：使用 cudaEventCreate() 创建 start 和 stop 事件。
记录事件：
- 在调用核函数之前，使用 cudaEventRecord(start) 记录开始时间。
- 在核函数调用之后，使用 cudaEventRecord(stop) 记录结束时间。
同步事件：使用 cudaEventSynchronize(stop) 确保核函数执行完成。
计算时间：使用 cudaEventElapsedTime(&milliseconds, start, stop) 计算两个事件之间的时间差，单位为毫秒。
输出时间：输出核函数的执行时间。
清理事件：使用 cudaEventDestroy() 清理事件，释放资源。

CUDA 设备同步：在记录结束事件后，确保使用 cudaEventSynchronize() 等待核函数完成，以获得准确的时间。
错误检查：在实际应用中，建议在每个 CUDA API 调用后添加错误检查，以确保没有发生错误。
多次测量：为了获得更稳定的性能数据，可以多次运行核函数并计算平均时间。

除了使用 CUDA 提供的 硬件性能计数器（如 CPI计时器）外，您还可以基于 CPU计时器 和 nvprof 工具进行核函数执行时间的计时。下面我会详细介绍这两种方法。

2. 基于 CPU 计时器计时

虽然 CUDA 核函数运行在 GPU 上，但我们仍然可以使用 CPU计时器 来测量 CUDA 程序的执行时间，尤其是对核函数调用前后以及数据传输的时间进行测量。常用的 CPU 计时器有 std::chrono 和 clock()，它们可以用于测量 CPU 时间。

使用 std::chrono 计时（C++11 或更高版本）

std::chrono 是 C++11 引入的时间库，提供高精度计时器，可以用来精确地测量 CUDA 核函数的执行时间。std::chrono::high_resolution_clock 是一个高精度时钟，它提供了较高的时间分辨率。

#include <iostream>
#include <chrono>
#include <cuda_runtime.h>__global__ void kernel_function() {int idx = threadIdx.x + blockIdx.x * blockDim.x;// 核函数中进行一些计算if (idx < 1000) {float value = idx * 2.0f;}
}int main() {// 使用 std::chrono 高精度计时器auto start = std::chrono::high_resolution_clock::now();// 设置线程块和网格大小int blockSize = 256;int numBlocks = (1000 + blockSize - 1) / blockSize;// 调用核函数kernel_function<<<numBlocks, blockSize>>>();// 等待核函数执行完毕cudaDeviceSynchronize();// 记录结束时间auto end = std::chrono::high_resolution_clock::now();// 计算执行时间std::chrono::duration<float> duration = end - start;std::cout << "Kernel execution time: " << duration.count() << " seconds." << std::endl;return 0;
}

std::chrono::high_resolution_clock::now()：用于获取当前的时间戳，具有较高的时间精度。
cudaDeviceSynchronize()：确保核函数执行完毕后再计算时间。
duration.count()：获取执行的时间，单位是秒。

这种方法适用于需要在 主机端（CPU）计时 CUDA 核函数的场景，但需要注意的是，它只能计时核函数的总执行时间，不能提供 GPU 上详细的硬件性能数据。

3. 使用 `nvprof` 计时

nvprof 是 NVIDIA Profiler，一个命令行工具，能够提供丰富的性能分析数据，帮助你了解 CUDA 程序的执行情况，包括内存传输、核函数执行时间、硬件性能计数等。使用 nvprof，你可以轻松地获取核函数的执行时间和其他性能指标。

使用 nvprof 计时

nvprof 可以用来记录 CUDA 核函数的执行时间、内存传输情况以及硬件级别的性能指标（如执行周期、指令数等）。它是 NVIDIA Profiler 工具的一部分，非常适用于性能分析。

编译 CUDA 程序：
首先，编译您的 CUDA 程序，确保使用了调试信息（-g 标志）。例如：
```
nvcc -g -G -o my_program my_program.cu
```
运行 nvprof：
使用 nvprof 命令运行您的 CUDA 程序并获取核函数执行时间：
```
nvprof --metrics time_elapsed ./my_program
```
这将显示核函数的执行时间（单位为微秒）。
获取更多性能指标：
nvprof 还可以显示有关硬件资源的其他信息，如执行周期数、指令数等。您可以通过 --metrics 选项获取多个指标：
```
nvprof --metrics sm__cycles_elapsed.avg,sm__inst_executed.avg ./my_program
```
- sm__cycles_elapsed.avg：执行的平均周期数。
- sm__inst_executed.avg：执行的平均指令数。
获取具体核函数的时间：
如果只关注某个特定的核函数，您可以使用以下命令：
```
nvprof --kernel <kernel_name> --metrics time_elapsed ./my_program
```
其中 <kernel_name> 替换为您程序中核函数的名称。

==12345== Profiling application: ./my_program
==12345== Metrics result:
==12345==   Metric 'time_elapsed' is 1500.0 ms
==12345==   Metric 'sm__cycles_elapsed.avg' is 2000000
==12345==   Metric 'sm__inst_executed.avg' is 1000000

4. 计算 CPI

如前所述，CPI（Cycles Per Instruction）可以通过以下公式计算：

$sm__cycles_elapsed.avg sm__inst_executed.avg \text{CPI} = \frac{\text{sm\_\_cycles\_elapsed.avg}}{\text{sm\_\_inst\_executed.avg}}$

在上面的例子中：

sm__cycles_elapsed.avg = 2000000
sm__inst_executed.avg = 1000000

所以：
$\text{CPI} = \frac{2000000}{1000000} = 2.0$

这意味着每条指令在该核函0数执行中平均消耗 2 个周期。

nvprof 提供了详细的性能数据，包括内存传输、核函数执行时间、硬件资源使用等。
nvprof 可以用于查看整个程序的性能，方便发现瓶颈。
nvprof 主要是一个命令行工具，不适合与程序中的计时逻辑紧密结合。
它通常用来进行后期的分析，而不是实时计时。

方法	优点	缺点
基于 CPU 计时器（如 std::chrono）	简单易用，适用于对 CUDA 核函数进行快速计时	只能测量核函数的总执行时间，无法提供硬件级别的性能数据
基于 `nvprof` 工具计时	提供详细的性能分析数据，支持多种硬件级别的计数器指标（如执行周期、指令数等）	主要是后期分析工具，不适合嵌入程序中实时计时，且有额外的运行开销

选择哪种计时方式取决于您的需求：

CPU计时器 更适用于简单的性能测量和快速开发。
nvprof 适合需要深入了解程序性能和瓶颈的情况，特别是在大规模程序调优时。
在 CUDA 编程中，网格（grid）和线程块（block）的配置对性能有显著影响。不同的网格和块数量会导致不同的性能表现，主要原因包括以下几个方面：

4. 不同的线程数量和块数拥有不同的性能

1. 资源利用率

GPU 资源限制：每个 GPU 有其特定的资源限制，包括每个线程块的最大线程数、共享内存、寄存器等。选择合适的线程块大小可以确保 GPU 资源的高效利用。
并行度：如果线程块数量过少，可能无法充分利用 GPU 的并行计算能力。相反，如果线程块数量过多，可能会导致资源竞争，降低性能。

2. 线程调度

线程块调度：GPU 使用线程调度器来管理线程块的执行。线程块的数量和大小会影响调度的效率。较小的线程块可能导致调度开销增加，而较大的线程块可能会导致资源浪费。
活跃线程数：为了保持 GPU 的高效运行，通常需要有足够数量的活跃线程。如果线程块数量不足，可能会导致 GPU 处于空闲状态，降低整体性能。

3. 内存访问模式

内存访问效率：线程块的配置会影响内存访问模式。合理的线程块大小可以提高内存访问的局部性，减少全局内存访问的延迟。
共享内存的使用：如果线程块的大小适当，可以利用共享内存来减少全局内存访问，从而提高性能。过小的线程块可能无法有效利用共享内存。

4. 计算与内存传输的平衡

计算与内存传输的比例：在 CUDA 程序中，计算和内存传输是两个主要的性能瓶颈。合理配置网格和块的数量可以帮助平衡计算和内存传输的比例，减少内存传输的影响。
内存带宽：如果线程块数量过多，可能会导致内存带宽的竞争，影响性能。适当的块数量可以帮助优化内存带宽的使用。

5. 线程块的大小

线程块的维度：线程块的维度（1D、2D、3D）也会影响性能。某些算法在特定维度上表现更好，合理选择线程块的维度可以提高性能。
线程块的大小：较大的线程块可能会导致更多的寄存器和共享内存的使用，影响其他线程块的调度。较小的线程块可能会导致调度开销增加。

6. 设备特性

GPU 架构：不同的 GPU 架构对线程块和网格的支持不同。某些架构可能对特定的线程块大小和数量有更好的优化。
计算能力：GPU 的计算能力（如 CUDA 计算能力）会影响可用的资源和性能表现。了解目标设备的特性可以帮助优化网格和块的配置。

7. 负载均衡

负载均衡：合理的网格和块配置可以确保每个线程块的工作量相对均匀，避免某些线程块过载而其他线程块空闲的情况。负载不均衡会导致性能下降。

不同的网格和块数量会影响 CUDA 程序的性能，主要是因为它们影响了资源利用率、线程调度、内存访问模式、计算与内存传输的平衡、线程块的大小、设备特性和负载均衡等因素。为了获得最佳性能，开发者需要根据具体的应用场景和目标 GPU 的特性，合理配置网格和块的数量。通常，进行性能测试和基准测试是找到最佳配置的有效方法。

5. 设备管理

在 CUDA 编程中，查询 GPU 设备信息、选择最佳 GPU 设备并进行设备管理是性能优化的重要步骤。以下是如何使用不同的 API 查询设备信息，选择最佳 GPU，使用 nvidia-smi 查询 GPU 信息以及在运行时设置设备的详细方法。

1. 使用 CUDA API 查询设备信息

CUDA 提供了多个 API 函数来查询 GPU 设备的各种信息，如设备数量、属性、内存、计算能力等。

1.1 查询设备数量

#include <iostream>
#include <cuda_runtime.h>int main() {int deviceCount;cudaError_t err = cudaGetDeviceCount(&deviceCount);if (err != cudaSuccess) {std::cerr << "Error getting device count: " << cudaGetErrorString(err) << std::endl;return -1;}std::cout << "Number of CUDA devices: " << deviceCount << std::endl;return 0;
}

cudaGetDeviceCount(&deviceCount)：返回可用的 CUDA 设备数量。

1.2 获取设备属性

每个 CUDA 设备都有一个 cudaDeviceProp 结构体，包含设备的各种信息。例如，内存大小、计算能力、每个线程块的最大线程数等。

#include <iostream>
#include <cuda_runtime.h>void printDeviceProperties(int deviceId) {cudaDeviceProp prop;cudaGetDeviceProperties(&prop, deviceId);std::cout << "Device " << deviceId << ": " << prop.name << std::endl;std::cout << "  Total Global Memory: " << prop.totalGlobalMem / (1024 * 1024) << " MB" << std::endl;std::cout << "  Shared Memory per Block: " << prop.sharedMemPerBlock / 1024 << " KB" << std::endl;std::cout << "  Max Threads per Block: " << prop.maxThreadsPerBlock << std::endl;std::cout << "  Compute Capability: " << prop.major << "." << prop.minor << std::endl;
}int main() {int deviceCount;cudaGetDeviceCount(&deviceCount);for (int i = 0; i < deviceCount; ++i) {printDeviceProperties(i);}return 0;
}

cudaGetDeviceProperties(&prop, deviceId)：查询指定设备的属性，存储在 cudaDeviceProp 结构体中。
prop.name：设备名称。
prop.totalGlobalMem：设备的全局内存总量（以字节为单位）。
prop.computeCapability：计算能力（如 6.1 表示 CUDA 6.1）。

1.3 获取当前设备

使用 cudaGetDevice() 可以获取当前选择的设备。

int currentDevice;
cudaGetDevice(&currentDevice);
std::cout << "Current device is: " << currentDevice << std::endl;

1.4 设置设备

使用 cudaSetDevice() 可以在程序中选择要使用的 GPU 设备。

int deviceId = 1;  // 假设选择设备 1
cudaSetDevice(deviceId);

2. 选择最佳 GPU 设备

选择最佳 GPU 设备通常基于多个因素，如内存大小、计算能力、使用的应用场景等。你可以选择具有最大内存或最高计算能力的设备。

例如，以下代码选择具有最大全局内存的设备作为最佳设备：

int bestDevice = 0;
size_t maxMemory = 0;
int deviceCount;
cudaGetDeviceCount(&deviceCount);for (int i = 0; i < deviceCount; ++i) {cudaDeviceProp prop;cudaGetDeviceProperties(&prop, i);if (prop.totalGlobalMem > maxMemory) {maxMemory = prop.totalGlobalMem;bestDevice = i;}
}std::cout << "Best device is: " << bestDevice << " with " << maxMemory / (1024 * 1024) << " MB memory." << std::endl;// 选择最佳设备
cudaSetDevice(bestDevice);

3. 使用 `nvidia-smi` 查询 GPU 信息

nvidia-smi 是 NVIDIA 提供的一个命令行工具，用于查询 GPU 状态和管理 GPU 资源。你可以通过 nvidia-smi 查看 GPU 的详细信息，如 GPU 使用情况、温度、内存使用量等。

3.1 查询 GPU 状态

在命令行中使用 nvidia-smi 查询 GPU 状态：

nvidia-smi

输出示例：

+-----------------------------------------------------------------------------+
| NVIDIA-SMI 460.32.03    Driver Version: 460.32.03    CUDA Version: 11.2     |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap| Memory-Usage | GPU-Util  Compute M. |
|===============================+======================+======================|
|   0  Tesla K80           Off  | 00000000:00:1E.0 Off |                    0 |
| N/A   39C    P8    29W / 149W |    0MiB / 11441MiB |      0%      Default |
+-------------------------------+----------------------+----------------------+

3.2 查询特定 GPU 的信息

你也可以查询特定 GPU 的信息，例如，查询设备 ID 为 0 的 GPU：

nvidia-smi -i 0

3.3 查看 GPU 内存和使用情况

要查看 GPU 的内存使用情况：

nvidia-smi --query-gpu=memory.used,memory.free,memory.total --format=csv