 		destroy<double><<<1,1000>>>( deviceAllocator, p, 1000 );
 		CUDA_CHECK_ERRORS();
 		CUDA_CALL( cudaDeviceSynchronize() );
+		test_device_memory_validity( p );
 		deviceAllocator.deallocate( p, 1000 );
-		test_device_memory_validity( p );
 	}
 
 	{

		ecuda::device_pitch_allocator<int> deviceAllocator3( deviceAllocator2 );
		//typename ecuda::device_pitch_allocator<double>::size_type pitch;
		typename ecuda::device_pitch_allocator<double>::pointer p( deviceAllocator1.allocate( 100, 100 ) );
		typename ecuda::device_pitch_allocator<double>::pointer q( p );
		for( std::size_t i = 0; i < 1000; ++i, ++q ) deviceAllocator.construct( q, static_cast<typename ecuda::device_pitch_allocator<double>::value_type>(i) );
		q = p;
		destroy<double><<<1,1000>>>( deviceAllocator, p, 1000 );
		CUDA_CHECK_ERRORS();
		CUDA_CALL( cudaDeviceSynchronize() );
		deviceAllocator.deallocate( p, 1000 );
	}

	return 0;

#include "../include/ecuda/allocators.hpp"
#include "../include/ecuda/type_traits.hpp"

#define THREADS 320

template<typename T,class Alloc>
__global__ void destroy( Alloc alloc, typename Alloc::pointer p, std::size_t n ) {
