 {
     const uint32 BLOCKDIM  = 128;
     const uint32 N_THREADS = BLOCKDIM * 1024;
-    NVBIO_VAR_UNUSED const uint32 N_TESTS   = 128;
+    NVBIO_VAR_UNUSED const uint32 N_TESTS = 128;
 
-    thrust::device_vector<uint32> dst( N_THREADS );
-    uint32* dst_dptr = thrust::raw_pointer_cast(&dst.front());
+    nvbio::vector<device_tag,uint32> dst( N_THREADS );
+    uint32* dst_dptr = nvbio::raw_pointer( dst );
 
     fprintf(stderr, "all test... started\n");
     all_test<N_THREADS,BLOCKDIM,2,N_TESTS>( dst_dptr );

    timer.start();

    for (uint32 i = 0; i < N_RUNS; ++i)
        all_test_kernel<BLOCKDIM,VECDIM,N_TESTS><<<BLOCKDIM, N_THREADS / BLOCKDIM>>>( dst );

    cudaThreadSynchronize();

    timer.stop();


    timer.start();

    for (uint32 i = 0; i < N_RUNS; ++i)
        any_test_kernel<BLOCKDIM,VECDIM,N_TESTS><<<N_THREADS / BLOCKDIM, BLOCKDIM>>>( dst );

    cudaDeviceSynchronize();

    timer.stop();

