c++ - OpenCL 重用具有不同 DEFINE (-D) 的 cl_kernel

Question

我的目标是减少调用 cl_kernel 时的调用开销我首先创建了名为 clfunctor

的仿函数类

该类由以“内核源代码”、“主要函数名称”和编译“选项”作为参数的构造函数组成 然后构造函数将调用 set_kernel_code(...) 以使用这些参数编译和构建代码。最终将得到“mykernel”作为“set_kernel_code”函数的输出

class clfunctor 
{
      // comile sourcecode to cl_kernel
      void set_kernel_code(const std::string& sourcecode, const std::string& 
                           program_name, const std::string& options="")
      { 
              char* kernel_source = new char[ sourcecode.size() + 1 ];
              std::copy( sourcecode.begin(), sourcecode.end(), kernel_source 
              );
              kernel_source[ sourcecode.size() ] = '\0'; // opencl need \0 to end kernel_code

              cl_program program;

              program = clCreateProgramWithSource( myclcontext, 1, (const char **)&kernel_source, NULL, &err);
              if( err != 0 ) echo_error( cl_error_string(err) ); // if error here.. call cl_init() ???

              //build program from opencl device
              clBuildProgram(program, 1, &device[iplatform][idevice], options.c_str(), NULL, NULL);
              if( err != 0 ) echo_error( cl_error_string(err) );

              // no have kernel
              mykernel = clCreateKernel( program, program_name.c_str(), &err );
              const size_t LOG_SIZ = 2040;


              clReleaseProgram( program );
              delete kernel_source;
       }





public:

cl_kernel           mykernel;
std::vector<size_t> local_nd;
std::vector<size_t> global_nd;



clfunctor(const std::string& sourcecode, const std::string& program_name, const std::string& options)
{
    set_kernel_code( sourcecode, program_name, options );
}




void operator()() const
{
    cl_int err;


    //-----------------------------------
    //        create GPU queue
    //----------------------------------
    cl_command_queue queue = clCreateCommandQueue( myclcontext, device[iplatform][idevice], 0, &err);
    if( err != 0 ) echo_error( cl_error_string(err) );



    cl_uint work_dim = local_nd.size(); // ND dimension ( N )
    err = clEnqueueNDRangeKernel( queue, kernel, work_dim, NULL, global_nd.data(), local_nd.data(), 0, NULL, NULL);
    if( err != 0 ) echo_error( cl_error_string(err) );
    clFinish( queue );




    clReleaseCommandQueue( queue );
}


virtual ~clfunctor()
{
    // relase kernel with object is destroyed
    clReleaseKernel( mykernel );
}

};

我有一个示例 cl 文件“nothing.cl”是

kernel void nothing()
{
    printf("[ echo ]: %d\n", MYID  );
}

之后我通过以下方式创建了这个仿函数对象

clfunctor0 hello( file2code("nothing.cl"), "nothing", "-DMYID=100");
hello.local_nd  = {1,1,1};
hello.global_nd = {1,1,1};
hello();   // print 100


clfunctor0 hello2( file2code("nothing.cl"), "nothing", "-DMYID=123");
hello2.local_nd  = {1,1,1};
hello2.global_nd = {1,1,1};
hello2(); // print 123


hello();  // THIS IS WRONG I expect it to print 100 but it print 123

如上代码，你可以看到我创建了两个对象。一个是“hello”，另一个是“hello2”，这两个对象共享来自“nothing.cl”的相同代码，但在 DEFINE (-D) 编译中有所不同。 'hello' 使用 -DMYID=100 ( MYID = 100 ) 但 hello2 使用 -DMYID=123 ( MYID = 123 )

我希望这些对象拥有不同的 mykernel，因为编译选项不同。首先，我调用 hello() 我正确地打印了“100”

然后我调用 hello2() 它也正确打印“123”

但是，当我再次调用 hello() 时，它不是打印“100”它打印来自 hello2 的参数 123

如何仅使用 -D(DEFINE) 编译参数为每个对象制作不同的内核。或者我做错了什么或对 opencl 内核有误解。

Answer 1

恐怕会发生缓存。编译器生成内部包含 123 的二进制文件并简单地重复使用。使用随机名称创建动态 cl 文件，您将始终具有正确的值。或者最好不要使用标志，而只是将参数传递给内核