 					*warp_buffer_float = query.z;
 				active_query.z = *warp_buffer_float;                            
 
-				int offset = NearestWarpKernel(batch.points + beg, batch.points_step, end - beg, active_query);                    
+				int offset = NearestWarpKernel<KernelPolicy::CTA_SIZE>(batch.points + beg, batch.points_step, end - beg, active_query);                    
 
 				if (active_lane == laneId)
 					result_idx = beg + offset;
 			}
 		}
 
+        template<int CTA_SIZE>
 		__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;
-			
+            __shared__ volatile float dist2[CTA_SIZE];
+            __shared__ volatile int   index[CTA_SIZE];
+			            			
 			dist2[tid] = pcl::device::numeric_limits<float>::max();
 
 			//serial step

#include "utils/boxutils.hpp"
#include "utils/scan_block.hpp"

#include "octree_iterator.hpp"


namespace pcl { namespace device { namespace appnearest_search
{   
	template<typename PointT>
	struct Batch
	{   
		typedef PointT PointType;
		const PointType* queries;

		const int *indices;


		__device__ __forceinline__ void processNode(int node_idx)
		{   
            __shared__ volatile int  per_warp_buffer[KernelPolicy::WARPS_COUNT];

			int mask = __ballot(node_idx != -1);                        

			while(mask)
			{                
