共享 Tesla K20 的两个进程同时执行

Question

当我启动 2 个内核实例以便在共享 GPU 资源的同时同时运行时，我遇到了一种奇怪的行为。

我开发了一个 CUDA 内核，旨在在单个 SM（多处理器）中运行，其中线程执行多次操作（带有循环）。

内核准备只创建一个块，因此只使用一个 SM。

simple.cu

#include <cuda_runtime.h>
#include <stdlib.h>
#include <stdio.h>
#include <helper_cuda.h>
using namespace std;

__global__ void increment(float *in, float *out)
{
    int it=0, i = blockIdx.x * blockDim.x + threadIdx.x;
    float a=0.8525852f;

    for(it=0; it<99999999; it++)
             out[i] += (in[i]+a)*a-(in[i]+a);
}

int main( int argc, char* argv[])
{
    int i;
    int nBlocks = 1;
    int threadsPerBlock = 1024;
    float *A, *d_A, *d_B, *B;
    size_t size=1024*13;

    A = (float *) malloc(size * sizeof(float));
    B = (float *) malloc(size * sizeof(float));

    for(i=0;i<size;i++){
            A[i]=0.74;
            B[i]=0.36;
    }

    cudaMalloc((void **) &d_A, size * sizeof(float));
    cudaMalloc((void **) &d_B, size * sizeof(float));

    cudaMemcpy(d_A, A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, size, cudaMemcpyHostToDevice);

    increment<<<nBlocks,threadsPerBlock>>>(d_A, d_B);

    cudaDeviceSynchronize();

    cudaMemcpy(B, d_B, size, cudaMemcpyDeviceToHost);

    free(A);
    free(B);

    cudaFree(d_A);
    cudaFree(d_B);

    cudaDeviceReset();

    return (0);
}

所以如果我执行内核：

time ./simple

我明白了

real 0m36.659s user 0m4.033s sys 0m1.124s

否则，如果我执行两个实例：

time ./simple & time ./simple

我为每个进程得到：

real 1m12.417s user 0m29.494s sys 0m42.721s

real 1m12.440s user 0m36.387s sys 0m8.820s

据我所知，执行应该同时运行一次（大约 36 秒）。但是，它们的持续时间是基准时间的两倍。我们知道 GPU 有 13 个 SM，每个 SM 应该执行一个块，因此内核只创建 1 个块。

它们是在同一个 SM 中执行的吗？

他们不应该在不同的 SM 中同时运行吗？

已编辑

为了让我更清楚，我将附上并发执行的配置文件，从 nvprof 获得：

简介，一审

个人资料，二审

现在，我想向您展示同一场景的行为，但同时执行 matrixMul 示例的两个实例：

简介，一审

个人资料，二审

如您所见，在第一种情况下，一个内核等待另一个内核完成。而在第二种情况（matrixMul）中，来自两个上下文的内核同时运行。

谢谢。

Answer 1

当您使用同一个 GPU 运行两个单独的进程时，它们每个都有自己的上下文。 CUDA 不支持在同一设备上同时拥有多个上下文。相反，每个上下文都以未定义的方式竞争设备，并带有驱动程序级别的上下文切换。这就是为什么执行表现得好像进程是序列化的——实际上它们是序列化的，但在驱动程序而不是 GPU 级别。

有一些可用的技术（MPS、Hyper-Q）可以做你想做的事，但你尝试做的方式是行不通的。

---

编辑以响应您问题中的更新：

您使用 MatrixMul 示例添加的示例并未显示您的想法。该应用程序运行 300 个短内核并计算这 300 次运行的平均值的性能数字。您的分析显示已设置为非常粗略的时间尺度分辨率，因此看起来像是有一个长时间运行的内核启动，而实际上它是一系列运行时间非常短的内核。

为了说明这一点，请考虑以下几点：

这是在 Kepler 设备上运行的单个 MatrixMul 进程的正常分析运行。请注意，有许多单独的内核一个接一个地直接运行。

这些是在同一 Kepler 设备上同时运行的两个 MatrixMul 进程的分析跟踪：

请注意，每个进程的配置文件跟踪中存在间隙，这是两个进程之间发生上下文切换的地方。该行为与您的原始示例相同，只是在更精细的时间粒度上。正如几个不同的人在本讨论过程中多次重复的那样——CUDA 不支持在示例设备上同时使用标准运行时 API 的多个上下文。 MPS 服务器确实允许通过添加一个守护程序来实现这一点，该守护程序使用大型共享内部 Hyper-Q 管道重新实现 API，但您没有使用它，它与您在此显示的结果无关问题。