CUDA带宽测试疑问：计算操作为何未被内存延迟完全掩盖？

安全阿岳

2026-9-24

CUDA显存带宽测试循环展开性能异常问题

我编写了一个测试GPU显存带宽的程序，测试过程中发现了几个异常现象。

其中一个现象是，直观实现的循环并非完全受内存带宽限制。原本的循环代码如下：

for (int *p = pStart; p < pEnd; p += shift)
    sum += *p;

我对上述循环手动展开，在循环体中添加了以下代码：

p += shift; sum += *p;
p += shift; sum += *p;
p += shift; sum += *p;

在GeForce RTX 2060上测试时，速度提升了20%，从183GB/s升至222GB/s（该显卡理论带宽为264GB/s）。

我对该现象存在以下疑问：

这类开销难道不应该被内存延迟隐藏吗？理想情况下这类程序的所有warp应该都在等待内存返回数据，warp在等待间隙执行的额外计算难道不会导致内存总线负载不足吗？

我使用NVidia Nsight Compute 2021.3.0分析了可执行文件，报告显示展开后的版本计算吞吐量更低、内存吞吐量更高，缓存占用可以忽略，这些结果都符合逻辑。但调度器的统计结果很值得关注，我无法理解这些数值的含义。我并非CUDA新手，知道warp可以在流水线中执行，也会因等待内存数据而停顿，但我不太清楚合格warp（eligible warps）、每条发射指令对应的warp周期是什么意思（展开版本中每个warp平均每114个GPU时钟周期才执行一条指令吗？）

Nsight Compute调度器统计结果
Base对应未展开的版本，main values对应展开后的版本。例如展开后每个调度器的活跃warp数为7.88，未展开版本为7.96——难道不是活跃warp数越高越好吗？

完整测试代码

重点查看__global__ void gpuReadMemory函数：

#include <assert.h>
#include <conio.h>
#include <stdio.h>
#include <cuda.h>
#define VC_EXTRALEAN
#include <windows.h>

#define CUDA_CHECK(err)     __cudaSafeCall(err, __FILE__, __LINE__)

inline void __cudaSafeCall(cudaError err, const char *file, const int line)
{
    if (err != cudaSuccess)
    {
        fprintf(stderr, "%s(%i): CUDA error %d (%s)\n",
                    file, line, int(err), cudaGetErrorString(err));
        throw "CUDA error";
    }
}

int getMultiprocessorCount()
{
    int num;

    CUDA_CHECK(cudaDeviceGetAttribute(&num, cudaDevAttrMultiProcessorCount, 0));
    return num;
}

__global__ void gpuWriteMemory(int *gpuArr, int dataSizeInMBs, 
                                                int packetShift, int passCount, int *gpuDebugArr)
{
    int *pStart = gpuArr + ((long long)packetShift * 
                (blockDim.y * blockIdx.x + threadIdx.y)) / sizeof(int) + threadIdx.x;
    int *pEnd = gpuArr + (((long long)dataSizeInMBs) << 20) / sizeof(int);
    int shift = gridDim.x * blockDim.y * packetShift / sizeof(int);

    for (int passInd = 0; passInd < passCount; passInd++)
        for (int *p = pStart; p < pEnd; p += shift)
            *p = blockIdx.x * 10000000 + threadIdx.y * 1000 + threadIdx.x;
}

__global__ void gpuReadMemory(int *gpuArr, int dataSizeInMBs, 
                                                int packetShift, int passCount, int *gpuDebugArr)
{
    int *pStart = gpuArr + ((long long)packetShift * 
                (blockDim.y * blockIdx.x + threadIdx.y)) / sizeof(int) + threadIdx.x;
    int *pEnd = gpuArr + (((long long)dataSizeInMBs) << 20) / sizeof(int);
    int shift = gridDim.x * blockDim.y * packetShift / sizeof(int);
    int sum = 0;
    int accessCount = 0;

    for (int passInd = 0; passInd < passCount; passInd++)
    {
#pragma unroll   // - doesn't have effect
        for (int *p = pStart; p < pEnd; p += shift)
        {
            sum += *p;

            p += shift; sum += *p;
            p += shift; sum += *p;
            p += shift; sum += *p;
        }
      *pStart = sum;   // Without it bandwidth reported is 3 times bigger than theoretical
    }

    // Suspiciously fast code:
    //for (int passInd = 0; passInd < passCount; passInd++)
    //  0x0000000500999d10               IADD3 R8, R8, 0x1, RZ  
    //  0x0000000500999d20               BSSY B0, 0x500999de0  
    //  for (int *p = pStart; p < pEnd; p += shift)
    //      0x0000000500999d30               ISETP.GE.AND.EX P0, PT, R7, UR6, PT, P0  
    //      for (int passInd = 0; passInd < passCount; passInd++)
    //          0x0000000500999d40               ISETP.GE.AND P1, PT, R8, c[0x0][0x170], PT  
    //          for (int *p = pStart; p < pEnd; p += shift)
    //              0x0000000500999d50          @P0  BRA 0x500999dd0  
    //              0x0000000500999d60               IMAD.MOV.U32 R3, RZ, RZ, R2  
    //              0x0000000500999d70               IMAD.MOV.U32 R5, RZ, RZ, R4  
    //              0x0000000500999d80               LEA R3, P0, R0, R3, 0x2  
    //              0x0000000500999d90               LEA.HI.X R5, R0, R5, RZ, 0x2, P0  
    //              0x0000000500999da0               ISETP.GE.U32.AND P0, PT, R3, UR4, PT  
    //              0x0000000500999db0               ISETP.GE.U32.AND.EX P0, PT, R5, UR5, PT, P0  
    //              0x0000000500999dc0          @!P0 BRA 0x500999d80  
    //              0x0000000500999dd0               BSYNC B0  

    // Normal code:
    //for (int *p = pStart; p < pEnd; p += shift)
    //  0x0000000500999da0               IMAD.MOV.U32 R8, RZ, RZ, R6  
    //  sum += *p;
    //0x0000000500999db0               IMAD.MOV.U32 R4, RZ, RZ, R8  
    //  0x0000000500999dc0               LDG.E.SYS R4, [R4]                // Reading from memory
    //  for (int *p = pStart; p < pEnd; p += shift)
    //      0x0000000500999dd0               IADD3 R11, P0, R11, UR5, RZ  
    //      0x0000000500999de0               IADD3 R8, P1, R8, UR5, RZ  
    //      0x0000000500999df0               IADD3.X R13, R13, UR4, RZ, P0, !PT  
    //      0x0000000500999e00               ISETP.GE.U32.AND P0, PT, R11, R2, PT  
    //      0x0000000500999e10               IADD3.X R5, R5, UR4, RZ, P1, !PT  
    //      0x0000000500999e20               ISETP.GE.U32.AND.EX P0, PT, R13, R3, PT, P0  
    //      sum += *p;
    //0x0000000500999e30               IMAD.IADD R9, R4, 0x1, R9  
    //  for (int *p = pStart; p < pEnd; p += shift)
    //      0x0000000500999e40          @!P0 BRA 0x500999db0  
    //      0x0000000500999e50               BSYNC B0  
    //      *pStart = sum;   // Lowers bandwidth being reported on GT 520M from 33 to 9.7 GB/s
    //0x0000000500999e60               STG.E.SYS [R6], R9  
}

class CMemorySpeedTester
{
public:
    CMemorySpeedTester()
    {
        m_multiprocessorCount = getMultiprocessorCount();

        CUDA_CHECK(cudaMalloc((void**)&gpuArr, ((long long)dataSizeInMBs) << 20));
        CUDA_CHECK(cudaMalloc((void**)&gpuDebugArr, debugArrLen * sizeof(int)));
        CUDA_CHECK(cudaMemset(gpuArr, -1, ((long long)dataSizeInMBs) << 20));

        debugArr = (int*)malloc(debugArrLen * sizeof(int));
        debugArr2D = (int (*)[1000])debugArr;       
        QueryPerformanceFrequency(&timerFreq);
    }

    void testBandwidth()
    {
        int threadPerBlock = 256;
        int blockCount = m_multiprocessorCount * 24;
        dim3 blocks(blockCount);
        dim3 threads1(32, threadPerBlock / 32);

        gpuWriteMemory<<<blocks, threads1>>>(gpuArr, dataSizeInMBs, 
                    32 * sizeof(int), 1, gpuDebugArr);
        CUDA_CHECK(cudaDeviceSynchronize());

        for (int passCount = 10; passCount <= 100; passCount *= 10)
        {
            int threadPerPacket = 32; 
            for (int packetShiftMult = 1; packetShiftMult <= 16; packetShiftMult *= 16)
            {
                int packetShift = threadPerPacket * sizeof(int) * packetShiftMult;
                dim3 threads(threadPerPacket, threadPerBlock / threadPerPacket);
        
                QueryPerformanceCounter(&t0);
                gpuReadMemory<<<blocks, threads>>>(gpuArr, dataSizeInMBs, 
                            packetShift, passCount, gpuDebugArr);
                CUDA_CHECK(cudaDeviceSynchronize());
                QueryPerformanceCounter(&t);

                double dt = double(t.QuadPart - t0.QuadPart) / timerFreq.QuadPart;

                printf("  %2d th./packet, packet shift %4d, %d pass(es): %.3f ms, %.2f GB/s\n",
                                threadPerPacket, packetShift, passCount,
                                dt * 1000, (double)(dataSizeInMBs) / packetShiftMult / (1 << 10) * passCount / dt);
            }
        }
    }

    void copyToHostDebugArr()
    {
        CUDA_CHECK(cudaMemcpy(debugArr, gpuDebugArr, debugArrLen * sizeof(int), cudaMemcpyDeviceToHost));
        CUDA_CHECK(cudaDeviceSynchronize());
    }

protected:
    static const int dataSizeInMBs = 800;
    static const int debugArrLen = 4000000;

    int m_multiprocessorCount;
    int *gpuArr, *gpuDebugArr;
    int *debugArr;
    int (*debugArr2D)[1000];
    LARGE_INTEGER timerFreq, t, t0;
};

int main(int argc, char **argv)
{
    printf("Started\n");

    CMemorySpeedTester runner;

    for (int runInd = 0; runInd < 5; runInd++)
    {
        printf("%d.\n", runInd);
        runner.testBandwidth();
    }

    printf("Finished. Press Enter...");
    getch();
}

测试环境

主测试环境：Windows 10 x64、Visual Studio 2019、CUDA Toolkit 10.2
兼容复现环境：GeForce GT 520M、Windows 7 x64、Visual Studio 2010、CUDA 7.5

内容的提问来源于stack exchange，提问作者Mikhail M

CUDA带宽测试疑问：计算操作为何未被内存延迟完全掩盖？

完整测试代码 #

测试环境 #

完整测试代码

测试环境