 			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;
+	public:                		
+		const Batch& batch;
 
 		int query_index;        
 		float3 query;  
 		int result_idx;
 
-		__device__ __forceinline__ Warp_appNearestSearch(const BatchType& batch_arg, int query_index_arg) 
+		__device__ __forceinline__ Warp_appNearestSearch(const Batch& batch_arg, int query_index_arg) 
 			: batch(batch_arg), query_index(query_index_arg){}
 
 		__device__ __forceinline__ void launch(bool active)

					*warp_buffer_float = query.z;
				active_query.z = *warp_buffer_float;                            

				int offset = NearestWarpKernel(batch.points + beg, batch.points_step, end - beg, active_query);                    

				if (active_lane == laneId)
					result_idx = beg + offset;
			}
		}

		__device__ __forceinline__ int NearestWarpKernel(const float* points, int points_step, int length, const float3& active_query)
		{                        
			const int STRIDE = warpSize;
			int tid = threadIdx.x;

			volatile float* dist2 = storage.cta_buffer_flt;
			volatile int*   index = storage.cta_buffer_int;
			
			dist2[tid] = pcl::device::numeric_limits<float>::max();

			//serial step

		if (__all(active == false)) 
			return;

		Warp_appNearestSearch search(batch, query_index);
		search.launch(active); 
	}

