 
         unsigned int* count_buf = (unsigned int*)buf.ptr(0);
 
-        cudaSafeCall(cudaMemcpyToSymbol(blocks_finished, &czero, sizeof(blocks_finished)));
         count_non_zero_kernel<256, T><<<grid, threads>>>(src, count_buf);
         count_non_zero_kernel_2ndstep<256, T><<<1, 256>>>(count_buf, grid.x * grid.y);
         cudaSafeCall(cudaThreadSynchronize());

        unsigned int* minloc_buf = (unsigned int*)locbuf.ptr(0);
        unsigned int* maxloc_buf = (unsigned int*)locbuf.ptr(1);

        cudaSafeCall(cudaMemcpyToSymbol(blocks_finished, &czero, sizeof(blocks_finished)));
        min_max_loc_kernel<256, T><<<grid, threads>>>(src, minval_buf, maxval_buf, minloc_buf, maxloc_buf);
        min_max_loc_kernel_2ndstep<256, T><<<1, 256>>>(minval_buf, maxval_buf, minloc_buf, maxloc_buf, grid.x * grid.y);
        cudaSafeCall(cudaThreadSynchronize());

    __constant__ int ctwidth;
    __constant__ int ctheight;

    // Global counter of blocks finished its work
    __device__ unsigned int blocks_finished = 0;


    // Estimates good thread configuration
