Function get_features

src/backend/cuda/kernel/fast.hpp:277–320

template<typename T>
__global__ void get_features(float *x_out, float *y_out, float *score_out,
                             const T *flags, const unsigned *d_counts,
                             const unsigned *d_offsets, const unsigned total,
                             const unsigned idim0, const unsigned idim1,
                             const unsigned edge) {
    const int xid = blockIdx.x * blockDim.x * 2 + threadIdx.x;
    const int yid = blockIdx.y * blockDim.y * 8 + threadIdx.y;
    const int tid = blockDim.x * threadIdx.y + threadIdx.x;

    const int xoff = blockDim.x;
    const int yoff = blockDim.y;

    const int xend = (blockIdx.x + 1) * blockDim.x * 2;
    const int yend = (blockIdx.y + 1) * blockDim.y * 8;

    const int bid = blockIdx.y * gridDim.x + blockIdx.x;

    __shared__ unsigned s_count;
    __shared__ unsigned s_idx;

    if (tid == 0) {
        s_count = d_counts[bid];
        s_idx   = d_offsets[bid];
    }
    __syncthreads();

    // Blocks that are empty, please bail
    if (s_count == 0) return;
    for (int y = yid; y < yend; y += yoff) {
        if (y >= idim1 - edge - 1 || y <= edge + 1) continue;
        for (int x = xid; x < xend; x += xoff) {
            if (x >= idim0 - edge - 1 || x <= edge + 1) continue;

            float v = flags[y * idim0 + x];
            if (v == 0) continue;

            unsigned id = atomicAdd(&s_idx, 1u);
            if (id >= total) return;
            y_out[id]     = x;
            x_out[id]     = y;
            score_out[id] = v;
        }
    }
}

template<typename T>
void fast(unsigned *out_feat, float **x_out, float **y_out, float **score_out,