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)可以通过以下公式计算:
CPI = sm__cycles_elapsed.avg sm__inst_executed.avg \text{CPI} = \frac{\text{sm\_\_cycles\_elapsed.avg}}{\text{sm\_\_inst\_executed.avg}} CPI=sm__inst_executed.avgsm__cycles_elapsed.avg
在上面的例子中:
sm__cycles_elapsed.avg
= 2000000sm__inst_executed.avg
= 1000000
所以:
CPI = 2000000 1000000 = 2.0 \text{CPI} = \frac{2000000}{1000000} = 2.0 CPI=10000002000000=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(¤tDevice);
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
输出示例:
memory.used [MiB], memory.free [MiB], memory.total [MiB]
0 MiB, 11264 MiB, 11441 MiB
3.4 使用 nvidia-smi
执行任务
您还可以使用 nvidia-smi
在命令行中启动或停止 GPU 任务。例如,查看 GPU 使用情况并限制其计算任务:
nvidia-smi -i 0 --persistence-mode=1
4. 运行时设置设备
CUDA 允许在运行时动态选择 GPU 设备。可以通过以下步骤在应用中进行设备选择:
4.1 获取设备数量
通过 cudaGetDeviceCount()
获取当前系统中的可用 GPU 数量。
4.2 根据设备特性选择设备
根据设备的性能指标(如计算能力、内存大小等),选择最佳的 GPU。
4.3 设置设备
通过 cudaSetDevice(deviceId)
选择指定的 GPU 设备进行计算。
4.4 同步设备
如果您的程序在多个设备上并行执行,您可以使用 cudaDeviceSynchronize()
来同步设备的执行,确保当前设备的所有任务完成后才进行下一步操作。
功能 | CUDA API | 命令行工具 (nvidia-smi ) |
---|---|---|
查询可用设备数量 | cudaGetDeviceCount(&deviceCount) | N/A |
查询设备属性 | cudaGetDeviceProperties(&prop, deviceId) | N/A |
获取当前设备 ID | cudaGetDevice(¤tDevice) | N/A |
选择设备 | cudaSetDevice(deviceId) | N/A |
查询设备内存和使用情况 | N/A | nvidia-smi --query-gpu=memory.used,memory.free,memory.total --format=csv |
获取设备信息 | cudaGetDeviceProperties() | nvidia-smi |
设备信息过滤 | cudaDeviceGetAttribute() (如最大线程数、内存等) | nvidia-smi -i <device_id> |
通过结合使用 CUDA API 和 nvidia-smi
,可以灵活地查询和选择 GPU 设备,在程序运行时进行设备管理和优化。这有助于提高程序的性能,尤其在多 GPU 系统中。