(四）、CUDA环境安装与程序编写

Wmpreturn

已于 2023-08-24 10:36:50 修改

阅读量498

点赞数

分类专栏： Linux系统和环境搭建简介文章标签： linux 运维服务器

于 2023-08-16 14:24:01 首次发布

本文链接：https://blog.csdn.net/weixin_48060069/article/details/132311796

版权

Linux系统和环境搭建简介专栏收录该内容

6 篇文章 2 订阅

订阅专栏

1、确定安装版本

        1、查看显卡支持的最高CUDA的版本，以便下载对应的CUDA安装包：

安装 NVIDIA 显卡驱动程序、查看当前系统中 NVIDIA 显卡的详细信息：nvidia-smi、

在 "CUDA 版本" 或 "Compute Capability" 部分查找显示的显卡支持的 CUDA 版本号。
        2、确定CUDA版本对应的cuDNN版本CUDA →上网搜索（建议参考 NVIDIA 提供的官方文档或发布说明）。

        3、Toolkit(nvidia):CUDA完整的工具安装包，其中提供了Nvidia驱动程序、开发CUDA程序相关的开发工具包等可供安装的选项。包括CUDA程序的编译器、IDE、调试器等，CUDA程序所对应的各式库文件以及它们的头文件。

2、安装方法

官网下载安装，并下载cuDNN。

3、CUDA环境变量配置

如果上面的nvcc-v命令提示错误，需要将以下命令放入环境变量配置文件~/.bashrc中（假
设环境中的CUDA版本为12.0)
则添加下面两行即可：
export PATH=$PATH:/usr/local/cuda-12.0/bin
export LD LIBRARY PATH=$LD LIBRARY PATH:/usr/local/cuda-12.0/1ib64

NVCC(NVIDIA CUDA Compiler)是由NVIDIA提供的用于编译CUDA代码的编译器。它将CUDA C/C++代码转换为可以在GPU上执行的二进制代码，从而实现GPU加速计算。NVCC不仅是编译器，还提供了一些用于管理编译和构建过程的选项。
NVCC支持将CUDA代码与普通的C/C++代码混合编译，允许开发者在同一文件中同时编写主机(CPU)代码和设备(GPU)代码。开发者可以使用CUDA扩展的C/C++语法来编写设备代码，包括CUDA核函数、线程和块的控制等。

4、NVCC的编译选项

设备选项用的不是很多。

5、使用NVCC编译简单的CUDA程序

1、编写CUDA源码
编写包含CUDA代码的源代码文件。通常，CUDA代码会包含主机代码（在CPU上运行)和设备代码（在GPU上运行）。
2、使用nvcc编译
3、运行可执行文件：使用生成的可执行文件来运行CUDA程序.
4、对程序进行性能分析

6、编码试验-checkdeviceinfo
#include<iostream>
#include<cuda.h>
#include<cuda_runtime.h>
int main() {
int dev = 0;
cudaDeviceProp devProp;
cudaGetDeviceProperties(&devProp, dev);
std::cout << "GPU Device Name" << dev << ": " << 
devProp.name << std::endl;
std::cout << "SM Count: " << 
devProp.multiProcessorCount << std::endl;
std::cout << "Shared Memory Size per Thread Block: " << 
devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
std::cout << "Threads per Thread Block: " << 
devProp.maxThreadsPerBlock << std::endl;
std::cout << "Threads per SM: " << 
devProp.maxThreadsPerMultiProcessor << std::endl;
std::cout << "Warps per SM: " << 
devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
return 0;
}
// Kernel定义
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) 
{ 
int i = blockIdx.x * blockDim.x + threadIdx.x; 
int j = blockIdx.y * blockDim.y + threadIdx.y; 
if (i < N && j < N) 
C[i][j] = A[i][j] + B[i][j]; 
}
int main() 
{ 
...
// Kernel 线程配置
 dim3 threadsPerBlock(16, 16); 
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
// kernel调用
 MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); 
...
}
编码试验-与串行代码的区别（设备部分）

在CUDA编程中，threadIdx、blockIdx、blockDim 和 gridDim 是内置变量，用于在并行计算中确定线程和线程块的索引和维度。用于操作GPU上的线程块和线程。

threadIdx.x、threadIdx.y 和 threadIdx.z 分别表示当前线程在线程块内的 x、y 和 z 方向的索引。

每个线程块内的线程都有自己的 threadIdx。

blockIdx.x、blockIdx.y 和 blockIdx.z 表示当前线程块在网格内的 x、y 和 z 方向的索引。每个线程块都有自己的 blockIdx。

blockDim.x、blockDim.y 和 blockDim.z 表示一个线程块内的线程数量的 x、y 和 z 方向的维度。

这是一个固定值，对于大多数情况下，都是相同的。

gridDim.x、gridDim.y 和 gridDim.z 表示整个网格的 x、y 和 z 方向的维度，即网格中线程块的数量。这是一个固定值，对于整个网格都是相同的

7、CUDA程序性能检测工具-nvprof

nvprof是NVIDIA提供的一个命令行工具，可以被用于收集关于CUDA应用程序的性能指标数据，例如GPU利用率、内存带宽、运行时间和延迟等信息。此外，nvprofi还可以在GPU代码执行期间跟踪CUDA APIi调用，以识别可能导致性能瓶颈的函数和代码路径。
nvprofj提供了多种分析选项，例如时间轴视图、函数摘要视图和指令分析视图等，以辅助更好地理解CUDA应用程序的性能瓶颈。

简单的时间轴视图
nvprof./my_cuda_app

指定统计信息：gpu使用率、占用情况

nvprof./my_cuda_appnvprof --metrics

gpu_utilization,achieved_occupancy ./my_cuda_app

指定输出文件
nvprof --output-profile
my_profile.nvvp ./my_cuda_app

使用分析视图
nvprof --analysis-metrics -o
my_profile.nvvp./my_cuda_app

8、块和线程索引的映射（矩阵相关）

通常情况下，矩阵是用行优先的方法在主机内存中线性存储的，CUDA程序中，可以建立二维网格(2,3)+二维块(4,2) ，使用其块和线程索引映射矩阵索引。

对于加法而言，结果矩阵的（i，j）即由两个n * m的矩阵对应坐标（i，j）相加而来，一共进行n * m次加法操作，一共需要n * m个线程，即至少要开辟的线程数 N = n * m。

对于乘法而言，结果矩阵的（i，j）是由第一个矩阵（n * k）的第 i 行和第二个矩阵（k * m）的第 j 列各个元素相加而来——每一个线程需要完成的任务。线程总数N = n * m。

当得到N (n * m)后，可以设置blockDim（可能有最大值1024的限制，通常取32的整数倍），系统自动生成gridDim。

在平时写的c语言中（i，j）表示第 i 行第 j 列。

在矩阵坐标中（ix，iy）表示第 iy 行，第 ix 列，即和 x、y轴一样的坐标系（x水平、y垂直）

`blockIdx.x` 是 CUDA 编程中的一个内置变量，用于表示当前线程块在 x 方向上的索引。

在 CUDA 中，线程以线程块的形式组织在网格（grid）中。每个线程块由若干个线程组成，线程块则被组织成一个三维网格。`blockIdx.x` 表示当前线程块在 x 方向上的索引，即在整个网格中的位置。

通过 `blockIdx.x` 可以在 CUDA 程序中进行条件判断或计算，来使不同线程块执行不同的操作或访问不同的数据。

int ix = blockIdx.x * blockDim.x + threadIdx.x;

int iy = blockIdx.y * blockDim.y + threadIdx.y;

这段代码计算了在二维网格中当前线程的全局索引。其中：

- `blockIdx.x` 和 `blockIdx.y` 分别表示当前线程块在 x 和 y 方向上的索引。
- `blockDim.x` 和 `blockDim.y` 分别表示每个线程块中线程的数量（在 x 和 y 方向上）。
- `threadIdx.x` 和 `threadIdx.y` 分别表示当前线程在所属线程块内的索引（在 x 和 y 方向上）。

通过这些值，可以计算出当前线程在整个网格中的全局索引。

具体来说，`ix` 的计算方式为：当前线程的 x 坐标是 `blockIdx.x * blockDim.x + threadIdx.x`。

同样地，`iy` 的计算方式为：当前线程的 y 坐标是 `blockIdx.y * blockDim.y + threadIdx.y`。

这样计算得到的 `ix` 和 `iy` 可以用于在 CUDA 程序中访问和处理不同线程的数据。

9、线程块内的线程数

线程块中的线程数最好配置为 32 的倍数，是因为 GPU 的硬件特性和并行计算模型的设计所决定的。

GPU 以线程块为单位执行并行计算任务，每个线程块内的线程可以协同工作，共享数据和通信。GPU 在执行计算任务时，会将线程块划分为更小的线程组（warps），每个线程组包含一组连续的线程。

GPU 的硬件是以线程组（warp）为单位调度和执行的，具体来说，每个时钟周期，GPU 会选择一个线程组（warp）来执行，而不是单独调度和执行每个线程。在同一个线程组（warp）内的线程需要以相同的指令流进行执行，称为 SIMT（Single Instruction, Multiple Threads）执行模型。

这就决定了线程块中的线程数最好配置为 32 的倍数，因为一个线程组（warp）内有固定数量的线程，通常是 32 个（具体数量可能因不同 GPU 架构而异）。如果线程块中的线程数不是 32 的倍数，就会导致线程组（warp）之间的负载不平衡。比如，如果线程块中的线程数是 40，那么就会有一个线程组（warp）包含 40 个线程，而另一个线程组（warp）只有 8 个线程，这样就会浪费 GPU 的计算资源。

因此，为了充分发挥 GPU 的并行计算能力，以及保持线程组的负载平衡，最好将线程块中的线程数配置为 32 的倍数。这样可以确保所有线程组（warp）都在同一时钟周期内被调度和执行，从而提高计算性能。

但是在实际工作中，很可能会出现手动配置参数创建的线程数不等于并行循环所需的线程数，例如：实际上需要执行1230次循环，但通常会配置2048个线程。

1. 设置配置参数，使线程总数超过实际工作所需的线程数。

2. 在向核函数传递参数时传递一个用于表示要处理的数据集总大小或完成工作所需的总线程数 N。

3. 计算网格内的线程索引后（使用 threadIdx + blockIdx*blockDim），判断该索引是否超过 N，只在不超过的情况下执行与核函数相关的工作。

注：适用于工作总量 N 和线程块中的线程数已知的情况。
// 假设N是已知的
int N = 100000;
// 把每个block中的thread数设为256
size_t threads_per_block = 256;
// 根据N和thread数量配置Block数量
size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;
//传入参数N
some_kernel<<<number_of_blocks, threads_per_block>>>(N);
number_of_block = (N + threads_per_block - 1) / threads_per_block;

这句话是对N / threads_per_block进行上取整。

向上取整的好处：使线程数量一定满足大于等于N，即一定能满足执行N次任务的需求且浪费也最多浪费一个block的线程数量。

10、核函数

定义核函数：

核函数（Kernel Function）是指在并行计算中在 GPU 上执行的函数。在 CUDA 编程模型中，核函数是由开发者编写的在 GPU 上执行的并行计算任务代码。

核函数在 CUDA 中通过 `__global__` 修饰符来标识，并且可以接受参数和返回值。在核函数中，可以使用特定的语法和功能来指定并行计算的方式，例如使用线程索引、线程块和网格等来管理并行执行。

在 CUDA 编程中，核函数会被调用多次以在不同的线程上执行，并且每个线程独立地执行相同的计算任务。在执行核函数时，线程索引、线程块索引和网格索引等信息可以用来确定每个线程在计算中的角色和任务。

核函数通常用于执行密集的数值计算任务，如向量加法、矩阵乘法等，并且能够充分利用 GPU 的并行计算能力，提高计算性能。通过编写适当的核函数，可以将计算任务划分为多个并行的线程块，在 GPU 上同时执行，从而加速计算过程。

需要注意的是，核函数中不能直接调用 CPU 上的函数或访问 CPU 内存中的数据，因为 GPU 和 CPU 是两个独立的计算设备。如果需要在核函数中使用 CPU 上的数据，需要通过将数据从主机（CPU）内存复制到设备（GPU）内存，并在核函数中访问设备内存来实现。同样地，如果需要将计算结果从设备内存复制回主机内存，也需要进行相应的数据传输操作。

总之，核函数是在 GPU 上执行的并行计算任务代码，通过充分利用 GPU 的并行计算能力，可以加速密集的数值计算任务。

计算分配：

将任务分配给线程块和线程。确定块的数量和每个块中线程的数量，这取决于问题的规模和GPU的性能。

数据访问和同步：

确保多个线程不会相互干扰。在核函数内，使用共享内存来减少全局内存的访问次数，以提高性能。需要注意线程之间的同步问题，特别是在使用共享内存时，使用 __syncthreads()来同步线程。

11、归约（reduction）算法

归约算法（Reduction Algorithm）是一种常见的并行计算算法，用于将一个规模较大的问题转化为规模较小的子问题，并通过合并子问题的结果来得到原问题的解。

在并行计算中，归约算法通常用于对一个数据集中的元素进行聚合操作，例如求和、求最大值、求最小值等。归约算法通过将数据集划分为多个部分，分配给不同的处理单元（例如线程、线程块或处理器核心）进行并行计算，并将每个部分的结果合并以得到最终的结果。

下面是一种常见的归约算法的示例，以求和操作为例：

1. 将输入数据集平均划分为多个小部分，并分配给不同的处理单元。
2. 每个处理单元分别计算其所分配的部分的局部和。
3. 各处理单元将局部和相加，得到一个全局和。
4. 重复上述步骤，直到所有部分都被合并为一个结果。

这种归约算法可以通过反复地将问题的规模减半来实现，从而达到较高的并行性能。在每个迭代中，将问题划分为更小的部分，并在每个部分上执行归约操作。这样，每次迭代的规模都会减半，而并行计算的规模则会增加。

归约算法在很多应用中都有广泛的应用，特别是在并行计算和并行编程模型（如CUDA、OpenMP等）中。它可以提高计算效率，并充分利用多个处理单元的并行计算能力，从而加速问题的求解过程。

需要注意的是，归约算法的性能与数据的划分和合并方式、并行计算的负载均衡以及通信开销等因素密切相关。在设计和实现归约算法时，需要考虑这些因素以保证算法的可扩展性和效率。

优化归约算法

·避免线程束分化：当一个线程束的线程执行判断语句的不同分支时，满足分支条件的线程会执行该分支的命令，不满足分支条件的线程就闲置且不能跳过。这样，整个线程束的执

行效率就会比没有分支的情况低一半。

·连续地址读取：CUDA中线程对数据的连续读取效率要比其它方式高，故我们可以将线程取址方式变为连续的。

·循环展开：操作的瓶颈可能在寻址和循环本身。将多个规约步骤合并到一个循环中，从而减少控制流的开销。

多步归约：使用多个归约阶段进行连续的归约操作。这些阶段可以使用不同的方法和参数，逐步减小数据规模，从而更有效地利用并行性。

不满足分支条件的线程就闲置且不能跳过的原因：

在GPU并行计算中，线程束（Thread Warp）是一组连续的线程，通常包含32个线程。这些线程将同时执行相同的指令，但可能具有不同的数据。当线程束中的线程执行判断语句的不同分支时，满足分支条件的线程会执行对应分支的指令，而不满足条件的线程则没有执行的需要。

由于GPU架构中的线程束是以SIMD（Single Instruction Multiple Data）方式工作的，即一条指令同时作用于线程束中的所有线程，因此在执行分支语句时，无法跳过不满足条件的线程。这是由GPU硬件设计所决定的，每个线程都必须按照指令流水线的节奏进行执行，无法单独选择执行或跳过某些指令。

当线程束中的线程执行不同分支时，不满足条件的线程会被闲置。这意味着它们并不执行任何实质性的计算任务，只是等待其他线程完成对应分支的指令执行。这种情况发生时，线程束的吞吐量会下降，因为部分线程没有有效的工作可以执行。

为了避免线程束分化和提高GPU的利用率，可以采取一些优化策略。例如，合理编写代码，尽量避免分支语句的出现；或者通过数据重排、数据预取等技术来增加线程束中线程的工作负载，减少闲置时间。此外，在GPU程序设计中，可以使用更大的线程束大小（如64或128个线程），以减少线程束分化的影响。