Cuda 编程中的主动扭曲

Question

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

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

我这样称呼我的内核：

VecAddBLOCK_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 个 8 个 warp 块。实现的入住率统计显示，您平均有 1 个经线处于活动状态，非常低。显而易见的问题是为什么这不是 8。

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

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