 		CUDA_CHECK(cudaStreamDestroy(*stream));
 		delete stream;
 	}
-}
-
-__global__ void print(double* d)
-{
-	int idx = blockIdx.x * blockDim.x + threadIdx.x;
-	printf("%d %f\n", idx, d[idx]);
 }
 
 CUDA_HOST

		return;
	}

	// We need to create a stream if no stream is specified since dereferencing
	// a null pointer is, well, not a good thing.

	bool destroyStream = false;

	if (!stream)
	{
		stream = new cudaStream_t;
		CUDA_CHECK(cudaStreamCreate(stream));
		destroyStream = true;
	}

	// Allocate memory on device for packed data

	unsigned char* d_p = 0;  // device-packed data

	{
		d_arr = arr;
	}

	int* d_b = 0;  // device-bitmap

