OpenCL built-in function select

Question

I am trying to use the select function to choose elements from v1 and v2 based on my 3rd argument, however i do not know a way to access the current components in v1 and v2.

Say if v[i] is more than 5, i want to select v2[i] into results[i] else select v1[i], but i can't access the components that way.

Any advice would be appreciated! I am a super beginner at this btw

__kernel void copy(__global int4* Array1, 
               __global int* Array2,
               __global int* output
                ) 
{

    int id = get_local_id(0);

    //Reads the contents from array 1 and 2 into local memory
    __local int4 local_array1;
    __local int local_array2;

    local_array1 = Array1[id];
    local_array2 = Array2[id];


    //Copy the contents of array 1 into an int8 vector called v
    int8 v;

    /*i have trouble here too, how do i copy into int8 v from int4 data type */ 
    v = vload8(0, Array1);


    //Copy the contents of array 2 into two int8 vectors called v1 and v2
    int8 v1, v2;
    v1 = vload8(0, Array2);
    v2 = vload8(1, Array2);

    //Creates an int8 vector in private memory called results

    int8 results;
    if (any(v > 5) == 1) {

        results = select(v2[what do i do to get current index], v1[i], isgreater(v[i], 5.0));*

        vstore8(results, 0, output);

    }
    else {
        results.lo = v1.lo;
        results.hi = v2.lo;
        vstore8 (results, 0, output);

    }

}

Answer 1

You're trying to access vector elements via the [] operator. That's illegal in OpenCL. It might work with some OpenCL compilers but it's still undefined behaviour.

The "official" way to access vector elements is by 1) vector.X or 2) vector.sX As you noticed, this does not allow dynamic access.

The reason is: a vector is not an array. A vector is supposed to map to a hardware "vector register" (or multiple registers). E.g. a "float8" will map to a single 256bit AVX register on AVX2 CPU, or two 128bit AVX registers on AVX1 CPU.

OpenCL does not have an operator to dynamically access vector elements. Perhaps a missing feature, but it reflects the reality of vectorized hardware: most instructions only operate on entire hardware vector registers, not on their individual elements. If you want to work on a vector element selected dynamically, you have to extract it from the vector. Here are several ways to do it.

Using vectors makes sense in some specific cases; IMO they're useful mainly for two things: 1) when you have a bunch of values logically tied together (e.g. colors in a pixel) and 99.99% of time you don't need to access individual values; 2) you have hardware with vector registers (e.g. a VLIW CPU or GPU) and your OpenCL compiler can't "autovectorize" the code, so you need to manually write vectorized code to get reasonable performance.

In your code, i'd simply change __global int4* Array1 to __global int* Array1, write the kernel without using vectors (simply indexing as a normal array), and see how it performs. If you're targeting modern Nvidia/AMD GPUs, you don't need vectors at all to get good performance.