【问题标题】:Why do multiple processes accessing single GPU increase performance?为什么多个进程访问单个 GPU 会提高性能?
【发布时间】:2019-10-04 04:16:11
【问题描述】:

我将 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()

【问题讨论】:

  • GPU 专为大规模并行而设计;这应该不足为奇。
  • 那么单个进程是不能耗尽显卡的?
  • 我想这取决于进程具体要求显卡做什么。如果它是像 Halo 这样的游戏,那么就会有相当多的处理过程。问题并不像你想象的那么简单。
  • 你是对的。通过增加内部的 for 循环运行来更改内核工作时,单个进程和多个进程之间的性能差异很小。但是,如果内核不是那么复杂,多个进程能够隐藏调用内核、传输数据等引入的延迟,我说得对吗?
  • 这很难说。如果有针对此类事情的分析器,它将准确地告诉您时间花费在哪里。

标签: opencl pyopencl


【解决方案1】:

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

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

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

正如你所说,缓冲区读/写和计算也是重叠的。

因此,如果内核在该分析窗口中花费了所有时间,则没有 显卡还剩多少容量?

GPU 也可以重叠多个计算,如果它有能力并且每个工作足够小,可以让一些正在运行的线程留给其他线程。运行中线程的数量可以高达 40*shaders。 40*3840 = 每个周期(或几个周期)发出/流水线化 153600 条指令,或者说 3.46 TFLOPS。

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

CPU 也无法以该速率预处理/后处理。所以当有多个进程时,缓冲区副本和 CPU 作为瓶颈相互隐藏。

【讨论】:

    猜你喜欢
    • 1970-01-01
    • 2012-03-05
    • 2010-12-09
    • 2021-01-30
    • 1970-01-01
    • 2012-10-03
    • 1970-01-01
    • 2017-11-19
    • 2013-07-21
    相关资源
    最近更新 更多