How to call boost_compute 'BOOST_COMPUTE_FUNCTION' defined function?

1k views Asked by At

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

1

There are 1 answers

1
kenba On BEST ANSWER

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