Search code examples
c++boostopenclboost-compute

How to call boost_compute 'BOOST_COMPUTE_FUNCTION' defined function?


I'm currently exploring boost_compute. Unfortunately there are less documentation pages and examples, than I need to understand what to do.

Given the following minified code:

BOOST_COMPUTE_FUNCTION(bool, add, (int* values, int* results, int constant),
{
    // Whats the indexing variable?
    // In opencl it would be get_global_id(0)
    int index = // ?

    results[index] = values[index] + values[index + 1] + values[index + 2] + constant;
});

void compute(float* results, compute::context* ctx, compute::command_queue* queue)
{
    compute::vector<float> device_values(100, *ctx);
    compute::vector<float> device_results(98, *ctx);

    compute::copy(
        parameters->values.begin(), parameters->values.end(), device_values.begin(), *queue
    );

    // Actual computation
    // HOW TO CALL 'add' for every device_results element?

    compute::copy(
        device_results.begin(), device_results.end(), results, *queue
    );
}

How to call the 'add' function and what's the iterating variable inside of this function? Furthermore I need this structure of code to make more complex calculation.

Kind Regards, Toni


Solution

  • In short boost:compute functions are not OpenCL kernel functions. They are more like OpenGL kernel functions.

    I believe that your function takes too many parameters to be used with the boost:compute algorithms.
    However, a slightly simpler function, just adding adjacent values without the constant, would be:

    BOOST_COMPUTE_FUNCTION(boost::compute::float_, add,
                            (boost::compute::float_ values0, boost::compute::float_ values1),
    {
      return values0 + values1;
    });
    

    And could be called using boost::compute::transform as @ddemidov suggested:

    boost::compute::transform(values.begin(), values.end() -1, // values0
                              values.begin() +1, // values1
                              results.begin(), // results
                              add, queue);
    

    It may be possible to implement your function using boost::compute::lambda functions. e.g.:

    using namespace boost::compute::lambda;
    
    float c = 1.234; // some constant
    
    boost::compute::transform(values.begin(), values.end() -1, // values0
                              values.begin() +1, // values1
                              results.begin(), // results
                              _1 + _2 + c, queue);
    

    But it's still short of a set of values...

    Your function could be written as an OpenCL kernel in boost:compute using the BOOST_COMPUTE_STRINGIZE_SOURCE macro:

    const char kernel_function_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
    
      kernel void add(global float* values, global float* results, global float* constant)
      {
        size_t index = get_global_id(0);
        results[index] = values[index] + values[index + 1] + values[index + 2] + *constant;
      }
    
    );
    

    After you've built your kernel program and created your kernel (using boost::compute::program), you can set the kernel arguments individually and call the boost::compute::command_queue enqueue_1d_range_kernel function:

    kernel.set_arg(0, values.get_buffer());
    kernel.set_arg(1, results.get_buffer());
    kernel.set_arg(2, &constant);
    queue.enqueue_1d_range_kernel(kernel, 0, count, 0);