 #include "assist.h"
 
 #undef DESCRIPTORS_FROM_UNBLURRED_IMAGE
+
+// override global setting
+// #undef USE_DYNAMIC_PARALLELISM
 
 /*************************************************************
  * V1: device side

    // const float sin_t = sinf(ang);
    float cos_t;
    float sin_t;
    sincosf( ang, &sin_t, &cos_t );

    const float csbp  = cos_t * SBP;
    const float ssbp  = sin_t * SBP;


    const float offsetptx = ix - 1.5f;
    const float offsetpty = iy - 1.5f;

    // The following 2 lines were the primary bottleneck of this kernel
    // const float ptx = csbp * offsetptx - ssbp * offsetpty + x;
    // const float pty = csbp * offsetpty + ssbp * offsetptx + y;
    const float ptx = fmaf( csbp, offsetptx, - fmaf( ssbp, offsetpty, x ) );
    const float pty = fmaf( csbp, offsetpty,   fmaf( ssbp, offsetptx, y ) );

    const float bsz = fabsf(csbp) + fabsf(ssbp);

    const int xmin = max(1,          (int)floorf(ptx - bsz));
    const int ymin = max(1,          (int)floorf(pty - bsz));
    const int xmax = min(width - 2,  (int)floorf(ptx + bsz));
    const int ymax = min(height - 2, (int)floorf(pty + bsz));

    const int wx = xmax - xmin + 1;
    const int hy = ymax - ymin + 1;
    const int loops = wx * hy;

    float dpt[9] = { 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f };
    // for (int i = 0; i < 9; i++) dpt[i] = 0.0f;

    for( int i = threadIdx.x; i < loops; i+=blockDim.x )
    {
        const int ii = i / wx + ymin;
        const int jj = i % wx + xmin;     

        const float dx = jj - ptx;
        const float dy = ii - pty;
        const float nx = fmaf( crsbp, dx,  srsbp * dy ); // crsbp * dx + srsbp * dy;
        const float ny = fmaf( crsbp, dy, -srsbp * dx ); // crsbp * dy - srsbp * dx;
        const float nxn = fabsf(nx);
        const float nyn = fabsf(ny);
        if (nxn < 1.0f && nyn < 1.0f) {
            const float2 mod_th = get_gradiant( jj, ii, layer );
            const float& mod    = mod_th.x;
