c++ - Cuda 编程中的主动扭曲

Question

我正在尝试使用 Nsight IDE 对我的代码进行性能分析。

我举了一个简单的矩阵加法的例子。

我这样称呼我的内核:

VecAdd<<<1,BLOCK_SIZEBLOCK_SIZE>>>(dA,dB,dC,BLOCK_SIZEBLOCK_SIZE);

这里 BLOCK_SIZE 是 16。

__global__ void VecAdd(float *dA, float *dB, float *dC, int N)
{
    int i = threadIdx.x;
    if (i < N)
        dC[i] = dA[i] + dB[i];
}

在进行占用分析时，

我得到的 Active warp 达到了 0.97。

我不知道为什么。

我附上了一份报告。有人可以解释为什么会这样吗？

Answer 1

Achieved Occupancy 是 active_warps/elapse_cycles/MAX_WARPS_PER_SM * 100 的百分比。

您的内核启动是 1 个 block ，每 block 8 个 warp。实现的占用统计显示您平均有 1 个 warp 事件，这是非常低的。显而易见的问题是为什么这不是 8。

由于您没有提供源代码，我假设您修改了 VecAdd CUDA SDK 示例，该示例执行 5 次常量读取、2 次 32 位全局加载、1 次 32 位写入以及一些用于索引和地址计算的基本数学运算。假设所有内存操作都在 L2 中命中，则每个 warp 大约需要 300 个周期。这很可能是因为您可能在启动之前将阵列从主机复制到设备。内核持续时间本身可能是 2-3 微秒。 8 * 300 个周期/2500 个周期 = 1 个 SM 上每个周期约 1 个事件扭曲。

启动开销、工作分配开销以及等待每个 warp 存储清除写入数据缓冲区的时间不算作 8 个 warp 处于事件状态的时间。如果增加每个 warp 的工作量，该值将增加到接近 8，这是给定启动线程数后可实现的最大值。如果您还增加网格大小以使设备饱和，您应该能够使每个 SM 的平均活跃 warp 接近 64 个。