     const int idx     = block_x + threadIdx.x;
     const int idy     = blockIdx.y;
 
-    if( idx >= dx.cols ) return;
-    if( idy >= dx.rows ) return;
+    if( outOfBounds( idx, idy, dx ) ) return;
 
     int32_t  dxVal  = dx.ptr(idy)[idx];
     int32_t  dyVal  = dy.ptr(idy)[idx];
     uint32_t magVal = mag.ptr(idy)[idx];
 
     // -1 if only is negative, 1 else
-    const int32_t signVal = (dxVal ^ dyVal) < 0 ? -1 : 1;
-
-    dxVal = (dxVal < 0 ) ? -dxVal : dxVal;
-    dyVal = (dyVal < 0 ) ? -dyVal : dyVal;
+    // const int32_t signVal = (dxVal ^ dyVal) < 0 ? -1 : 1;
+    const int32_t signVal = d_sign( dxVal ^ dyVal );
+
+    dxVal = d_abs( dxVal );
+    dyVal = d_abs( dyVal );
 
     // 0 - the pixel can not belong to an edge
     // 1 - the pixel might belong to an edge


        dst.ptr(idy)[idx] = d_thinning_lut[log];
#endif
    }
}

#if 0
__global__


    thinning
        <<<grid,block,0,_stream>>>
        ( _d_edges, cv::cuda::PtrStepSzb(_d_intermediate) );

    uint32_t edge_counter = 0;
    POP_CUDA_MEMCPY_ASYNC( _d_edge_counter, &edge_count, sizeof(uint32_t), cudaMemcpyHostToDevice, _stream );

    thinning_and_store
        <<<grid,block,0,_stream>>>
        ( cv::cuda::PtrStepSzb(_d_intermediate), _d_edges, _d_edge_counter, _edge_max, _d_edgelist );

    POP_CUDA_MEMCPY_ASYNC( &edge_counter, _d_edge_counter, sizeof(uint32_t), cudaMemcpyDeviceToHost, _stream );

    // Note: right here, Dynamic Parallelism would avoid blocking.
    cudaStreamSynchronize( _stream );
    // Try to figure out how that can be done with CMake.

    const uint32_t nmax      = params._distSearch;
    const int32_t  threshold = params._thrGradientMagInVote;
    block.x = 2;
    block.y = 32;
    block.z = 0;
    grid.x  = edge_counter / 32 + ( edge_counter & 0x1f != 0 ? 1 : 0 );

    {
        uint32_t dummy = 0;
        POP_CUDA_MEMCPY_ASYNC( _d_edge_counter, &dummy, sizeof(uint32_t), cudaMemcpyHostToDevice, _stream );
    }

    gradiant_descent
        <<<grid,block,0,_stream>>>
        ( _d_edgelist, edge_counter, _d_edgelist_2, _d_edge_counter, _d_edges, nmax, _d_dx, _d_dy, threshold, _d_out_points );

    POP_CUDA_MEMCPY_ASYNC( &edge_counter, _d_edge_counter, sizeof(uint32_t), cudaMemcpyDeviceToHost,_stream );

    // Note: right here, Dynamic Parallelism would avoid blocking.
    cudaStreamSynchronize( _stream );
    // Try to figure out how that can be done with CMake.

    /*
     * After this:
     * - it is unclear whether a sweep over the entire edge plane
     *   is still an efficient choice, or whether it would be better
     *   to reduce the edge points into a flat array first.
     * - if reduction is the thing to do, it would be better to
     *   combine it with the second thinning step.
     *
     * cctagDetectionFromEdges takes as parameters:
     * - markers     : output, list of marker coordinates
     * - points      : a vector that contains all edge points (input)
     * - sourceView  : image at this scale (input)
     * - cannyGradX  : _d_dx (input)
     * - cannyGradY  : _d_dy (input)
     * - edgesMap    : _d_edges (input)
     * - frame       : running counter for video frames
     * - level       : this layer of the pyramid
     * - scale       : probably a coordinate multiplier
     * - params      : parameters
     *
     * calls vote
     *
     * vote takes as input:
     * - points     : a vector that contains all edge points (input)
     * - seeds      : unknown (output)
     * - edgesMap   : _d_edges (input)
     * - winners    : map of winners (output)
     * - cannyGradX : _d_dx (input)
     * - cannyGradY : _d_dy (input)
     * - params     : parameters
     *
     * calls gradientDirectionDescent
     *
     * gradientDirectionDescent takes as input
     * - canny       : _d_edges (input)
     * - p           : coordinate of one edge point (input)
     * - dir         : direction, 1 or -1
     * - nmax        : global parameter the distance from edge point for searching
     * - cannyGradX  : _d_dx
     * - cannyGradY  : _d_dy
     * - thrGradiant : global parameter for gradient thresholding
     * returns: NULL or a new point
     */
#if 0
    // very costly printf-debugging
    debug_gauss
