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

Copying one array into another using OpenCL on CPU is much slower than C++ code

本文关键字:数组 另一个 复制 C++ 代码 一个 CPU OpenCL      更新时间:2023-10-16

我比较了在CPU上运行的OpenCL代码的性能,该代码只需将数据从一个2D数组复制到另一个2D阵列,而纯C++代码也可以做同样的事情。我在OpenCL代码中使用了一个工作组来进行公平的比较。我使用了英特尔的OpenCL驱动程序和英特尔编译器。OpenCL代码大约比CPU代码慢5倍。编译器为复制循环提供以下消息:

loop was transformed to memset or memcpy.

关于如何使用C++代码加快OpenCL代码的速度,有什么建议吗?

感谢

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);
}

英特尔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。所以这很符合要求。我认为国际刑事法院甚至不可能将速度提高5倍。

我在你的代码中注意到你有平台。在(1),你的计算机中平台0是什么?

请记住,如果您根本不关心本地内存(在内核中不调用get_local_id),则应该将enqueeNDRange的本地大小视为一个简单的魔术参数。要么将其保留为NULL,要么尝试查找生成最快结果的值。

OpenCL代码,即使经过优化,它仍然会执行1by1(逐个工作项)的复制。因为OpenCL编译器只允许在每个工作项的基础上进行优化。而C++的情况可能会被编译器优化为memcpy()调用(正如编译器告诉你的那样)。

如果禁用编译器优化,它将在GPU中执行得更快。

BTW这是有原因的吗?为此,C++中有memcpy(),OpenCL中有clEnqueueCopyBuffer()。我认为后一个是你应该使用的。