NVIDIA GPU线程数达核心数3倍时耗时不变的原因及算力疑问

测试代码 #

#include<stdio.h>
#include<stdint.h>
#include <chrono>
#include <cuda.h>

__global__ void test(int base, int* out)
{
    int curTh = threadIdx.x+blockIdx.x*blockDim.x;

    {
        int tmp = base * curTh;
#pragma unroll
        for (int i = 0; i<1000*1000*100; ++i) {
            tmp *= tmp;
        }
        out[curTh] = tmp;
    }
}

typedef std::chrono::high_resolution_clock Clock;
int main(int argc, char *argv[])
{
    cudaStream_t stream;
    cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

    int data = rand();
    int* d_out;
    void* va_args[10] = {&data, &d_out};
    int nth = 10;

    if (argc > 1) {
        nth = atoi(argv[1]);
    }

    int NTHREADS = 128;
    printf("nth: %d\n", nth);
    cudaMalloc(&d_out, nth*sizeof(int));

    for (int i = 0; i < 10; ++i) {
        auto start = Clock::now();

        cudaLaunchKernel((const void*) test,
                nth>NTHREADS ? nth/NTHREADS : 1,
                nth>NTHREADS ? NTHREADS : nth, va_args, 0, stream);

        cudaStreamSynchronize(stream);
        printf("use :%ldms\n", (Clock::now()-start)/1000/1000);
    }

    cudaDeviceReset();
    printf("host Hello World from CPU!\n");

    return 0;
}

运行结果 #

线程数=1 #

root@d114:~# ./cutest 1 
nth: 1
use :255ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms

线程数=13056（核心数×3） #

root@d114:~# ./cutest 13056
nth: 13056
use :272ms
use :223ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms
use :214ms

线程数=21760（核心数×5） #

root@d114:~# ./cutest 21760
nth: 21760
use :472ms
use :424ms
use :424ms
use :424ms
use :424ms
use :424ms
use :424ms
use :424ms
use :424ms
use :428ms

问题解答 #

1. 为什么线程数到核心数3倍时耗时不变？
   RTX 2080Ti采用Turing架构，包含68个流式多处理器（SM），每个SM有64个CUDA核心，总核心数4352。但GPU的并行能力不是仅由核心数决定，而是由SM的线程驻留能力和线程束调度机制共同决定：

• 每个SM最多可同时驻留2048个线程（64个线程束，每个线程束32个线程），通过快速切换线程束保持CUDA核心持续运行，避免空闲。
• 你的kernel是纯计算密集型任务，无内存依赖或同步操作，每个线程的计算完全独立。当线程数增加到13056时，这些线程可均匀分配到所有SM中，每个SM仅需承载约192个线程（13056÷68），远低于2048的驻留上限。此时所有CUDA核心都满负荷运行，同时完成所有线程的计算任务，因此总耗时与单线程时一致（单线程仅占用1个核心，其他核心空闲；多线程时所有核心并行工作，同时完成所有线程的计算）。

2. 这是否意味着NVIDIA GPU的计算能力是核心数的3倍？
   不是。GPU的最大计算能力由核心数、核心频率、指令吞吐量共同决定，线程数的增加只是让GPU的并行资源被充分利用。实际上，RTX 2080Ti的最大并发线程数可达139776（68×2048），远大于你测试的13056——只要线程数不超过这个上限，且任务是纯计算密集型，GPU都能以相同的耗时完成所有线程的计算（因为所有核心持续满负荷运行）。

3. 为什么线程数到21760时耗时翻倍？
   从测试结果看，21760线程的耗时约为13056线程的2倍，说明此时GPU无法一次性并行处理所有线程，需要分两批执行。可能的原因包括：

• kernel的寄存器使用量过高，导致每个SM只能驻留更少的线程，无法同时承载所有线程任务；
• 测试环境中GPU同时运行其他后台任务，导致调度效率下降。

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