VexCL vector of structs?

Question

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.

Answer 1

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