在 CPU 上使用 OpenCL 将一个数组复制到另一个数组比 C++ 代码慢得多

Question

我比较了在 CPU 上运行的 OpenCL 代码的性能，该代码只是将数据从一个 2D 数组复制到另一个数组中，并将其转换为执行相同操作的纯 C++ 代码。我在 OpenCL 代码中使用了一个工作组来进行公平比较。我使用了英特尔的 OpenCL 驱动程序和英特尔编译器。OpenCL 代码大约比 CPU 代码慢 5 倍。编译器为复制循环提供以下消息：

loop was transformed to memset or memcpy.

关于如何使 OpenCL 代码与 C++ 代码同步的任何建议？

谢谢

OpenCL 主机代码：

#include <cstdio>
#include <cstdlib>
#include <iostream>
#include <fstream>
#include <cmath>
#include <ctime>
#include <CL/cl.hpp>

int main(int argc, char **argv)
{
    // Create the two input vectors
    const int N = 8192;
    double *in = new double[N*N]; 
    double *out = new double[N*N];

    for(int i = 0; i < N; i++)
        for (int j=0; j < N; j++) {
            in[i*N + j] = i + j;
            out[i*N + j] = 0.;
    }


    double time;
    std::clock_t start;
    int niter = 100;

    cl_int cl_err;

    std::vector<cl::Platform> platforms;
    cl_err = cl::Platform::get(&platforms);

    std::vector<cl::Device> devices;
    cl_err = platforms.at(1).getDevices(CL_DEVICE_TYPE_CPU,
                                        &devices);

    cl_context_properties context_properties[3] = {CL_CONTEXT_PLATFORM,
                                    (cl_context_properties)(platforms.at(1)()),
                                                   0};
    cl::Context context = cl::Context(devices, 
                                      context_properties, 
                                      NULL, NULL, &cl_err);

    cl::Buffer buffer_in = cl::Buffer(context, 
                                      CL_MEM_USE_HOST_PTR | CL_MEM_READ_ONLY,
                                      N*N*sizeof(double), 
                                      in, &cl_err);

    cl::Buffer buffer_out = cl::Buffer(context, 
                                       CL_MEM_USE_HOST_PTR | CL_MEM_WRITE_ONLY, 
                                       N*N*sizeof(double),
                                       out, &cl_err);

    cl::CommandQueue queue = cl::CommandQueue(context, devices.at(0), 0, &cl_err);

    std::ifstream sourceFile("vector_copy.cl");
    std::string sourceCode((std::istreambuf_iterator<char>(sourceFile)),
                            std::istreambuf_iterator<char>());
    cl::Program::Sources source(1, std::make_pair(sourceCode.c_str(),
                                sourceCode.length()+1));

    cl::Program program(context, source, &cl_err);

    cl_err = program.build(devices, NULL, NULL, NULL);

    cl::Kernel kernel(program, "vector_copy", &cl_err);

    cl_err = kernel.setArg(0, buffer_in); 
    cl_err = kernel.setArg(1, buffer_out);
    cl_err = kernel.setArg(2, N);

    cl::NDRange global(N);
    cl::NDRange local(N);

    start = std::clock();
    for (int n=0; n < niter; n++) {
        cl_err = queue.enqueueNDRangeKernel(kernel,
                                            cl::NullRange,
                                            global,
                                            local,
                                            NULL, NULL);

        cl_err = queue.finish();
    }

    time =  (std::clock() - start)/(double)CLOCKS_PER_SEC;
    std::cout << "Time/iteration OpenCL (s) = " << time/(double)niter << std::endl;

    return(0);
}

OpenCL内核代码：

__kernel void vector_copy(__global const double* restrict in, 
                          __global double* restrict out,
                         const int N) 
{

    int i = get_global_id(0);
    int j;

    for (j=0; j<N; j++) {
        out[j + N*i] = in[j + N*i];
    }

}

C++代码：

#include <cstdio>
#include <cstdlib>
#include <iostream>
#include <fstream>
#include <cmath>
#include <ctime>

const int N = 8192;

int main(int argc, char **argv)
{
    double *in = new double[N*N]; 
    double *out = new double[N*N];
    // Create the two input vectors
    for(int i = 0; i < N; i++)
        for (int j=0; j < N; j++) {
            in[j + N*i] = i + j;
            out[j + N*i] = 0.;
    }


    std::clock_t start;
    int niter = 100;

    start = std::clock();
    for (int n=0; n < niter; n++) {
        for (int i=0; i<N; i++)
            for (int j=0; j<N; j++) {
                out[j + N*i] = in[j + N*i];
            }

    }

    double time =  (std::clock() - start)/(double)CLOCKS_PER_SEC;
    std::cout << "Time/iteration C = " << time/(double)niter << std::endl;

    return(0);
}

Answer 1

英特尔 OpenCL 编译器能够跨工作组进行矢量化。例如，基本上单个函数在不同的 SSE 寄存器中同时运行 8 个线程。

您的特定内核不会这样做。但这并不重要。我使用 Visual Studio 2010 和最新的英特尔 OpenCL 应用程序测试了您的程序。我被迫将 N 从 8192 减少到 4096，因为我拥有的集成 GPU 将最大 OpenCL 缓冲区大小减少到 128MB，即使只使用 CPU。

我的结果：您的 OpenCL 内核为我提供了大约 6956MB/s 的带宽。一个简单更改的内核（使用 N*N 作为全局大小和 NULL 作为本地大小调用它，因为如果我们根本不关心本地内存，那么对于 CPU 的我们应该保持未定义）。

__kernel void vector_copy2(__global const double* restrict in, 
                      __global double* restrict out) 
{
  int i = get_global_id(0);
  out[i] = in[i];
}

得到大致相同的结果（7006MB/s）。该内核实际上是跨线程矢量化的，这可以使用英特尔 OpenCL 内核编译器进行验证。它为多个（如 4）生成一个内核，为单个线程生成一个内核。然后它只运行向量化内核，直到它必须为最后几个工作项运行单线程内核。

C++ 代码给出了 6494MB/s。所以它是完全一致的。我认为 ICC 甚至不可能让它快 5 倍。

我注意到在你的代码中你有platforms.at(1)，你电脑中的平台0是什么？

请记住，如果您根本不关心本地内存（您不在内核中调用 get_local_id），您应该将 enqueueNDRange 的本地大小视为一个简单的魔术参数。要么将其保留为 NULL，要么尝试找到产生最快结果的值。