     __constant__ int ctwidth;
     __constant__ int ctheight;
 
-    static const unsigned int czero = 0;
-
     // Global counter of blocks finished its work
-    __device__ unsigned int blocks_finished;
+    __device__ unsigned int blocks_finished = 0;
 
 
     // Estimates good thread configuration


        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);
        cudaSafeCall(cudaThreadSynchronize());



        unsigned int* count_buf = (unsigned int*)buf.ptr(0);

        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());
