Gas OpenCL - Field resolution

[danw]

I'm currently building "baby's first OpenCL kernel", to try and speed up setting border voxels in a field to a constant value.  Plenty familiar with vex and such, but first time poking around with OpenCL code.

I've got it working, except that I can't work out how best to pass in the resolution of the DOPs field I'm operating on.  I thought that the stride_x variables might be it... but I've worked out I can only derive x and y resolution from them (dividing stride_z by stride_y), but not z res.

I can calculate it using "Include Size" and "Include Voxel Size" - but it seems odd to be doing floating-point maths to work back to such a basic bit of info, and I don't know how much that'd slow down an OpenCL kernel.

Is there any way to just pass the info straight in?

 

(Edit: should probably chuck up the code I've gotten so far, in case anyone spots any other glaring mistakes while I'm at it :-P)

kernel void windsandwich(
    int stride_x, 
    int stride_y, 
    int stride_z, 
    int stride_offset,
    global float *vel_x,
    global float *vel_y,
    global float *vel_z
)
{
    size_t x = get_global_id(0);
    size_t y = get_global_id(1);
    size_t z = get_global_id(2);
    size_t idx = stride_offset + x * stride_x + y * stride_y + z * stride_z;
    float3 v = 0;

    if (z<4) {
        v = (float3)(0.0f, 0.0f, 0.125f);
    } else if (z>(  /*x_res_attribute*/  -5)) {
        v = (float3)(0.0f, 0.0f, -0.125f);    
    } else {
        v = (float3)(vel_x[idx], vel_y[idx], vel_z[idx]);    
    }
    
    vel_x[idx] = v.x;
    vel_y[idx] = v.y;
    vel_z[idx] = v.z;
}

 

Edited January 17, 2019 by danw

[danw]

Found it.

Needed to use the get_global_size() function.

size_t x_res = get_global_size(0);
size_t y_res = get_global_size(1);
size_t z_res = get_global_size(2);