标题： 利用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 背景知识

实时性要求

硬实时软件的核心需求：硬实时软件有两个主要需求：高性能和确定性。高性能通常通过平均执行时间或其倒数吞吐量来衡量，而确定性则通过抖动（最大执行时间与平均执行时间之差）来衡量。在实时系统中，优化平均时间通常比优化抖动更容易。
帧时长决定可接受值：可接受的平均时间和抖动值取决于应用的实时帧时长。对于软实时系统（如60Hz，帧时长16.666ms），标准的CUDA启动/同步方法可能足够。但对于硬实时系统（如1KHz/1000µs到10KHz/100µs），必须积极地最小化平均时间和抖动。
CUDA计算在帧中的占比：需要注意的是，CUDA计算几乎不可能占据整个帧的全部计算时间。能分配给CUDA计算的帧时间比例因应用而异。
非常规编程方法的采用：为了实现实时目标，一些通常被认为不理想的编程方法是完全可以接受的，例如：使用忙等待、强制系统资源在非实时应用需要时保持完全未使用状态，以及禁用电源管理。

CUDA 架构

CUDA编程模型：在CUDA编程模型中，最小的并行计算单元是线程。线程被组织成用户指定大小的线程块（或简称块）。块内的线程有多种机制可以进行快速通信和同步。块之间被设计为相互独立工作。这些块进一步组织成一个网格，网格包含了整个并行计算任务。在这些线程、块和网格中执行的源代码和目标代码被称为CUDA核函数（kernel）。
CUDA硬件模型：在CUDA硬件模型中，最小的计算单元是lane。Lanes被组织成称为warp的组，其大小在硬件中是固定的（在所有CUDA架构中都是32）。在这个架构层级，warp的功能类似于经典SIMD架构的变体，由warp解释指令并将其并行应用于warp内的所有lane。然而，CUDA架构支持因数据依赖的执行流导致的线程分歧——代价是并行度降低——因此该架构被称为SIMT（单指令多线程）。
线程到lane的映射：每个线程映射到一个lane。通常情况下（但并非总是），一个块包含的线程数多于一个warp内的lane数，因此CUDA硬件会将块细分为多个warp。如果一个块小于一个warp，或者块的大小不是warp大小的整数倍，那么一些lane将会处于非活动状态。
流式多处理器（SM）：在warp的层级之上是流式多处理器（SM）。SM之间大部分是独立运行的。块由CUDA硬件和设备驱动程序映射到单个SM上。最后，所有SM的集合构成了设备。
常驻Warp与延迟隐藏：一个SM可以支持比其能同时执行的warp数量更多的warp。以Jetson TX2为例，该架构有2个SM，每个SM能同时执行2个warp，每个warp包含32个lane。因此，整个设备能同时执行22=4个warp，即2232=128个lane。然而，该设备支持26432=4096个常驻线程，即264=128个常驻warp。其原因是它对停顿（stall）的处理方式。如果一个正在活跃执行的warp执行了某个需要超过一个周期才能完成的操作（通常是全局内存加载），该warp就会被停顿并从活跃执行中移除。在下一个周期，如果有一个可用的不同warp，就会执行它。SM利用常驻warp集合提供的额外并行性来隐藏停顿。值得注意的是，上下文切换完全在CUDA硬件内部处理，并且可以在一个周期到下一个周期之间发生，没有性能损失。

A2 方法细节

CUDA 核函数启动/同步

传统CUDA执行流程：执行CUDA GPU工作的传统方法是首先使用CUDA C（实际上是C++的变种）编写一个CUDA核函数。在CUDA设备上执行的代码由CUDA函数限定符 __global__ 标识。CPU将数据缓冲区复制到设备，并使用特殊的 <<<...>>> 语法执行一个类似C语言的调用，以表明它正在执行CUDA启动。CUDA启动是异步的，CPU在之后会继续执行。这允许它并行执行不相关的工作。当CPU准备好等待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);   
}

零拷贝固定内存技术：然而，像Jetson TX2这样的系统，CPU和GPU共享单一内存库，因此内存复制通常效率低下。可以改用一种称为零拷贝固定内存（zero-copy pinned memory）的技术。这允许CPU和GPU使用相同的缓冲区。唯一的注意事项是该内存是无缓存的（uncached）。这种内存使用 cudaHostAlloc 配合 cudaHostAllocMapped 标志进行分配。使用此模型，启动/同步的源代码可能如下形式：

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

传统方法的局限性：这种编程风格简单而强大。不幸的是，启动（launch）和同步（synchronize）操作的确定性水平对于许多硬实时应用来说是不可接受的。

持久化线程

核心思想：持久化线程模型通过在应用开始时仅启动一次CUDA核函数，并使其一直运行直到应用结束，从而避免了这些确定性问题。
方法缺点：这种方法有一些缺点：
- 无法在整个应用程序中启动异构的核函数。核函数始终在运行。不过，可以使用一个switch语句从多个预定义的操作中进行选择。
- 当等待CPU的工作负载时，核函数必须执行忙等待（busy wait）。在硬实时系统中，这并非罕见解决方案，但仍值得注意。它可能比让CUDA GPU空闲消耗更多能量。同样，CPU在等待GPU完成时也必须执行忙等待。
- 核函数使用的块和线程数量受限于设备支持的常驻线程数（在Jetson TX2上为4096）。
实现多样性：持久化线程的概念非常广泛，允许多种CPU和GPU之间的协调方法，以及在GPU上分配工作的方法。本文描述的方法紧密模仿了启动/同步方法，它被有意设计得简单以提供最佳的确定性。
同步机制需求：由于没有显式的启动来开始一个工作负载，也没有同步来检测工作负载何时完成，因此需要其他的同步原语。这可以手动实现。然而，RedHawk CUDA持久化线程（CuPer）API提供了一个简单的API来抽象这些原语操作。

RedHawk CUDA持久化线程（CuPer）API

API接口：标准接口在 <cuper.h> 头文件中提供。所有元素都在 Cuper::Std 命名空间内声明。其中定义了三个类：Cpu、Cuda1Block 和 CudaMultiBlock。Cpu 类的对象在CPU源代码中创建。一个典型的用法如下所示：

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();   
}

CPU端工作流程：在此示例中，名为 Persistent 的CUDA核函数在进入主循环前只启动一次。任何在正常操作期间用于来回传递用户数据的缓冲区都必须在此时指定。此外，Cuper::Std::Cpu 对象提供了一个 token()，其值也必须传递给核函数。p 对象用于控制CUDA核函数内工作负载的执行。在主循环内，p.startCuda 通知CUDA GPU输入缓冲区已准备好，应开始执行其工作负载，这类似于CUDA核函数启动。p.waitForCuda 使CPU等待GPU上的工作完成，这类似于CUDA同步。
终止流程：如果希望主循环能够退出，可以调用 p.terminateCuda 来请求终止。
GPU端工作流程：同时，此示例中的 Persistent 核函数可能形式如下：

_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的所有情况下都使用了零拷贝固定内存。

性能

性能考量方面：为了衡量CuPer持久化线程的性能，需要考虑多个方面：
- CuPer API开销（即同步和栅栏时间）
- GPU内存I/O
- CPU内存I/O
- 算法（超出了本文的范围，但在“单块工作负载”部分提供了一些示例）
性能叠加的复杂性：然而，需要注意的是，计算一个完整工作负载的时间，并不能简单地将这些时间相加。由于CUDA使用常驻warp来隐藏停顿，许多这些方面的时间将会交织在一起。

CUDA动态并行

一种不同的持久化核函数方法：还有另一种使用持久化核函数的方法，与前面的例子有很大不同。在这种方法中，持久化核函数只执行单个线程。在该线程从 p.waitForWork 返回后，它实际上根本不执行工作负载。相反，它使用CUDA动态并行（CDP）启动另一个核函数，等待该核函数完成，然后调用 p.completeWork。
与传统模型的比较和缺点：这与传统的CUDA启动/同步编程模型非常相似。主要区别在于它将启动和同步操作移到了CUDA GPU上。不幸的是，它会遭受长周期的抖动尖峰。因此，对于单块工作负载，它与上面描述的持久化线程方法相比没有竞争力。注意，这种方法需要比正常情况稍长的时间才能使其计时行为稳定下来。在下一节呈现的结果中，前两次迭代的计时被丢弃。
代码示例：使用此方法的代码形式如下：

_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(); 
        . 
    } 
}

单块工作负载

示例工作负载：为了展示单块实现的性能，本文提供了两个常见的工作负载：
- 32x32矩阵乘法：该任务可以完美地适配一个块中的32*32=1024个线程。这是一个标准的基于共享内存的实现。注意，在启动/同步的情况下，性能需要很长时间才能稳定。因此，对于这些情况，不是丢弃第一次迭代，而是丢弃前512次迭代。
- 1024元素向量求和：这是一个归约（reduction）算法的例子。这是一个使用逐次减小步长的标准算法，但当步长小到足以容纳在warp的32个lane内时，会切换到使用warp级原语。

大规模工作负载

处理大规模工作负载的持久化线程技术：单块工作负载提供了良好的性能，可以集成到紧凑的实时帧需求中。然而，有些情况下需要更大的工作负载，并且帧需求可以容纳它们。有多种持久化线程技术可用于更大的工作负载：
- 多块核函数
- 带线程组的单块核函数
- 混合模式：带线程组的多块核函数
- 执行CUDA动态并行（CDP）启动/同步的单线程核函数
多块核函数：多块核函数方法在前面的一些案例中已经探讨过。它仅仅是配置核函数使用超过1个块。然而，对于持久化线程方法，线程数存在一个硬性上限：最大常驻线程数，在Jetson TX2上是4096。因此，如果一个块被配置为使用其最大尺寸1024，那么块的数量不能超过4个。所以这种方法必然是受限的。而且，根据配置不同，它还会带来性能损失。使用这种技术的代码安排如下：

__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(); 
    }   
}

带线程组的单块核函数：另一种方法是使用1个块，但将工作负载划分为多个线程组。线程组是一个软件概念，取代了CUDA硬件的块概念。并且，CUDA块的集合将由CUDA硬件调度器调度，而每个线程组必须由单个正在运行的块手动执行。虽然可以实现更复杂的调度方法，但它们会产生性能开销，因此对于实时应用程序来说，最简单的方法是迭代遍历线程组。使用线程组，对于一个16K的例子，代码安排如下：

__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(); 
    } 
}

混合模式：带线程组的多块核函数：混合方法是指持久化核函数配置了多个块，但这些块不足以处理整个工作负载。因此，也使用了线程组。在这种情况下，每个线程组包含网格中的所有块。所以，对于一个16K的例子，配置可能是1024线程/块 * 2块/组 * 8组 = 16K。在组之间划分块有几种不同的方法，但使用这种配置，代码可能安排如下：

__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(); 
        … 
    }   
}

使用CUDA动态并行的单线程核函数：尽管前面描述的使用CUDA动态并行的方法对于单块工作负载没有竞争力，但对于更大的工作负载它可能是有效的，即使它会遭受长周期的抖动尖峰。在这种情况下，CUDA硬件调度器负责调度子核函数的所有块。

异构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]); 
    }   
}

此方法的限制：这种方法的限制是：
- 持久化核函数只能有一个块和网格配置。如果某个特定函数期望其他配置，必须使用胶水代码进行转换。在上面的例子中，matrixMultiply 期望二维的 threadIdx 坐标。twoDimIdx 从持久化核函数的一维坐标进行转换。然后 matrixMultiply 必须使用它而不是普通的 threadIdx。
- 缓冲区必须能够容纳持久化核函数的所有可能行为。例如，在上面的情况中，需要3个缓冲区，并且每个缓冲区的大小是每个行为需求的最大值。
- 如果任何行为需要使用CUDA <<<...>>> 启动配置中的第三个参数来动态分配共享内存，它必须是所有行为所需的最大值。

双缓冲

提高GPU利用率：之前介绍CuPer时提到，在执行GPU工作的同时可以执行不相关的CPU工作。这种额外的并行化程度是有用的。然而，当CPU没有明确分配工作负载时，CUDA GPU仍然处于空闲状态。如果应用程序可以接受1帧的响应延迟，可以通过双缓冲来改进这一点。
双缓冲概念：双缓冲的思想是，CPU最初在GPU上启动一个工作负载。然后，在每一帧中，它启动一个工作负载，然后使用前一个工作负载的结果。例如，在CPU的第101次迭代中，它启动工作负载101，并使用工作负载100的结果。这可以被看作是一个2级流水线。
CuPer双缓冲接口与流程：CuPer在 Cuper::DoubleBuffer 命名空间内提供了一个备用接口来支持这一点。它与 Cuper::Std 下的接口非常相似，但增加了一些额外的函数。工作流程可以用下面这张关于工作负载102和103附近的流程图来可视化：

双缓冲工作流程
方法要点：
- 只要CPU能够跟上已完成的工作负载，GPU就可以保持忙碌。
- 在CPU上，每个工作负载都在一帧中启动，并在下一帧中使用。
- 在GPU上，每个工作负载都正常处理。
- 为了支持这一点，必须有两套缓冲区，并且它们的使用在帧与帧之间交替。（这就是“双缓冲”这个术语的来源。）Cuper::DoubleBuffer 实现会自动跟踪正确的缓冲区。
代码示例：一个 Cuper::DoubleBuffer 的使用示例可能如下，与 Cuper::Std 的不同之处已高亮显示：

__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 实验环境

硬件配置：
- 平台：NVIDIA Jetson TX2
- 系统时钟：CPU、GPU (gp10b) 和内存 (emc) 的时钟被设置为允许的最大值。
- 风扇：系统风扇也设置为最大值。
软件与系统配置：
- 操作系统：RedHawk Linux RTOS
- CUDA版本：8.0
- 实时调度策略：
  - 使用 SCHED_FIFO 调度器。
  - 使用实时优先级80。
  - 将测试进程绑定到特定的CPU核心3。
  - 通过RedHawk Linux的特定功能，将CPU 3与其他所有进程、中断和ltmr中断隔离（Shielding）。
- 其他配置：
  - 禁用了X服务器（Xorg / lightdm）。
  - 对于涉及CUDA启动/同步的案例，调用了cudaSetDeviceFlags并设置了cudaDeviceScheduleSpin标志。
测试方法：
- 非实时负载模拟：使用开源的stress-1.0.4软件包，配置了5个CPU worker、5个VM worker、5个I/O worker和5个HDD worker（命令为stress -c 5 -m 5 -i 5 -d 5）来模拟非实时开销。
- 计时：所有测试用例的计时都持续了至少2小时，以确保平均值稳定并对最大值有合理的置信度。对于某些特定的CuPer持久化线程案例，运行时间更长以获得更高的置信度。
- 数据预处理：每个测试的第一次迭代的计时数据被丢弃，因为预计它会因加载各种文本页和数据对象到缓存而耗时过长。

A4 实验结果

实验一：API开销（空工作负载）
- 实验内容：在没有实际计算任务的情况下，测量CuPer API与传统启动/同步方法的自身开销。
- 实验结果：
  - 对于单块（1-block）情况（表1，图1），传统启动/同步方法的抖动非常大（最大时间远超平均时间）。CuPer方法的抖动极小，但其平均时间会随着线程数的增加而线性增长。
  - 对于多块（multi-block）情况（表2），CuPer的开销显著增加，尤其是在需要网格范围同步的completeWork调用中。当线程总数达到4096时，CuPer的平均性能甚至劣于启动/同步方法。
- 结论：对于小工作负载，推荐使用单块配置和Cuda1Block。
实验二：API开销 + GPU内存I/O（向量增量）
- 实验内容：在API开销的基础上，加入了一个简单的向量增量任务，以衡量包含GPU内存I/O时的性能。比较了CuPer（零拷贝）、启动/同步（零拷贝、cudaMemcpy、UVM）等多种方法。
- 实验结果（表3，图2）：即使包含了GPU内存I/O时间，持久化线程方法在平均时间和抖动方面仍然表现出色。UVM（统一虚拟内存）方法表现出长周期的抖动问题。
- 结论：在包含GPU内存I/O的场景下，持久化线程方法依然保持了其性能优势。
实验三：CPU内存I/O
- 实验内容：评估CPU访问零拷贝固定内存（uncached）所带来的开销。由于此开销分散在代码中难以隔离，实验使用memcpy来模拟CPU的读写操作。
- 实验结果（表4，图3，图4）：与访问普通缓存内存（malloc分配）相比，CPU访问零拷贝内存的速度较慢。这是使用零拷贝技术的固有成本，而非持久化线程方法本身造成的。
- 结论：开发者在评估整体实时性能时，必须考虑到这一部分由CPU内存I/O引入的额外时间成本。
实验四：单块工作负载性能
- 实验内容：比较了CuPer、启动/同步和CuPer+CDP（CUDA动态并行）在两个典型单块任务上的性能：32x32矩阵乘法和1K向量求和。
- 实验结果（表5, 6，图5, 6）：在这两个任务中，CuPer持久化线程方法在平均时间和抖动方面都显著优于启动/同步方法和CDP方法。
- 结论：对于单块工作负载，CuPer是性能最佳的选择。
实验五：大规模工作负载性能
- 实验内容：比较了四种处理大规模工作负载的技术（多块、单块多组、混合模式、CDP）在不同任务上的表现。
- 32K向量增量（表7，图7）：CuPer单块多组（1-block multi-group）方法在平均性能和确定性上表现最佳。
- 64x64矩阵乘法（表8，图8）：CuPer单块多组方法再次胜出。
- 128x128矩阵乘法（表9，图9）：一个有趣的案例，CDP方法虽然确定性最差（抖动最大），但其平均性能极好，导致其观察到的最大耗时反而最低。
- 32K向量求和（表10，图10）：混合2块（hybrid 2-block）实现表现最佳，它成功地利用了第二个SM带来的性能增益，且收益超过了使用多块带来的开销和抖动成本。
- 结论：对于大规模工作负载，不存在一种普遍最优的方法。最佳实现取决于具体的算法，需要通过实验来确定。