PDA

View Full Version : How to port existing C++ function using vectors?



Narada
09-14-2012, 04:10 AM
Hi all,

I'm new to C, C++ and OpenCL. I'm learning them as best as I can at the moment.

First let me present you with an example C++ function.


#include <vector>

using namespace std;

class Test {

private:

double a;
vector<double> b;
vector<long> c;
vector<vector<double> > d;

public:

double foo(long x, double y) {
// uses x, y, a, b, c, d
// uses b.size()
}

}

This is an example function from an existing codebase. I'm trying to port it to OpenCL using either the C or C++ binding. The initial objective is to do a literal port. On this note I had a few questions.

The function takes two input arguments and returns a value. However in order to calculate its return value it accesses innumerable class members as well as calling functions on its members like vector.size().

My primary broad question is: how do I provide this member data to the kernel so that it is available when I port the body of the function to OpenCL?

Specifically here are my underlying questions.

1. How do I make an existing vector<T> and a vector<vector<t>> available to a kernel? Is there any way to just pass pointers to them into the kernel and just read them using indices in OpenCL? What's the best way to make sure the memory contents of these members in the host are kept in sync with the OpenCL runtime?
2. The C++ vectors are dynamic data structures. What do I specify for their size when creating buffers to pass in as kernel arguments? Do I just do sizeof(c)? What if the vector changes part way through?
3. How would I pass in the nested vector<vector<double>> type?
4. What would the kernel argument types in OpenCL itself for such dynamic C++ types as vectors?
5. How do I access functions like b.size() within the OpenCL kernel itself as the C++ function accesses this function? Do I have to pass this as an additional argument to the kernel? What if it changes in the meantime?

I should add that rewriting the function is not an option as it is an existing codebase and I have some constraints.

I appreciate these questions may be a bit basic but as I said I'm learning at the moment. When answering I would really appreciate some examples using either the C or C++ binding and OpenCL.

Many thanks.

chippies
09-14-2012, 07:03 AM
1. How do I make an existing vector<T> and a vector<vector<t>> available to a kernel? Is there any way to just pass pointers to them into the kernel and just read them using indices in OpenCL? What's the best way to make sure the memory contents of these members in the host are kept in sync with the OpenCL runtime?

You can't just pass pointers to the vectors unfortunately. You need to convert the vectors into arrays, i.e. T* for vector<T>. Handling vector<vector<T> > is more difficult since each vector<T> element can have a different length. A possible solution is to assemble the whole vector<vector<T> > into an array (T*) but also pass in an array that gives the starting location of each vector<T> element and the number of T's in each vector. Indexing into the array is then a two step procedure.


2. The C++ vectors are dynamic data structures. What do I specify for their size when creating buffers to pass in as kernel arguments? Do I just do sizeof(c)? What if the vector changes part way through?

You would have to work out the number of elements in your vector and multiply by the size of an element, e.g. c.size() * sizeof(T). If an element in one of the vectors changes while a kernel is executing, the kernel will not receive the new value until your code enqueues a transfer and that transfer is completed. This might or might not happen while the kernel is executing, depending on the implementation and the options used when setting up the OpenCL context.


3. How would I pass in the nested vector<vector<double>> type?
4. What would the kernel argument types in OpenCL itself for such dynamic C++ types as vectors?

I think I've answered these above.


5. How do I access functions like b.size() within the OpenCL kernel itself as the C++ function accesses this function? Do I have to pass this as an additional argument to the kernel? What if it changes in the meantime?

None of the functions supplied by the vector class will be directly available. You will have to pass the size of the vector as one of the parameters of your kernel. If it changes while the kernel is running, too bad, the kernel will not find out.


I should add that rewriting the function is not an option as it is an existing codebase and I have some constraints.

I'd suggest that you rewrite the functions that this function calls so that they transform the vectors into arrays, send them off to the GPU and transform the results back into vectors. That's definitely not efficient since you will be executing a lot of unnecessary copies between the GPU and CPU and wasting time converting from vectors to arrays, but if you cannot rewrite this function then you will have to make some compromises.

Narada
09-14-2012, 07:23 AM
Thank you very much chippies. Regarding (4) what would the opencl data type be for the arrays passed in? I'm familiar with the scalar and vector data types but in this case the array sizes are unknown to me. For example for long arrays passed in how would I correct the following?


__kernel void test(__global long *b, __global long *c) {
// access using b[0] and c[1] for example?
}

chippies
09-16-2012, 11:17 AM
Hi Narada



__kernel void test(__global double *b, long bLen, __global long *c, long cLen) {
b[0] = 1.0; //Assign this value to the first element in b.
b[bLen-1] = 1.0; //Assign a value to the last element of b.
c[cLen-1] = cLen; //Assign the length of c to the last element of c, for some reason :) .
}


In your code b is defined as vector<double> b, hence you would convert that to an array double* bHost, followed by copying that to a device pointer cl_mem bDev using clEnqueuWriteBuffer. Your kernel parameter is them bDev. You have to do the same with c. Note that I also pass in bLen and cLen. These are constants containing the length - number of elements - of b and c, since your kernel has no way of knowing what those values are unless you explicitly pass them as parameters.