在一个 GPU 上运行多个 CUDA 应用程序

Question

CUDA文档没有具体说明多少个CUDA进程可以共享一个GPU。例如，如果系统中只安装了一张GPU卡，同一用户启动多个CUDA程序，效果如何？能保证执行的正确性吗？这种情况下GPU如何调度任务？

Answer 1

来自独立主机进程的 CUDA 活动通常会创建独立的 CUDA

上下文，每个进程一个。 因此，从单独的主机进程启动的 CUDA 活动将在同一设备上的单独 CUDA 上下文中进行。

不同上下文中的 CUDA 活动将被序列化。 GPU 将执行一个进程中的活动，当该活动空闲时，它可以并且将会上下文切换到另一个上下文，以完成从另一个进程启动的 CUDA 活动。

未指定详细的上下文间调度行为。 （在单个 GPU 上运行多个上下文通常也不能违反基本 GPU 限制，例如设备分配的内存可用性。）请注意，上下文间切换/调度行为未指定，并且也可能因计算机设置而异。 随意观察或微基准测试可能表明较新设备上不同进程的内核可以同时运行（在 MPS 之外），但这并不正确。 较新的机器设置可能具有时间切片而不是循环行为，但这并没有改变这样的事实：在任何给定的时刻，只能运行来自一个上下文的代码。

这种情况的“例外”（来自独立主机进程的 GPU 活动的序列化）是 CUDA 多进程服务器。 简而言之，

MPS 充当“漏斗”来收集来自多个主机进程的 CUDA 活动，并像来自单个主机进程一样运行该活动。 主要好处是避免“内核序列化，否则内核可能会同时运行”。 规范的用例是启动多个 MPI 级别，所有这些级别都打算使用单个 GPU 资源。 请注意，上述描述适用于处于“默认”计算模式

的 GPU。 “独占进程”或“独占线程”计算模式下的 GPU 将拒绝在一台设备上创建多个进程/上下文的任何尝试。 在其中一种模式下，其他进程尝试使用已在使用的设备将导致 CUDA API 报告失败。 在某些情况下，可以使用

nvidia-smi 实用程序修改计算模式。

我是这个主题的新手。但我发现仅在一个 GPU 上模拟多个 GPU 是可能的。 “针对多个 GPU 进行开发将允许模型根据额外的资源进行扩展。如果在具有单个 GPU 的系统上进行开发，我们可以使用虚拟设备模拟多个 GPU。这可以轻松测试多 GPU 设置，而无需额外的资源。”

Answer 2

来源：

https://www.tensorflow.org/guide/gpu#allowing_gpu_memory_growth

也许使用这种技术，我们可以在这些虚拟 GPU 之一上运行每个模型（至少用于推理）。

我有一个关于类似主题的问题。我的研究领域是 GPU 分区，但我是这个领域的新手。我在一些学术论文中看到他们通过这个代码获取当前的SM ID：

Answer 3

__device__ __forceinline__ unsigned int get_smid() {
unsigned int smid;
asm("mov.u32 %0, %%smid;" : "=r"(smid));    
return smid;
} 

然后，通过一个简单的 if 条件，运行在所需 SM ID 上分配的块，并通过一个偏移值（具体取决于专用于来自其他块的每个内核函数调用线程的目标 SM 的数量），并在特定 SM 上运行该内核。通过这种方式，它们可以在一个 GPU 上并行运行多个内核。我的问题是这种方法是否正确，以及是否可以通过调用具有偏移值而不是线程的块来在协作网格组内的更高级别上实现它。

您没有指定您的用例，但只需编写不同的例程并在运行时选择要做什么就很简单了：

Answer 4

__device__ void RunProgram1() { 
    printf("T:%i B:%i program1\n", threadIdx.x, blockIdx.x); 
}

__device__ void RunProgram2() { 
    printf("T:%i B:%i program2\n", threadIdx.x, blockIdx.x); 
}


__global__ void StartDifferentTasks() {
    TaskID = blockIdx.x / 2; //assign two blocks to each task
    //you can partition this as needed
    switch (TaskID) {
        case 0: RunProgram1(); break;
        case 1: RunProgram2(); break;
        default: assert(false);
    }
}

int main() {
    StartDifferentTasks<<<4, 32>>>();
    cudaDeviceSynchronize(); //todo: add error checking  
}