     // const float sin_t = sinf(ang);
     float cos_t;
     float sin_t;
-    sincosf( ang, &sin_t, &cos_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;
    const float ptx = csbp * offsetptx - ssbp * offsetpty + x;
    const float pty = csbp * offsetpty + ssbp * offsetptx + y;

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

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

    const int32_t wx = xmax - xmin + 1;
    const int32_t hy = ymax - ymin + 1;
    const int32_t 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+=32)
    {
        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 = crsbp * dx + srsbp * dy;
        const float ny = crsbp * dy - srsbp * dx;
        const float nxn = fabs(nx);
        const float nyn = fabs(ny);
        if (nxn < 1.0f && nyn < 1.0f) {
            const float2 mod_th = get_gradiant( jj, ii, layer );
            const float& mod    = mod_th.x;

                           Descriptor*   descs,
                           Plane2D_float layer )
{
    const int width  = layer.getWidth();
    const int height = layer.getHeight();

    // int bidx = blockIdx.x & 0xf; // lower 4 bits of block ID
    const int ix   = threadIdx.y; // bidx & 0x3;       // lower 2 bits of block ID
