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
- ==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
如前所述，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 = 2000000
- sm__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;
使用 cudaSetDevice() 可以在程序中选择要使用的 GPU 装备。
- int deviceId = 1; // 假设选择设备 1
- cudaSetDevice(deviceId);
选择最佳 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);
查询 GPU 信息
nvidia-smi
是 NVIDIA 提供的一个命令行工具，用于查询 GPU 状态和管理 GPU 资源。你可以通过 nvidia-smi
查看 GPU 的具体信息，如 GPU 使用情况、温度、内存使用量等。
3.1 查询 GPU 状态
在命令行中使用 nvidia-smi
查询 GPU 状态：
输出示例：
- +-----------------------------------------------------------------------------+
- | 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 |
- +-------------------------------+----------------------+----------------------+
你也可以查询特定 GPU 的信息，比方，查询装备 ID 为 0 的 GPU：
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
执行任务
您还可以使用 nvidia-smi
在命令行中启动或制止 GPU 任务。比方，查看 GPU 使用情况并限制其计算任务：
- nvidia-smi
- -i 0 --persistence-mode=1
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获取当前装备 IDcudaGetDevice(¤tDevice)N/A选择装备cudaSetDevice(deviceId)N/A查询装备内存和使用情况N/Anvidia-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 系统中。
免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。 |
