     /* When this kernel finishes, the best point does not
      * exist or it is stored in point_buffer[0]
      */
+    POP_CHK_CALL_IFSYNC;
     POP_CUDA_MEMCPY_TO_HOST_ASYNC( cctag_pointer_buffer,
                                    point_buffer,
                                    sizeof(popart::NearbyPoint),
                                    _stream );
+    POP_CHK_CALL_IFSYNC;
 
     cudaStreamSynchronize( _stream );
     if( not cctag_pointer_buffer->readable ) {

          cut_buffer,
          signals,
          vCutsSize );

    int numPairs = vCutsSize*(vCutsSize-1)/2;
    block.x = 32; // we use this to sum up signals
    block.y = 32; // we can use some shared memory/warp magic for summing
    block.z = 1;
    grid.x  = grid_divide( numPairs, 32 );
    grid.y  = 1;
    grid.z  = 1;

    // _meta.toDevice( Identification_result, 0.0f, _stream );
    // _meta.toDevice( Identification_resct,  0,    _stream );

    popart::identification::idComputeResult
        <<<grid,block>>>

    }
    mT.condition( center );

    const int i = blockIdx.x * 32 + threadIdx.x;
    const int j = blockIdx.y * 32 + threadIdx.y;
    if( i >= gridNSample ) return;
    if( j >= gridNSample ) return;


    NearbyPoint* nPoint;
    identification::CutSignals* signals;

    {
        const int idx = j * gridNSample + i;
        nPoint  = &point_buffer[idx];
        signals = &sig_buffer[idx * vCutsSize];
    }

    nPoint->point = make_float2( center.x - halfWidth + i*stepSize,
                                center.y - halfWidth + j*stepSize );

    popart::geometry::matrix3x3  mInvT;
    mT.invert( mInvT ); // note: returns false if it fails
    mInvT.condition( nPoint->point );

    nPoint->result   = 0.0f;
    nPoint->resSize  = 0;
