0

我有一个 CUDA 程序，其中内核寄存器将最大理论达到的占用率限制为 %50。因此，我决定对那些在块线程之间保持不变并且在整个内核运行期间几乎是只读的变量使用共享内存而不是寄存器。我不能在这里提供源代码；我所做的在概念上是这样的：

我的初始程序：

__global__ void GPU_Kernel (...) {
    __shared__ int sharedData[N]; //N:maximum amount that doesn't limit maximum occupancy
    int r_1 = A; //except for this first initialization, these registers don't change anymore
    int r_2 = B;
    ...
    int r_m = Y;

    ...   //rest of kernel;
}

我将上面的程序更改为：

__global__ void GPU_Kernel (...) {
    __shared__ int sharedData[N-m];
    __shared__ int r_1, r_2, ..., r_m;
    if ( threadIdx.x == 0 ) {
        r_1 = A;
        r_2 = B;
        ...
        r_m = Y; //last of them
    }
    __syncthreads();
    ...   //rest of kernel
}

现在块内的扭曲线程执行广播读取以访问新创建的共享内存变量。同时，线程不会使用过多的寄存器来限制实现的占用。

第二个程序的最大理论占用率等于 %100。在实际运行中，第一个程序的平均占用率为 ~%48，第二个程序的平均占用率为 ~%80。但问题是净加速的提高大约是 %5 到 %10，这比我预期的提高入住率的预期要低得多。为什么这种相关性不是线性的？

考虑下面来自 Nvidia 白皮书的图片，我一直在想，当实现占用率为 %50 时，例如，一半的 SMX（在较新的架构中）内核一次处于空闲状态，因为其他内核请求的资源过多会阻止它们积极的。我的理解有问题吗？还是对上述现象的解释不完整？还是添加__syncthreads();和共享内存访问成本？

在此处输入图像描述

4

1 回答 1

2

为什么这种相关性不是线性的？

如果您已经受到内存带宽限制或计算限制，并且这些限制中的任何一个都接近设备的理论性能，那么提高占用率可能无济于事。当这些都不是您的代码性能的限制因素时（即您未处于或接近峰值内存带宽利用率或峰值计算），提高占用率通常会有所帮助。由于您没有为您的程序提供任何代码或任何指标，因此没有人可以告诉您为什么它没有加快速度。分析工具可以帮助您找到性能限制因素。

您可能对几个网络研讨会感兴趣：

CUDA 优化：识别性能限制器 Paulius Micikevicius 博士

CUDA Warps and Occupancy Considerations+ 与 NVIDIA 的 Justin Luitjens 博士一起直播

特别是，查看第二次网络研讨会的幻灯片 10 。

于 2013-12-02T22:39:38.423 回答