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

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

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代码相结合，可以为应用程序提供强大的确定性特征。