如何用内核编译opencl项目

Question

我完全是opencl的初学者,我在互联网上搜索并为opencl项目找到了一些"helloworld"演示.通常在这种最小的项目中,有一个*.cl文件包含某种opencl内核,而*.c文件包含main函数.那么问题是如何使用命令行编译这种项目.我知道我应该在linux上使用某种-lOpenCL标志,在mac上使用-framework OpenCL.但我不知道将*.cl内核链接到我的主源文件.感谢您提供任何意见或有用的链接.

Answer 1

在OpenCL中,.cl通常在运行时编译和构建包含设备内核代码的文件.它意味着您的主机OpenCL程序中的某个位置,您必须编译和构建您的设备程序才能使用它.此功能可实现最大的可移植性.

让我们考虑一下我从两本书中收集的一个例子.下面是一个非常简单的OpenCL内核,它从两个全局数组中添加两个数字并将它们保存在另一个全局数组中.我将此代码保存在一个名为的文件中vector_add_kernel.cl.

kernel void vecadd( global int* A, global int* B, global int* C ) {
    const int idx = get_global_id(0);
    C[idx] = A[idx] + B[idx];
}

下面是用C++编写的利用OpenCL C++ API的主机代码.我将它保存在一个名为ocl_vector_addition.cpp保存.cl文件的文件旁边的文件中.

#include <iostream>
#include <fstream>
#include <string>
#include <memory>
#include <stdlib.h>

#define __CL_ENABLE_EXCEPTIONS
#if defined(__APPLE__) || defined(__MACOSX)
#include <OpenCL/cl.cpp>
#else
#include <CL/cl.hpp>
#endif

int main( int argc, char** argv ) {

    const int N_ELEMENTS=1024*1024;
    unsigned int platform_id=0, device_id=0;

    try{
        std::unique_ptr<int[]> A(new int[N_ELEMENTS]); // Or you can use simple dynamic arrays like: int* A = new int[N_ELEMENTS];
        std::unique_ptr<int[]> B(new int[N_ELEMENTS]);
        std::unique_ptr<int[]> C(new int[N_ELEMENTS]);

        for( int i = 0; i < N_ELEMENTS; ++i ) {
            A[i] = i;
            B[i] = i;
        }

        // Query for platforms
        std::vector<cl::Platform> platforms;
        cl::Platform::get(&platforms);

        // Get a list of devices on this platform
        std::vector<cl::Device> devices;
        platforms[platform_id].getDevices(CL_DEVICE_TYPE_GPU|CL_DEVICE_TYPE_CPU, &devices); // Select the platform.

        // Create a context
        cl::Context context(devices);

        // Create a command queue
        cl::CommandQueue queue = cl::CommandQueue( context, devices[device_id] );   // Select the device.

        // Create the memory buffers
        cl::Buffer bufferA=cl::Buffer(context, CL_MEM_READ_ONLY, N_ELEMENTS * sizeof(int));
        cl::Buffer bufferB=cl::Buffer(context, CL_MEM_READ_ONLY, N_ELEMENTS * sizeof(int));
        cl::Buffer bufferC=cl::Buffer(context, CL_MEM_WRITE_ONLY, N_ELEMENTS * sizeof(int));

        // Copy the input data to the input buffers using the command queue.
        queue.enqueueWriteBuffer( bufferA, CL_FALSE, 0, N_ELEMENTS * sizeof(int), A.get() );
        queue.enqueueWriteBuffer( bufferB, CL_FALSE, 0, N_ELEMENTS * sizeof(int), B.get() );

        // Read the program source
        std::ifstream sourceFile("vector_add_kernel.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()));

        // Make program from the source code
        cl::Program program=cl::Program(context, source);

        // Build the program for the devices
        program.build(devices);

        // Make kernel
        cl::Kernel vecadd_kernel(program, "vecadd");

        // Set the kernel arguments
        vecadd_kernel.setArg( 0, bufferA );
        vecadd_kernel.setArg( 1, bufferB );
        vecadd_kernel.setArg( 2, bufferC );

        // Execute the kernel
        cl::NDRange global( N_ELEMENTS );
        cl::NDRange local( 256 );
        queue.enqueueNDRangeKernel( vecadd_kernel, cl::NullRange, global, local );

        // Copy the output data back to the host
        queue.enqueueReadBuffer( bufferC, CL_TRUE, 0, N_ELEMENTS * sizeof(int), C.get() );

        // Verify the result
        bool result=true;
        for (int i=0; i<N_ELEMENTS; i ++)
            if (C[i] !=A[i]+B[i]) {
                result=false;
                break;
            }
        if (result)
            std::cout<< "Success!\n";
        else
            std::cout<< "Failed!\n";

    }
    catch(cl::Error err) {
        std::cout << "Error: " << err.what() << "(" << err.err() << ")" << std::endl;
        return( EXIT_FAILURE );
    }

    std::cout << "Done.\n";
    return( EXIT_SUCCESS );
}

我在Ubuntu 12.04的机器上编译这段代码,如下所示:

g++ ocl_vector_addition.cpp -lOpenCL -std=c++11 -o ocl_vector_addition.o

它产生一个ocl_vector_addition.o,当我运行时,显示成功的输出.如果你看一下编译命令,你会发现我们没有传递任何关于我们.cl文件的信息.我们只使用了-lOpenCLflag来为我们的程序启用OpenCL库.此外,不要被-std=c++11命令分心.因为我std::unique_ptr在主机代码中使用过,所以我不得不使用这个标志来成功编译.

那么这个.cl文件在哪里使用？如果你看一下主机代码,你会发现我在下面重复编写的四个部分:

// 1. Read the program source
std::ifstream sourceFile("vector_add_kernel.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()));

// 2. Make program from the source code
cl::Program program=cl::Program(context, source);

// 3. Build the program for the devices
program.build(devices);

// 4. Make kernel
cl::Kernel vecadd_kernel(program, "vecadd");

在第1步中,我们读取保存设备代码的文件的内容并将其放入std::string命名sourceCode.然后我们创建一对字符串及其长度并将其保存到source具有该类型的字符串中cl::Program::Sources.在我们准备好代码之后,我们创建了一个cl::program以program该命名的对象,context并将源代码加载到程序对象中.第三步是OpenCL代码编译(和链接)的步骤device.由于设备代码是在第3步中构建的,因此我们可以创建一个名为的内核对象vecadd_kernel,并将其vecadd内部命名的内核与我们的cl::kernel对象相关联.这几乎是编译.cl程序中文件所涉及的一组步骤.

我展示和解释的程序从内核源代码创建了设备程序.另一种选择是使用二进制文件.使用二进制程序增强了应用程序加载时间并允许程序的二进制分发,但限制了可移植性,因为在一个设备上正常工作的二进制文件可能无法在另一个设 使用源代码和二进制文件创建程序也分别称为脱机和在线编译(此处有更多信息).我在这里跳过它,因为答案已经太久了.

Answer 2

我的回答晚了四年。不过，我要补充一些补充@Farzad 的答案，如下所示。

令人困惑的是，在 OpenCL 实践中，编译这个动词被用来表示两种不同的、不兼容的东西：

• 在一种用法中，编译意味着您已经认为它意味着什么。这意味着在构建时构建，如从 *.c 源生成用于构建时链接的 *.o 对象。
• 然而，在另一种用法中——这个其他用法你可能不熟悉——编译意味着在运行时解释，如来自 *.cl 源，产生 GPU 机器代码。

一种发生在构建时。另一个发生在运行时。

引入两个不同的动词可能不会那么令人困惑，但这不是术语的演变方式。按照惯例，动词to compile用于两者。

如果不确定，那么试试这个实验：重命名你的 *.cl 文件，这样你的其他源文件就找不到它，然后构建。

看？它构建得很好，不是吗？

这是因为在构建时不咨询 *.cl 文件。只有稍后，当您尝试执行二进制可执行文件时，程序才会失败。

如果有帮助，您可以将 *.cl 文件视为数据文件或配置文件，甚至是脚本。它可能不是字面上的数据文件、配置文件或脚本，因为它最终会被编译成某种机器代码，但机器代码是 GPU 代码，它不是由 *.cl 程序文本生成的直到运行时。此外，在运行时，不涉及您的 C 编译器。相反，它是您的 OpenCL 库进行构建。

我花了相当长的时间在脑海中理顺这些概念，主要是因为——和你一样——我早就熟悉 C/C++ 构建周期的各个阶段；因此，我以为我知道像编译这样的词是什么意思。一旦您理清了单词和概念，各种 OpenCL 文档就开始变得有意义，您就可以开始工作了。