c++structopenclvexcl

VexCL vector of structs?


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.


Solution

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