     int block = pcl::device::appnearest_search::KernelPolicy::CTA_SIZE;
     int grid = (batch.queries_num + block - 1) / block;    
 
-    cudaSafeCall( cudaFuncSetCacheConfig(pcl::device::appnearest_search::KernelAN<PointType>, cudaFuncCachePreferL1) );
+    cudaSafeCall( cudaFuncSetCacheConfig(pcl::device::appnearest_search::KernelAN, cudaFuncCachePreferL1) );
 
     pcl::device::appnearest_search::KernelAN<<<grid, block>>>(batch);
     cudaSafeCall( cudaGetLastError() );

			LOG_WARP_SIZE = 5,
			WARP_SIZE = 1 << LOG_WARP_SIZE,
			WARPS_COUNT = CTA_SIZE/WARP_SIZE,                    
		};

		struct SmemStorage
		{         
			volatile int   per_warp_buffer[WARPS_COUNT];
			volatile int   cta_buffer_int[CTA_SIZE];
			volatile float cta_buffer_flt[CTA_SIZE];
		};
	};

	__shared__ KernelPolicy::SmemStorage storage;


	template<typename PointType>
	struct Warp_appNearestSearch
	{   
	public:                
		typedef Batch<PointType> BatchType;

		const BatchType& batch;

		int query_index;        
		float3 query;  
		int result_idx;

		__device__ __forceinline__ Warp_appNearestSearch(const BatchType& batch_arg, int query_index_arg) 
			: batch(batch_arg), query_index(query_index_arg){}

		__device__ __forceinline__ void launch(bool active)

				int active_lane = __ffs(mask) - 1; //[0..31]                        
				mask &= ~(1 << active_lane);   

				volatile int* warp_buffer = &per_warp_buffer[warpId];

				//broadcast beg
				if (active_lane == laneId)
