So I know that it is possible to use custom types with OpenCL. But I haven't been able to use them with VexCL. Creating a device vector of structs works fine, but I can't perform any operations.
As I haven't found any examples using custom types with VexCL my question is is that even possible? Thanks in advance.
VexCL does not support operations with vectors of structs out of the box. You will need to help it a bit. First, you need to tell VexCL how to spell the type name of the struct. Let's say you have the following struct defined on the host side:
struct point2d {
double x;
double y;
};
You need to provide a specification of the vex::type_name_impl struct that will generate a string corresponding to the type name of the struct. Remember that the code you are generating is C99:
namespace vex {
template <> struct type_name_impl<point2d> {
static std::string get() { return "struct point2d"; }
};
}
You will also need to make sure every generated kernel knows about your struct. This may be achieved with vex::push_program_header() function after the VexCL context has been initialized:
vex::push_program_header(ctx, "struct point2d { double x; double y; };");
This will allow you to declare vectors of the struct, and to pass the vectors to custom functions. That should be general enough. Here is the complete example:
#include <vexcl/vexcl.hpp>
// Host-side definition of the struct.
struct point2d {
double x, y;
};
// We need this for code generation.
namespace vex {
template <>
struct type_name_impl<point2d> {
static std::string get() { return "struct point2d"; }
};
}
int main() {
const size_t n = 16;
vex::Context ctx(vex::Filter::Env);
std::cout << ctx << std::endl;
// After this, every kernel will have the struct declaration in header:
vex::push_program_header(ctx, "struct point2d { double x; double y; };");
// Now we may define vectors of the struct:
vex::vector<point2d> x(ctx, n);
vex::vector<double> y(ctx, n);
// We won't be able to use the vectors in any expressions except for
// custom functions, but that should be enough:
VEX_FUNCTION(point2d, init, (double, x)(double, y),
struct point2d p = {x, y}; return p;
);
VEX_FUNCTION(double, dist, (point2d, p),
return sqrt(p.x * p.x + p.y * p.y);
);
x = init(3,4);
y = dist(x);
std::cout << y << std::endl;
}
And here is the kernel that will be generated for the assignment operation of y = dist(x);:
struct point2d { double x; double y; };
double dist
(
struct point2d p
)
{
return sqrt(p.x * p.x + p.y * p.y);
}
kernel void vexcl_vector_kernel
(
ulong n,
global double * prm_1,
global struct point2d * prm_2
)
{
for(ulong idx = get_global_id(0); idx < n; idx += get_global_size(0))
{
prm_1[idx] = dist( prm_2[idx] );
}
}
