CUDA——矢量求和

最新推荐文章于 2024-07-12 14:57:56 发布

UCAS_HMM

最新推荐文章于 2024-07-12 14:57:56 发布

阅读量770

点赞数

分类专栏： CUDA 文章标签： c++ 数据结构

本文链接：https://blog.csdn.net/ucas_hmm/article/details/126523901

版权

CUDA 专栏收录该内容

10 篇文章 3 订阅

订阅专栏

CUDA-矢量求和

介绍
- 问题描述
- 代码实现和运行
性能分析
总结

介绍

问题描述

对包含nElem元素的两个矢量A和B进行求和，结果放在C中， ${\bf A+B=C}$ 。

代码实现和运行

基本实现逻辑：在主机上创建主机内存和设备内存，并初始化主机矢量h_arrA和h_arrB，然后将数据拷贝到设备矢量h_arrA和h_arrB中，调用核函数sumArr进行计算，得到结果d_arrC，将结果拷贝回主机矢量h_arrC上，调用check函数检查运算结果是否正确。

//- 只展示了关键代码段，并不是全部完整的代码
__global__ void sumArr(int *arrA, int *arrB, int *arrC, int nElem)
{
    int Index = blockIdx.x * blockDim.x + threadIdx.x;
    if (Index < nElem)
    {
        arrC[Index] = arrA[Index] + arrB[Index];
    }
}

int main(int argc, char **argv)
{
    int nElem = 1<<24;
    dim3 block(atoi(argv[1]));  //- block的值通过命令行输入
    dim3 grid((nElem+block.x-1)/block.x);
    sumArr<<<grid, block>>>(d_arrA, d_arrB, d_arrC, nElem);
    cudaDeviceSynchronize();
    return 0;
}

编译和运行结果

nvcc -arch=sm_70 -O0 -c main.cu -o test

[xxxx@k052 chapter2]$ ./test 512
gridDim:(32768,1,1) blockDim:(512,1,1)
result is correct!

性能分析

通过nsight compute对该核函数进行分析。

Launch Statistics

核函数发布配置

Threads：该矢量的大小为 $2^{24}=16777216$ ，为32的整数倍，因此总共会发布这么多线程；
Block Size：用户命令行输入数据，为512；
Grid Size：thread和block都是一维配置，因此可以计算出共有32768个block数；
Registers Per Thread：每个线程所需的寄存器数量为16，具体如何从代码中确定寄存器数量还不会；
Static Shared Memory Per Block：因为共享内存是一个block内创建并共享的，共享内存是以block为基本统计单位的。这是静态内存，通过__shared__表示符申明。
Dynamic Shared Memory Per Block：动态共享内存，通过extern __shared__标识符申请，具体大小在发布核函数时传递进来；
Driver Shared Memory Per Block：
Shared Memory Configuration Size：
其中，共享内存和L1/TEX缓存共享片上内存，可以通过cudaFuncSetCacheConfig来进行用户自定义分配。

在这个信息页中提供的有用信息比较少，主要是看寄存器和共享内存的消耗情况，因为这两个指标对占有率的影响非常大。

Occupancy

占有率的定义是实际常驻warp数和SM支持的理论最大warp数的比值。高的占有率不一定意味着高效，但是一般来说最好位置占有率处于较高的水平。

SM支持的理论最大warp数是通过算例标准规定的，可以查询官方文档。不同的设备有不同的计算能力，同样也就对应不同的硬件上限，本计算平台是V100，算力7.0，因此后面的分析都是基于这个硬件来的。

Properties	Value
每个线程的寄存器上限	255
每个块的寄存器上限	64K
每个SM的寄存器上限	64K
每个块的共享内存上限	96KB
每个SM的共享内存上限	96KB
每个SM的驻留块上限	32
每个SM的驻留线程上限	2048

而实际的SM常驻warp数量是通过木桶原理来计算出来的，影响这个数据的三个指标分别为：

Block Limit Registers：该核函数每个线程消耗16个寄存器，一个block有512个线程，因此一个block需要8192个寄存器，而一个SM中最多可使用的寄存器为64K个，因此最多可以容纳8个这样的block；
Block Limit Shared Mem：由于该核函数不需要共享内存，因此这个指标不构成限制，直接取上限32；
Block Limit SM：这个指标直接给出的是32；
Block Limit Warps：因为SM最大允许的驻留线程数为2048，即64个warps，按照当前512个线程一个block计算，block将会被限制在4个。

因此，取这些参数的最小值4即为每个SM常驻block数量，对应为64个warp，那么理论占有率应该为100%。然而实际达到的占有率只有87.96%，这个暂时不清楚。

从这个线程配置计算可以推论出，如果block包含的线程数量很多，即大block，可能会导致每个block消耗的寄存器和共享内存很多，使得常驻block数变少。而小block则会容易触碰SM对Block数的限制，导致最终常驻的线程数量较少。

这部分还给出了线程消耗寄存器、共享内存以及block线程数单变量分析：

单变量分析

Memory Workload Analysis

Memory Chart

Memory Chart
从这张图里面我们可以看到数据的流动路径，逻辑上，核函数对全局内存进行读写，物理上，核函数依次通过L1、L2、设备内存进行读写。在本例子中，A、B和C的数据均为 $2^{24}\times 4byte = 64MB$ ，因此整个过程需要加载的数据量为128MB，需要写入的数据量为64MB。

L1/TEX Cache

L1/TEX缓存

核函数一共执行了1572864次针对Global的读写指令，其中包括1048576次Global Load指令和524288次Global Store指令。通过简单计算可以知道，每次Global Load指令加载了128个字节，每次Global Store指令同样也是写入了128个字节。Sectors/Req为4，说明一次可以加载4个sector，此时为最优情况，对齐合并访问。

Device Memory

Device Memory
这里重点观察%peak这个指标，与Speed of Light中的内存利用率是一致的。

Scheduler Statistics

Scheduler Statistics
V100的SM中包含4个warp调度器，每个调度器可以管理16个warp，因此每个SM能够常驻64个warps。从图中可以看到，每个调度器中活动的warp数还比较接近16，这说明SM的占有率是比较高的。但是平均每个周期内，就绪的warp（Eligible warp）数却非常少，这说明warp池中的warp大部分是出于阻塞状态的（stalled warp），同时每个周期由调度器发射的warp数（issued warp或Selected warp）也比较少。这就属于典型的占有率很高，但是计算核心长期处于闲置状态，利用率很低的情况。

Warp State Statistics

在这里插入图片描述
从这张图里面就能了解为什么warp长时间处于stall状态。因为大量的时间在等待Long Scoreboard和LG Throttle这两种指令。说明该warp是因为对全局内存的访问而阻塞。

Long Scoreboard：等待局部、全局、表面、纹理内存操作；
LG Throttle：等待本地和全局（LG）内存操作的L1指令队列未满。通常，仅当非常频繁地执行本地或全局内存指令时，才会发生此暂停。如果适用，请考虑将多个较宽的内存操作组合为较少的较宽内存操作，并尝试交错内存操作和数学指令。
Short Scoreboard：因为共享内存的频繁操作、特殊功能函数等。
Drain：
Not Selected：
Dispatch Stall：
IMC Miss：

Compute Workload Analysis

Compute Workload Analysis
这里有提到了SM的忙碌程度为8.67%，指令发射槽的忙碌程度也是8.67，这和前面的现象是对应的，因为大部分warp是阻塞的，因此尽管warp池占满的，但是就绪的warp很少，大家都在等待数据传输，所以计算单元大部分时间是闲置的。这里列出了一些基本操作：

LSU:
ADU:
FMA:
LSU:
CBU:
FP16:
FP64:
TEX:
Tensor(FP):
XU:

总结

该核函数的特点是计算简单，数据读写负载大。对全局内存数据的访问延时大约为150个指令周期，而算术指令大约为3个周期，因此需要每个调度器拥有50个常驻warp才能掩盖延迟。但是本设备每个调度器最多只包含16个warp，因此必然导致大部分时间所有的常驻内存是阻塞的，导致计算单元闲置。这种核函数没有更好的优化方式，在实际应用中，这一类核函数最好是和其他核函数合并，使数据加载之后能够被充分利用起来，增加算术指令的耗时占比，这样才能降低所需要的warp数量。

在这里插入图片描述
例如上述蓝色表示的是增加了对加载数据的计算复杂度，和紫色的naive比较，平均下来的每个warp的阻塞周期大幅缩短，同时就绪的warp也增多。