从缓冲区复制的速度差异

Question

我无法理解这些示例中复制速度变化如此之大的原因。我从他们那里得到了几乎不清晰的图像。 “更快”变体的计算时间也更快

没有帮助： 将'slow'变体的所有变量移入内核，各种内存标志几乎不会改变结果。

原来问题出在内核，但究竟是什么问题？

警告！我粘贴了整个文件

import pyopencl as cl
import numpy as np
from PIL import Image
import time

更快的变体。从缓冲区复制大约需要 0.15 秒

width = 800
height = 800
X = 0
Y = 0
R = 2
maxiter = 80000

xmin = X - R
xmax = X + R
ymin = Y - R
ymax = Y + R

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

r1 = np.linspace(xmin, xmax, width, dtype=np.float64)
r2 = np.linspace(ymin, ymax, height, dtype=np.float64)
q = r1 + r2[:, None] * 1j
q = np.ravel(q)
output = np.empty(width*height, dtype=np.uint8)

mf = cl.mem_flags
q_opencl = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=q)
output_opencl = cl.Buffer(ctx, mf.WRITE_ONLY, output.nbytes)

prg = cl.Program(ctx, """
    __kernel void mandelbrot(__global double2 *q,
                     __global uchar *output, ushort const maxiter)
    {
        int gid = get_global_id(0);
        double nreal, real = 0;
        double imag = 0;
        output[gid] = 0.0;

        int curiter = 0;
        for(curiter = 0; curiter < maxiter; curiter++) {
            nreal = real*real - imag*imag + q[gid].x;
            imag = 2* real*imag + q[gid].y;
            real = nreal;
            if (real*real + imag*imag > 4.0f){
                break;
            }
        }
        if (curiter < maxiter) {
            output[gid] = curiter*64;
        }
    }
    """).build()

prg.mandelbrot(queue, output.shape, None, q_opencl, output_opencl, np.uint16(maxiter))

t0 = time.time()
cl.enqueue_copy(queue, output, output_opencl).wait()
print(time.time()-t0, 'copy')

output = output.reshape((width, height))

较慢的变体。从缓冲区复制大约需要 0.78 秒

size = (800, 800)
X = 0
Y = 0
R = 2
maxiter = 80000

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

output = np.empty(size[0]*size[1], dtype=np.uint8)

mf = cl.mem_flags
output_cl = cl.Buffer(ctx, mf.WRITE_ONLY, output.nbytes)

prg = cl.Program(ctx, """
    __kernel void mandelbrot(
        __global uchar *out,
        int width,
        int height,
        double real,
        double imag,
        double const radius,
        int const maxiter) {
            int id = get_global_id(0);

            int i = id % width;
            int j = id / width;

            double window_radius = (width < height) ? width : height;
            double x0 = real + radius * (2 * i - (float)width) / window_radius;
            double y0 = imag - radius * (2 * j - (float)height) / window_radius;
            double x = 0;
            double y = 0;

            int n = 0;
            double x_temp = 0;
            for(n = 0; n < maxiter; n++)
            {
                x_temp = x*x - y*y + x0;
                y = 2 * x*y + y0;
                x = x_temp;
                if (x*x + y*y > 4.0f){
                    break;
                }
            }
            if (n < maxiter) {
                out[id] = n*64;
            }
            else {
                out[id] = 0;
            }
    }
""").build()

prg.mandelbrot(queue, output.shape, None,
                output_cl,
                np.int32(size[0]),
                np.int32(size[1]),
                np.float64(X),
                np.float64(Y),
                np.float64(R),
                np.int32(maxiter),
                )

t0 = time.time()
cl.enqueue_copy(queue, output, output_cl).wait()
print(time.time() - t0, 'copy')

output = output.reshape((size[1], size[0]))

Answer 1

我无法理解为什么在这些示例中复制速度变化如此之大。

原因是：你没有测量复制命令的时间。

您说“prg.mandelbrot() 执行内核并进行所有计算” - 这不是它的作用。它排队内核。然后你enqueue复制命令，然后你调用wait()。一些实现在入队后立即开始执行，但有些直到你调用 clFinish/clFlush/clWaitForEvents 才开始执行（最后一个是 PyOpenCL 的 Event.wait() 所做的 - 在你的代码中，cl.enqueue_copy() 返回一个事件）。

问题在于，您正在尝试使用主机 CPU 时间来测量 OpenCL (GPU) 时间，这属于初学者的错误。它永远不会起作用。您必须通过 OpenCL 事件分析来测量 GPU 上的时间。 Here's怎么办。

Answer 2

在第一个版本中，您将 80000 给 maxiter 这是无符号短。最大 65535。溢出并环绕到 14k ish 值。

你有 int 的第二个版本，它是 32 位的。 80k 正确通过。

14.5k / 80.0k 次迭代

0.15 / 0.78 秒

一致