Improving Real-Time Performance with CUDA Persistent Threads (CuPer) on the Jetson TX2

发表时间: 2018-03

标题： 利用CUDA持久化线程（CuPer）提升Jetson TX2的实时性能
作者： Todd Allen
机构： Concurrent Real-Time

速读

一句话结论 本文针对硬实时应用中传统 CUDA 启动与同步机制带来的严重抖动问题，在 Jetson TX2 平台上提出了一种名为 CuPer 的持久化线程编程模型及 API，通过一次启动、持续运行的方式消除了调度开销，成功满足了低至 100µs 的严苛帧时长约束。

要解决什么问题 硬实时系统（如帧时长低至 100µs 到 1000µs 的应用）对性能确定性有极高要求，其核心痛点在于控制最大执行时间与平均执行时间之差（即抖动）。在传统的 CUDA 编程模型中，CPU 需要在每一帧异步启动核函数，并在计算完成后执行同步操作以获取结果。尽管可以通过零拷贝固定内存（无缓存的共享内存）绕开 CPU 与 GPU 之间的显式内存拷贝开销，但核函数的启动和同步操作本身在底层调度时会引入极大的不确定性。这种由 CUDA 运行时机制带来的长周期抖动尖峰，直接导致系统无法稳定满足亚毫秒级的硬实时帧时长约束。

怎么做的 核心思路是采用持久化线程模型：在应用程序启动时仅触发一次 CUDA 核函数启动，并让该核函数在整个应用生命周期内持续运行。由于绕过了传统模型中每一帧都要经历的底层启动与同步调度，这种方法从根本上消除了由 CUDA 运行时引入的性能抖动。为了将这种机制工程化，作者开发了 CuPer API，利用忙等待机制完成轻量级同步。关键设计由 CPU 端和 GPU 端两个控制部件构成。在 CPU 端，开发者实例化 `Cuper::Std::Cpu` 对象，在每一帧的主循环中，准备好输入数据后调用 `startCuda()` 发出信号，随后可执行其他 CPU 任务，最后调用 `waitForCuda()` 阻塞等待 GPU 结果。在 GPU 端，核函数内实例化 `Cuper::Std::Cuda1Block` 对象，进入死循环并调用 `waitForWork()` 进行忙等待；一旦接收到 CPU 信号便开始执行计算，完成后调用 `completeWork()` 通知 CPU。针对超出单个线程块容量的大规模工作负载，由于持久化线程受限于设备的最大常驻线程数（Jetson TX2 上为 4096 个），CuPer 提供了带线程组的单块核函数（将工作负载在软件层面划分为多个组，由单个块循环处理）以及混合模式（多块结合线程组）来扩展并行能力。此外，为了进一步压榨硬件利用率，CuPer 还设计了双缓冲接口，允许 CPU 在当前帧启动计算并消费上一帧的结果，从而形成两级流水线，掩盖 GPU 等待 CPU 分配任务时的空闲时间。

效果如何 实验在 NVIDIA Jetson TX2 平台（锁定最高时钟频率）和 RedHawk Linux 实时操作系统上搭建，通过核心隔离和最高实时优先级确保测试环境纯粹，并注入模拟的非实时负载以验证稳定性。对比基线主要包括：代表传统路线的启动/同步方法（分别测试了显式拷贝、零拷贝和统一虚拟内存三种机制），以及代表将调度下放给 GPU 的 CUDA 动态并行（CDP）方法。量化结果表明，在单块工作负载（如 32x32 矩阵乘法和 1024 元素向量求和）下，CuPer 持久化线程在平均耗时和最大耗时上均呈现压倒性优势，彻底消除了传统方法中出现的长周期抖动尖峰。在处理大规模工作负载时，最优策略因任务而异：对于 32K 向量增量和 64x64 矩阵乘法，CuPer 的单块多组方法在确定性上表现最佳；对于 32K 向量求和，混合 2 块实现取得了最优平衡；而在 128x128 矩阵乘法中，CDP 方法虽然抖动最大，但凭借极高的平均吞吐量，反而实现了最低的最大绝对耗时。该方法也存在明确的代价与局限：首先，强制 CPU 读写无缓存的零拷贝内存会带来额外的 CPU 内存 I/O 延迟；其次，由于核函数持续运行，系统无法动态切换异构核函数，只能通过在单一核函数内预置分支语句和按最大需求预分配显存来妥协；最后，忙等待机制会阻止 GPU 进入低功耗状态，且当配置的线程数逼近硬件常驻上限时，CuPer 的多块同步开销会急剧增加，甚至劣于传统方法。

A1 主要贡献

本文旨在解决将图形处理单元（GPU）用于具有亚毫秒级帧时长约束的硬实时应用时遇到的性能确定性差的问题。传统的CUDA启动/同步编程模型会引入显著的抖动，不适用于帧时长可能低至100µs的严苛实时系统。

为了应对这一挑战，本文的核心贡献如下：
1. 提出并验证了持久化线程编程模型：该模型通过在应用启动时仅启动一次CUDA核函数，并使其持续运行，从而避免了传统模型中导致性能不确定性的启动和同步操作。
2. 开发了一个简单的API（CuPer）：本文提出并介绍了一个名为“RedHawk CUDA Persistent Threads (CuPer)”的API，它为持久化线程编程风格提供了简洁的接口，简化了CPU和GPU之间的同步协调。
3. 提供了详尽的性能评测：在NVIDIA Jetson TX2平台和RedHawk Linux实时操作系统上，本文详细记录并分析了使用CuPer API的计时结果，并与传统的启动/同步方法进行了对比，证明了其在降低延迟和抖动方面的显著优势。
4. 给出了最佳实践建议：基于评测结果，本文为开发者如何根据不同的工作负载（从小规模单块任务到大规模多块任务）选择和实现最合适的持久化线程技术提供了具体建议。

A3 背景知识

实时性要求

CUDA 架构

A2 方法细节

CUDA 核函数启动/同步

__global__ void CudaKernel (float* A) { … }   
void cpuFunction (float* h_A)   
{ 
    cudaMemcpy(d_A, h_A, N); 
    CudaKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A); 
    cudaDeviceSynchronize(); 
    cudaMemcpy(h_A, d_A, N);   
}

__global__ void CudaKernel (float* A) { … }   
cpuFunction (float* h_A)   
{ 
    … 
    cudaHostGetDevicePointer(&d_A, h_A); 
    CudaKernel<<<blocksPerGrid, threadsPerBlock>>>(d_A); 
    cudaDeviceSynchronize();   
}

持久化线程

RedHawk CUDA持久化线程（CuPer）API

void cpuFunction (…)   
{ 
    Cuper::Std::Cpu p; 
    cudaHostGetDevicePointer(&d_A, h_A); 
    Persistent<<<blocksPerGrid, threadsPerBlock>>>(p.token(), d_A); 
    for (…) { 
        … initialize h_A … 
        p.startCuda(); 
        … possibly do unrelated CPU work … 
        p.waitForCuda(); 
        … use results in h_A … 
    } 
    p.terminateCuda();   
}

_global__ void Persistent (Cuper::Std::Token token, float* A) 
{ 
    Cuper::Std::Cuda1Block p(token); 
    for (…) { 
        p.waitForWork(); 
        if (p.isTerminated()) break; 
        … perform workload … 
        p.completeWork(); 
    } 
}

进入CUDA核函数后，它创建一个 Cuper::Std::Cuda1Block 类的对象 p，并使用从CPU传递的 token 值将其与对应的 Cuper::Std::Cpu 对象关联起来。该对象接收来自 Cpu 对应对象的命令，并在CUDA核函数内协调执行。核函数调用 p.waitForWork，这使CUDA GPU等待来自CPU的工作负载。从该调用返回后，可以安全地使用一致的用户缓冲区执行工作负载。计算完成后，它调用 p.completeWork 以此通知CPU。
* 终止检查：此外，如果需要终止核函数的能力，可以调用 p.isTerminated 来确定CPU是否已请求终止。
* 工作流程图：工作流程可以通过下图来可视化：

工作流程图
工作流程图

对角线箭头显示：
* Persistent<<<...>>> 启动了CUDA核函数
* p.startCuda 释放了阻塞的 p.waitForWork
* p.completeWork 释放了阻塞的 p.waitForCuda
* Cuda1Block与CudaMultiBlock：此示例展示了 Cuda1Block 类的使用。它适用于任何只需一个块即可执行的CUDA工作负载。单块方法通常比多块解决方案更高效。但是，如果需要多个块，可以用 CudaMultiBlock 替换 Cuda1Block。
* 零拷贝内存的使用：该方法的一个主要目标是避免在主循环内调用CUDA库函数，包括 cudaMemcpy。因此，在使用CuPer的所有情况下都使用了零拷贝固定内存。

性能

CUDA动态并行

_global__ void Algorithm (float* Result, float* A) { … } 
__global__ void Persistent (Cuper::Std::Token token, float* Result, float* A) 
{ 
    for (…) { 
        p.waitForWork(); 
        Algorithm<<<16, 1024>>>(Result, A); 
        cudaDeviceSynchronize(); 
        p.completeWork(); 
    } 
} 

void cpuFunction (…) 
{ 
    Persistent<<<1, 1>>>(p.token(), d_Result, d_A); 
    for (…) { 
        … 
        p.startCuda(); 
        p.waitForCuda(); 
        . 
    } 
}

单块工作负载

大规模工作负载

__global__ void Algorithm (Cuper::Std::Token token, float* Result, float* A)   
{ 
    int index = blockIdx.x * blockDim.x + threadIdx.x; 
    for (…) { 
        p.waitForWork(); 
        … 
        p.completeWork(); 
    }   
}   
void cpuFunction (…)   
{ 
    Algorithm<<<4, 1024>>>(p.token(), d_Result, d_A); 
    for (…) { 
        … p.startCuda(); 
        p.waitForCuda(); 
    }   
}

__global__ void Algorithm (Cuper::Std::Token token, float* Result, float* A) {
    for (…) { 
        p.waitForWork(); 
        for (unsigned int group = 0; group < 16; group++) { 
            int index = group * blockDim.x + threadIdx.x; 
            … 
        } 
        p.completeWork(); 
    } 
} 

void cpuFunction (…) { 
    Algorithm<<<1, 1024>>>(p.token(), d_Result, d_A); 
    for (…) { 
        … 
        p.startCuda(); 
        p.waitForCuda(); 
    } 
}

__global__ void Algorithm (Cuper::Std::Token token, float* Result, float* A)   
{ 
    for (…) { 
        p.waitForWork(); 
        for (unsigned int group = 0; group < 8; group++) { 
            int pseudoBlock = group * gridDim.x + blockIdx.x; 
            int index = pseudoBlock * blockDim.x + threadIdx.x; 
            … 
        } 
        p.completeWork(); 
    }   
}   

void cpuFunction (…)   
{ 
    Algorithm<<<2, 1024>>>(p.token(), d_Result, d_A); 
    for (…) { 
        … 
        p.startCuda(); 
        p.waitForCuda(); 
        … 
    }   
}

异构GPU行为

enum Command { MatMul_32x32x32, VecSum_32K };   

__global__ void Persistent (Cuper::Std::Token token, unsigned int* Command, float* Buf0, float* Buf1, float* Buf2);   
{ 
    for (…) { 
        p.waitForWork(); 
        unsigned int cmd = *Command; 
        switch (cmd) { 
            case MatMul_32x32x32: // C = Buf0, A = Buf1, B = Buf2 
                dim3 twoDimIdx(thread.Idx.x & 0x1f, threadIdx.x>>5); 
                matrixMultiply(twoDimIdx, Buf0, Buf1, Buf2); 
                break; 
            case VecSum_32K: // Result = Buf0, Vector = Buf1 
                vectorSum(Buf0, Buf1); 
                break; 
        } 
        p.completeWork();
    }
}

void cpuFunction (…)   
{ 
    Persistent<<<1, 1024>>>(p.token(), d_Command, d_Buf0, d_Buf1, d_Buf2); 
    for (…) { 
        … 
        initializeMatrixA(h_Buf1); 
        initializeMatrixB(h_Buf2); 
        *h_Command = MatMul_32x32x32; 
        p.startCuda(); 
        p.waitForCuda(); 
        useMatrixC(h_Buf0);

        … 
        initializeVector(h_Buf1); 
        *h_Command = VecSum_32K; 
        p.startCuda(); 
        p.waitForCuda(); 
        useResult(h_Buf0[0]); 
    }   
}

双缓冲

__global__ void Persistent (Cuper::DoubleBuffer::Token token, float* A0, float* A1) 
{ 
    Cuper::DoubleBuffer::Cuda1Block p(token); 
    for (…) { 
        p.waitForWork(); 
        unsigned int which = p.claimBuffer(); 
        float* A = which ? A1 : A0; 
        … 
        p.completeWork(); 
    } 
}

void cpuFunction (…)   
{ 
    Cuper::DoubleBuffer::Cpu p; 
    Persistent<<<…>>>(p.token(), d_A0, d_A1);

    // 用初始工作负载启动流水线
    unsigned int which = p.claimBuffer(); 
    float* h_A = which ? h_A1 : h_A0; 
    // ...为工作负载0初始化h_A...
    p.startCuda();

    for (unsigned int i = 1; …; i++) { 
        which = p.nextBuffer(); 
        float* h_A = which ? h_A1 : h_A0; 
        // ...为工作负载i初始化h_A...
        p.startCuda(); 
        … // 在这里可能做不相关的CPU工作
        p.waitForCuda(); 
        which = p.claimBuffer(); 
        float* h_A = which ? h_A1 : h_A0; 
        // ...消费工作负载(i-1)在h_A中的结果...
    } 

    // 从流水线中冲刷最后一个工作负载
    p.waitForCuda();   
    which = p.claimBuffer();   
    float* h_A = which ? h_A1 : h_A0;   
    // ...消费工作负载i在h_A中的结果...

    p.terminateCuda();
}

A4 实验环境

A4 实验结果

TABLE 1: API开销，1块（空）

FIGURE 1: API开销，1块（空）

TABLE 2: API开销，多块（空）

TABLE 3: API开销 + GPU内存I/O，1块（向量增量）

FIGURE 2: API开销 + GPU内存I/O，1块，最大值（向量增量）

TABLE 4: CPU内存I/O

FIGURE 3: CPU内存I/O：存储到输入缓冲区，最大值

FIGURE 4: CPU内存I/O：从输出缓冲区加载，最大值

TABLE 5: 32X32矩阵乘法 + API开销 + GPU内存I/O

FIGURE 5: 32X32矩阵乘法 + API开销 + GPU内存I/O

TABLE 6: 1K向量和 + API开销 + GPU内存I/O

FIGURE 6: 1K向量和 + API开销 + GPU内存I/O

TABLE 7: 32K向量增量 + API开销 + GPU内存I/O

FIGURE 7: 32K向量增量 + API开销 + GPU内存I/O

TABLE 8: 64X64矩阵乘法 + API开销 + GPU内存I/O

FIGURE 8: 64X64矩阵乘法 + API开销 + GPU内存I/O

TABLE 9: 128X128矩阵乘法 + API开销 + GPU内存I/O

FIGURE 9: 128X128矩阵乘法 + API开销 + GPU内存I/O

TABLE 10: 32K向量和 + API开销 + GPU内存I/O

FIGURE 10: 32K向量和 + API开销 + GPU内存I/O

A5 结论

本文展示的结果证明，对于具有严格亚毫秒级帧时长要求的硬实时应用，持久化线程编程模型是一种使用CUDA的可行方法。CuPer API提供了一种实现此模型的简便方式。将此方法与使用RedHawk Linux特性的CPU代码相结合，可以为应用程序提供强大的确定性特征。