 				work.maxNum(),
 				blocksize,
 				input->ptr());
-
-			work.readByNum(&buffer[0]);
 
 			if (innerBlockPerGrid <= 1) {
 				cudaMemcpy(tmp.ptr(), work.ptr(), work.bytes(), cudaMemcpyDeviceToDevice);

}

namespace idaten {
	void compact()
	{
		int f[] = { 3, 1, 7, 0, 4, 1, 6, 3, 3, 1, 7, 0, 4, 1, 6, 3, 3, 1, 7, 0, 4, 1, 6, 3, 3, 1, 7, 0, 4, 1, 6, 3 };
		//int f[] = { 3, 1, 7, 0, 4, 1, 6, 3, 3, 1 };
		//int f[] = { 3, 1, 7, 0, 4, 1, 6, 3 };
		//int f[] = { 0, 25, 25, 25 };

		//int c = aten::nextPow2(AT_COUNTOF(f));
		int c = AT_COUNTOF(f);

		std::vector<int> x(c);
		memcpy(&x[0], f, sizeof(int) * AT_COUNTOF(f));

		idaten::TypedCudaMemory<int> src;
		src.init(x.size());
		src.writeByNum(&x[0], x.size());

		idaten::TypedCudaMemory<int> dst;
		dst.init(x.size());

		std::vector<int> buffer(1024);

		int blocksize = 8;
		int blockPerGrid = (x.size() - 1) / blocksize + 1;

		exclusiveScan<<<blockPerGrid, blocksize / 2, blocksize  * sizeof(int)>>>(
			dst.ptr(), 
			dst.maxNum(),
			blocksize, 
			src.ptr());

		dst.readByNum(&buffer[0]);

		idaten::TypedCudaMemory<int> incr;
		incr.init(blockPerGrid);

				input->ptr(),
				work.ptr());

			// swap.
			auto p = input;
			input = output;
			output = p;

			elementNum = innerBlockPerGrid;
			count++;
		}

#if 1
		input = &tmp;
		output = &incr;

		for (int i = count - 1; i >= 0; i--) {
			// blocks per grid.
			auto bpg = stackBlockPerGrid[i];

			auto threadPerBlock = (output->maxNum() + bpg - 1) / bpg;

			incrementBlocks << <bpg, threadPerBlock >> > (
				output->ptr(),
				output->maxNum(),
				input->ptr());

			// swap.
			auto p = input;
			input = output;
			output = p;
		}

		idaten::TypedCudaMemory<int>* incrResult = (count & 0x1 == 0 ? &tmp : &incr);
#endif

		incrementBlocks << <blockPerGrid, blocksize >> > (
			dst.ptr(),
			dst.maxNum(),
			incrResult->ptr());
	}

	void compact()
	{
		int f[] = { 3, 1, 7, 0, 4, 1, 6, 3, 3, 1, 7, 0, 4, 1, 6, 3, 3, 1, 7, 0, 4, 1, 6, 3, 3, 1, 7, 0, 4, 1, 6, 3, 3, 1, 7, 0, 4, 1, 6, 3 };
		//int f[] = { 3, 1, 7, 0, 4, 1, 6, 3, 3, 1 };
		//int f[] = { 3, 1, 7, 0, 4, 1, 6, 3 };
		//int f[] = { 0, 25, 25, 25 };

		//int c = aten::nextPow2(AT_COUNTOF(f));
		int c = AT_COUNTOF(f);

		std::vector<int> x(c);
		memcpy(&x[0], f, sizeof(int) * AT_COUNTOF(f));

		idaten::TypedCudaMemory<int> src;
		src.init(x.size());
		src.writeByNum(&x[0], x.size());

		idaten::TypedCudaMemory<int> dst;
		dst.init(x.size());

		scan(src, dst);
	}
}
