logo标签列表

在多个GPU上运行OpenCL内核？

openclmultiple-gpu
3
我现在已经编写了几个算法，可以在一个GPU上并行运行，但它们都有同样的问题。当我尝试在多个GPU上执行它们时（例如3个），代码在一个GPU上执行的时间与在3个GPU上执行的时间完全相同（没有加快）。我尝试使用更多数据来执行，尝试执行不同的任务，但什么也没有帮助。最后，我试图运行最简单的任务，如元素求和，仍然出现这个可怕的错误。这就是为什么我不相信这是特定算法的问题，我觉得我的代码存在错误（甚至是在将代码并行化到多个GPU上的方法上）。
这是我Parallel.cpp类的头文件：
#ifndef PARALLEL_H
#define PARALLEL_H

#define __NO_STD_VECTOR // Use cl::vector and cl::string and
#define __NO_STD_STRING // not STL versions, more on this later
#include <CL/cl.h>

class Parallel
{
    public:
        Parallel();
        int executeAttachVectorsKernel(int*, int*, int*, int);
        static void getMaxWorkGroupSize(int*, int*, int*);
        virtual ~Parallel();
    protected:
    private:
        char* file_contents(const char*, int*);
        void getShortInfo(cl_device_id);
        int init(void);
        cl_platform_id platform;
        cl_device_id* devices;
        cl_uint num_devices;
        cl_command_queue* queues;
        int* WGSizes;
        int* WGNumbers;
        cl_context context;
        cl_program program;
        cl_kernel kernel;
        cl_mem input1;
        cl_mem input2;
        cl_mem output;
};

#endif // PARALLEL_H

这里是初始化方法init：

int Parallel::init() {
cl_int err;

//Connect to the first platfrom
err = clGetPlatformIDs(1, &platform, NULL);
if (err != CL_SUCCESS) {
    cerr << "Error occured while executing clGetPlatformIDs" << endl;
    return EXIT_FAILURE;
}

//Get devices number
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &num_devices);
if (err != CL_SUCCESS) {
    cerr << "Error: Failed to create a device group:" << endl;
    return EXIT_FAILURE;
}

cout << "NUM DEVICES =" << num_devices << endl;

devices = new cl_device_id[num_devices];
//Get all the GPU devices
err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, num_devices, devices, NULL);

//Create one context for all the devices
context = clCreateContext(NULL, num_devices, devices, NULL, NULL, &err);
if (!context) {
    cerr << "Error: Failed to create a compute context!" << endl;
    return EXIT_FAILURE;
}

queues = new cl_command_queue[num_devices];
WGNumbers = new int[num_devices];
WGSizes = new int[num_devices];


for(int i = 0; i < num_devices; i++) {
    //Create a command queue for every device
    queues[i] = clCreateCommandQueue(context, devices[i], 0, &err);
    if (!queues[i]) {
        cerr << "Error: Failed to create a command commands!" << endl;
        return EXIT_FAILURE;
    }

    cl_ulong temp;
    clGetDeviceInfo(devices[i], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(temp), &temp, NULL);
    WGSizes[i] = (int)temp;

    clGetDeviceInfo(devices[i], CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(temp), &temp, NULL);
    WGNumbers[i] = (int)temp;
}

//Translate kernel code into chars
int pl;
size_t program_length;
string path = "./kernel/kernel_av.cl";

char* cSourceCL = file_contents(path.c_str(), &pl);
program_length = (size_t)pl;

//Create a program
program = clCreateProgramWithSource(context, 1,
                  (const char **) &cSourceCL, &program_length, &err);

if (!program) {
    cerr << "Error: Failed to create compute program!" << endl;
    return EXIT_FAILURE;
}

//Create an executable
err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
if (err != CL_SUCCESS)
{
    size_t len;
    char buffer[2048];

    cerr << "Error: Failed to build program executable!" << endl;
    exit(1);
}

// Create the compute kernel in the program
kernel = clCreateKernel(program, "calculate2dim", &err);
if (err != CL_SUCCESS)
{
    cerr << "Error: Failed to create compute kernel!" << endl;
    exit(1);
}
}

执行内核的方法在这里：
int Parallel::executeAttachVectorsKernel(int* data1, int* data2, int* results, int vectors_num) {

cl_int err;
size_t global;  // global domain size for our calculation
size_t local;   // local domain size for our calculation

int partition = vectors_num/num_devices;
unsigned int count = partition;
input1 = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * count, NULL, NULL);
input2 = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(int) * count, NULL, NULL);
output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(int) * count, NULL, NULL);
if (!input1 || !input2 || !output) {
    cerr << "Error: Failed to allocate device memory!" << endl;
    exit(1);
}

int** data1_apart = new int*[num_devices];
int** data2_apart = new int*[num_devices];
int** results_apart = new int*[num_devices];

for(int i = 0; i < num_devices; i++) {
    cout << "Executing parallel part on GPU " << i + 1 << endl;
    cout << "Partition size = " << partition << endl;
    data1_apart[i] = new int[partition];
    data2_apart[i] = new int[partition];
    results_apart[i] = new int[partition];

    for(int j = i*partition, k = 0; k < partition; j++, k++) {
        data1_apart[i][k] = data1[j];
        data2_apart[i][k] = data2[j];
    }

    //Transfer the input vector into device memory
    err = clEnqueueWriteBuffer(queues[i], input1,
                               CL_TRUE, 0, sizeof(int) * count,
                               data1_apart[i], 0, NULL, NULL);

    err = clEnqueueWriteBuffer(queues[i], input2,
                               CL_TRUE, 0, sizeof(int) * count,
                               data2_apart[i], 0, NULL, NULL);

    if (err != CL_SUCCESS)
    {
        cerr << "Error: Failed to write to source array!" << endl;
        exit(1);
    }

    int parameter4 = count/WGNumbers[i];

     //Set the arguments to the compute kernel
    err = 0;
    err  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &input1);
    err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &input2);
    err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &output);
    err |= clSetKernelArg(kernel, 3, sizeof(int), &parameter4);
    if (err != CL_SUCCESS)
    {
        cerr << "Error: Failed to set kernel arguments! " << err << endl;
        exit(1);
    }

    global = WGNumbers[i];
    local = WGSizes[i];

    if(local > global) {
        local = global;
    }
    cout << "global = " << global << " local = " << local << endl;

    err = clEnqueueNDRangeKernel(queues[i], kernel,
                                 1, NULL, &global, &local,
                                 0, NULL, NULL);
    if (err)
    {
        cerr << "Error: Failed to execute kernel!" << endl;
        return EXIT_FAILURE;
    }
}

for(int i = 0; i < num_devices; i++) {
    //Wait for all commands to complete
    clFinish(queues[i]);

    //Read back the results from the device to verify the output

    err = clEnqueueReadBuffer(queues[i], output,
                               CL_TRUE, 0, sizeof(int) * count,
                               results_apart[i], 0, NULL, NULL );
    if (err != CL_SUCCESS)
    {
        cerr << "Error: Failed to read output array! " <<  err << endl;
        exit(1);
    }

    for(int j = 0; j < partition; j++) {
        results[i*partition + j] = results_apart[i][j];
    }

    delete [] data1_apart[i];
    delete [] data2_apart[i];
    delete [] results_apart[i];
}

clReleaseMemObject(input1);
clReleaseMemObject(input2);
clReleaseMemObject(output);
delete [] data1_apart;
delete [] data2_apart;
}

在将这个问题发布到stackoverflow之前，我已经为这个问题奋斗了2-3周，现在我真的需要某人的帮助，因此我非常感谢任何想法和答案！

3个回答
2
这是我认为正在发生的事情。您为每个参与的OpenCL设备调用一次clEnqueueNDRangeKernel。此时，尚未调用clFlush，因此没有任何内核开始执行。接下来，您为每个队列进行clFinish。第一个clFinish调用导致第一个排队的工作组运行。它还等待其完成。一旦第一个工作组完成，clFinish将控制返回到您的应用程序。然后，您的应用程序为下一个队列调用clFinish。这触发第二个工作组运行，并等待其完成。因此，工作按顺序运行。解决方案可能就是在每次调用clEnqueueNDRangeKernel之后立即调用clFush。这就是我的AMD系统的行为方式。我很快会发布一个可行的示例。
谢谢，我会尝试这个！如果可能的话，请在这种方法上提供更多评论！ - Vladimir
这是一个工作示例[链接]http://notabs.org/blcutil/wip/blcutil_devel-018.7z以下是如何在一个设备上运行，然后在另一个设备上运行，最后同时运行的方法： [链接]http://notabs.org/blcutil/wip/sample-output.htm命令行选项-opencl将使程序列出opencl设备。命令行选项-opencl=允许您选择一个或多个列出的设备进行使用。 - user1940376
是的，我认为这是正确的。OpenCL规范对此非常明确：“请注意，当回调（或其他代码）将命令排队到命令队列时，命令不需要立即开始执行，直到队列被刷新。” - doug65536
1

您的所有设备都使用相同的缓冲区。当内核被执行时，数据将在设备之间移动。如果没有适当的同步，结果将是未定义的。

如果可能，请考虑为每个设备分配一个不同的缓冲区集。

这就是我现在正在做的事情，为重负载（约130毫秒工作）提供几乎即时（约1毫秒差异）的执行内核。 - huseyin tugrul buyukisik
0
你使用哪些GPU？我有一张GTX590，它显示为两个GPU设备。当我尝试在两个设备上运行时，似乎会等待每个设备完成才会移动到下一个设备（即使不应该这样）。我不知道Nvidia是否已经解决了这个问题。
阅读一些消息后，我认为在我阅读Nvidia网站时看到他们建议为每个设备创建单独的上下文，并在不同的线程中运行它们。这就是我所做的，它非常有效。我使用了pthreads（或SDL_threads）来实现。设置起来非常容易。
网页内容由stack overflow 提供, 点击上面的
可以查看英文原文，
原文链接