label.cl

///
/// @file label.cl
///
/// Kernel to map channel confluences and identify their major & minor upstream pixels.
///
/// @author CPS
/// @bug No known bugs
///

#ifdef KERNEL_LABEL_CONFLUENCES

#define CHECK_INFLOWS(here_idx,nbr_vec) { \
   nbr_idx = get_array_idx(nbr_vec); \
   if ( !mask_array[nbr_idx] && (mapping_array[nbr_idx]&IS_THINCHANNEL) ) { \
       if ( link_array[nbr_idx]==here_idx ) { \
           /* The nbr pixel flows into here */ \
           inflows_list[n_inflows++] = nbr_idx; \
       } \
   } \
}
// Check in all 8 pixel-nbr directions
#define CHECK_E_INFLOWS(idx,vec)  CHECK_INFLOWS(idx,(float2)(vec[0]+1.0f, vec[1]      ))
#define CHECK_NE_INFLOWS(idx,vec) CHECK_INFLOWS(idx,(float2)(vec[0]+1.0f, vec[1]+1.0f ))
#define CHECK_N_INFLOWS(idx,vec)  CHECK_INFLOWS(idx,(float2)(vec[0],      vec[1]+1.0f ))
#define CHECK_NW_INFLOWS(idx,vec) CHECK_INFLOWS(idx,(float2)(vec[0]-1.0f, vec[1]+1.0f ))
#define CHECK_W_INFLOWS(idx,vec)  CHECK_INFLOWS(idx,(float2)(vec[0]-1.0f, vec[1]      ))
#define CHECK_SW_INFLOWS(idx,vec) CHECK_INFLOWS(idx,(float2)(vec[0]-1.0f, vec[1]-1.0f ))
#define CHECK_S_INFLOWS(idx,vec)  CHECK_INFLOWS(idx,(float2)(vec[0],      vec[1]-1.0f ))
#define CHECK_SE_INFLOWS(idx,vec) CHECK_INFLOWS(idx,(float2)(vec[0]+1.0f, vec[1]-1.0f ))

///
/// Flag if a pixel IS_MAJORCONFLUENCE and if so flag which upstream pixel IS_MAJORINFLOW
///   or IS_MINORINFLOW.
///
/// Compiled if KERNEL_LABEL_CONFLUENCES is defined.
///
/// @param[in]  seed_point_array       (float2 *, RO):
/// @param[in]  mask_array             (bool *,   RO):
/// @param[in]  uv_array               (float2 *, RO):
/// @param[in]  slt_array              (uint *,   RO):
/// @param[in,out]  mapping_array      (uint *,   RW):
/// @param[in]  count_array            (uint *,   RO):
/// @param[in]  link_array             (uint *,   RO):
///
/// @returns void
///
/// @ingroup structure
///
__kernel void label_confluences(
        __global const float2 *seed_point_array,
        __global const bool   *mask_array,
        __global const float2 *uv_array,
        __global const float  *slt_array,
        __global       uint   *mapping_array,
        __global const uint   *count_array,
        __global const uint   *link_array
   )
{
    // For every (redesignated) thin channel pixel...

    const uint global_id = get_global_id(0u)+get_global_id(1u)*get_global_size(0u);
    if (global_id>=N_SEED_POINTS) {
        // This is a "padding" seed, so let's bail
        return;
    }
#ifdef VERBOSE
    // Report how kernel instances are distributed
    if (global_id==0 || global_id==get_global_offset(0u)) {
        printf("\n  >>> on GPU/OpenCL device: id=%d offset=%d ",
                get_global_id(0u),
                get_global_offset(0u));
        printf("#workitems=%d x #workgroups=%d = %d=%d\n",
                get_local_size(0u), get_num_groups(0u),
                get_local_size(0u)*get_num_groups(0u),
                get_global_size(0u));
    }
#endif
    __private uchar n_inflows=0u, n_equal_dominant_inflows=0u;
    __private uint i, idx, nbr_idx, inflows_list[8], equal_dominant_inflows_list[8],
                   dominant_slt_index=0u;
    __private float dominant_slt=-MAXFLOAT;
    __private float2 vec=seed_point_array[global_id];

    // Remember here
    idx = get_array_idx(vec);
    // Check upstream neighbors
    CHECK_N_INFLOWS(idx,vec);
    CHECK_S_INFLOWS(idx,vec);
    CHECK_E_INFLOWS(idx,vec);
    CHECK_W_INFLOWS(idx,vec);
    CHECK_NE_INFLOWS(idx,vec);
    CHECK_SE_INFLOWS(idx,vec);
    CHECK_NW_INFLOWS(idx,vec);
    CHECK_SW_INFLOWS(idx,vec);
    if (n_inflows>1) {
        atomic_or(&mapping_array[idx],IS_MAJORCONFLUENCE);
        for (i=0;i<n_inflows;i++) {
            if ( count_array[idx]!=(count_array[inflows_list[i]]+1) ) {
                atomic_or(&mapping_array[inflows_list[i]],IS_MINORINFLOW);
            } else {
                equal_dominant_inflows_list[n_equal_dominant_inflows++]= inflows_list[i];
                if (slt_array[inflows_list[i]]>dominant_slt) {
                    dominant_slt_index = inflows_list[i];
                    dominant_slt = slt_array[dominant_slt_index];
                }
            }
        }
#ifdef DEBUG
        if (n_equal_dominant_inflows==0) {
            printf(
         "\nn_equal_dominant_inflows=0 @ %g,%g  idx=%d  mapping=%d & %d= is_thin=%d mask=%d \n",
                    vec[0],vec[1],
                    idx,
                    mapping_array[idx],IS_THINCHANNEL,(mapping_array[idx]&IS_THINCHANNEL),
                    mask_array[idx]);
            for (i=0;i<n_inflows;i++) {
                printf("ifidx=%d   count=%d => %d @ idx=%d\n",inflows_list[i],
                        count_array[inflows_list[i]], count_array[idx],idx);
            }
            for (i=0;i<n_inflows;i++) {
                printf("slt=%g : domslt=%g\n",slt_array[inflows_list[i]]+1, dominant_slt);
            }
        }
#endif
        for (i=0;i<n_equal_dominant_inflows;i++) {
            if (equal_dominant_inflows_list[i]==dominant_slt_index) {
                atomic_or(&mapping_array[equal_dominant_inflows_list[i]],IS_MAJORINFLOW);
            } else {
                atomic_or(&mapping_array[equal_dominant_inflows_list[i]],IS_MINORINFLOW);
            }
        }
    }
    return;
}
#endif