为什么访问单个 GPU 的多个进程会提高性能？

Question

我正在结合使用 PyOpenCl 和 Python 3.7.

当使用多个进程调用同一个内核时，每个进程都有自己的上下文指向同一个 GPU 设备，我得到的性能改进几乎与进程数量成线性关系。

我可以想象并行进程的执行使得一些重叠传输成为可能，其中进程 A 的内核被执行，而进程 B 将数据发送到图形卡。但这不应该是性能提升的原因。

附件是一个代码示例，我在其中实现了一个虚拟应用程序，其中一些数据被解码。

当设置 n_processes=1 我得到大约 12 Mbit/sec，而当设置 n_processes=4 我得到 45 Mbit/sec.

我正在使用单个 AMD Radeon VII 显卡。

有人对这种现象有很好的解释吗？

更新：我使用 CodeXL 分析了脚本。似乎在内核执行之间浪费了很多时间，并且多个进程能够利用它。

import logging
import multiprocessing as mp
import pyopencl as cl
import pyopencl.array as cl_array
from mako.template import Template
import numpy as np
import time

logging.basicConfig(level=logging.INFO,
                    format='%(asctime)s %(process)d %(levelname)-8s [%(filename)s:%(lineno)d] %(message)s')

kernelsource = """
float boxplus(float a,float b)
{
float boxp=log((1+exp(a+b))/(exp(a)+exp(b)));
return boxp;
}

void kernel test(global const float* in,
                global const int* permutation_vector,
                global float* out)
{
 int gid = get_global_id(0);
 int p = gid; // permutation index
 float t = 0.0;
 for(int k=1; k<10;k++){
    p = permutation_vector[p];
    t= boxplus(in[p],in[gid]);
 }
 out[gid] = t;
}
"""


class MyProcess(mp.Process):
    def __init__(self, q):
        super().__init__()
        self.q = q

    def run(self) -> None:
        platform = cl.get_platforms()
        my_gpu_devices = [platform[0].get_devices(device_type=cl.device_type.GPU)[0]]
        ctx = cl.Context(devices=my_gpu_devices)
        queue = cl.CommandQueue(ctx)
        tpl = Template(kernelsource)
        rendered_tp = tpl.render()
        prg = cl.Program(ctx, str(rendered_tp)).build()

        size = 100000  # shape of random input array
        dtype = np.float64
        output_buffer = cl_array.empty(queue, size, dtype=dtype)
        input_buffer = cl_array.empty(queue, size, dtype=dtype)

        permutation = np.random.permutation(size)
        permutation_buffer = cl_array.to_device(queue, permutation.astype(np.int))

        def decode(data_in):
            input_buffer.set(data_in)
            for i in range(10):
                prg.test(queue, input_buffer.shape, None,
                         input_buffer.data,
                         permutation_buffer.data,
                         output_buffer.data)
            queue.finish()
            return output_buffer.get()

        counter = 1
        while True:
            data_in = np.random.normal(size=size).astype(dtype)
            data_out = decode(data_in)
            if counter % 100 == 0:
                self.q.put(size * 100)
                counter = 1
            else:
                counter += 1


def run_test_multi_cpu_single_gpu():
    q = mp.Queue()
    n_processes = 4
    for i in range(n_processes):
        MyProcess(q).start()
    t0 = time.time()
    symbols_sum = q.get()
    i = 0
    while True:
        i += 1
        print('{} Mbit/sec'.format(1 / 1e6 * symbols_sum / (time.time() - t0 + 1e-15)))
        symbols = q.get()
        symbols_sum += symbols

if __name__ == '__main__':
    run_test_multi_cpu_single_gpu()

Answer 1

内核循环工作太少。它必须几乎可以与内核启动开销相媲美。内核启动开销也与 Python 中的函数调用开销相当。

 for(int k=1; k<10;k++){
    p = permutation_vector[p];
    t= boxplus(in[p],in[gid]);
 }

这个延迟可能隐藏在另一个进程的内核启动延迟后面，它的内核启动延迟可能隐藏在第三个进程的函数调用开销后面。而 GPU 可以承担更多，复杂度为 O(N) 的 for 循环只有 10 次循环。即使是低端 GPU 也会因为至少数千次复杂度为 O(N*N) 的迭代而饱和。

缓冲区 read/writes 和计算也如您所说重叠。

So if the kernel takes all time in that profiling window, there is no capacity left on the graphic card?

GPU 也可以重叠多个计算，如果它有能力并且每个工作足够小以让一些运行中的线程为其他线程保留。运行中的线程数可高达 40*shaders。每个周期（或几个周期）40*3840 = 153600 条指令 issued/pipelined 或者说 3.46 TFLOPS。

3.46 TFLOPS 每个 64 位数据元素甚至有 1000 FLOP，它可以以 3.46 GB/s 速率传输数据。这没有在内核中流水线化任何内容（读取元素 1、计算、写入结果、读取元素 2）。但它是流水线，在开始第一个元素计算后，下一批项目被映射到相同的着色器上，加载新数据，它可能需要数百GB/s，这比 PCI-e 带宽

]

而且 CPU 无法 preprocess/post 以该速度处理。所以当有多个进程时，缓冲区副本和 CPU 作为瓶颈相互隐藏。

为什么访问单个 GPU 的多个进程会提高性能？

Why do multiple processes accessing single GPU increase performance?

opencl

pyopencl