VexCL: set values common to all vector elements

Question

I work with ray tracing and use GPU to calculate pixel colours. I was using NVIDIA CUDA and now want to go to VexCL. I'm trying to use such code:

struct Ray;
vex::Context ctx(...);

...

unsigned int frame_width, frame_height;
std::array<float, 4> camera_direction, camera_up;
float camera_fov;

...

// initialize values and store them in GPU memory too

...

vex::vector<Ray> rays(ctx, frame_width * frame_height);

and something like

rays = some_expression_to_calculate_ray(vex::element_index(), frame_width,
                                        camera_direction, camera_up, camera_fov);

So my question is: how can I explain to VexCL that some values must be common for all vector elements?

I was trying VEX_CONSTANT, vex::raw_pointer but it is not what I need.

Can you clarify a bit what exactly you want to do? You can pass a scalar value to any vexcl expression, including function call, and any vector element in the expression will use that value. Why this does not work for you? — ddemidov, Aug 20 '15 at 18:12
Also, this could be related: http://stackoverflow.com/questions/30842071/vexcl-vector-of-structs — ddemidov, Aug 20 '15 at 18:22
Thanks @ddemidov, but it is not what I actually want. When using CUDA I was allocating camera_direction, camera_up, etc with keyword `__constant__` or `__device__` and then was using them in kernels. These values are time to time changing by main program code and updating using `cudaMemcpy()`, but when calculating a frame they are constant and common for each pixel. How can I do the same using VexCL? I don't want to make a copy of such values for each pixel. — zedrian, Aug 20 '15 at 19:41

score 1 · Accepted Answer · edited May 23 '17 at 12:29

If you change the type of camera_direction and camera_up from std::array<float,4> to cl_float4, then you would be able to directly use those in an expression:

#include <vexcl/vexcl.hpp>

int main() {
    vex::Context ctx(vex::Filter::Env);

    VEX_FUNCTION(float, dummy, (size_t, idx)(cl_float4, dir)(cl_float4, up)(float, fov),
            // whatever
            return idx + length(dir - up) + fov;
            );

    cl_float4 camera_dir = {1, 2, 3, 4}, camera_up = {1, 0, 0, 0};
    float camera_fov = 42;

    vex::vector<float> rays(ctx, 1024);

    rays = dummy(vex::element_index(), camera_dir, camera_up, camera_fov);
}

(I've changed rays to be a vector of floats for simplicity, see the linked question for how to work with structs in VexCL.) camera_dir, camera_up, and camera_fov are defined host-side, and they are passed to the kernel as parameters. So no unnecessary copies are being made. Here is the generated OpenCL kernel:

float dummy(ulong idx, float4 dir, float4 up, float fov) {
  return idx + length(dir - up) + fov;
}
kernel void vexcl_vector_kernel(
  ulong n,
  global float * prm_1,
  ulong prm_2,
  float4 prm_3,
  float4 prm_4,
  float prm_5
)
{
  for(ulong idx = get_global_id(0); idx < n; idx += get_global_size(0))
  {
    prm_1[idx] = dummy( (prm_2 + idx), prm_3, prm_4, prm_5 );
  }
}

VexCL: set values common to all vector elements

1 Answers1