 }
 
 namespace idaten {
-	void compact()
+
+	void scan(
+		idaten::TypedCudaMemory<int>& src,
+		idaten::TypedCudaMemory<int>& dst)
 	{
-		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(), 
+		int blockPerGrid = (dst.maxNum() - 1) / blocksize + 1;
+
+		exclusiveScan << <blockPerGrid, blocksize / 2, blocksize * sizeof(int) >> > (
+			dst.ptr(),
 			dst.maxNum(),
-			blocksize, 
+			blocksize,
 			src.ptr());
 
-		dst.readByNum(&buffer[0]);
+		if (blockPerGrid <= 1) {
+			return;
+		}
 
 		idaten::TypedCudaMemory<int> incr;
 		incr.init(blockPerGrid);

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

			output->readByNum(&buffer[0]);

			auto tmp = input;
			input = output;
			output = tmp;

			elementNum = innerBlockPerGrid;
			count++;
		}

		tmp.readByNum(&buffer[0]);

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

		for (int i = count - 1; i >= 0; i--) {
			auto bpg = stackBlockPerGrid[i];

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

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

			input->readByNum(&buffer[0]);
			output->readByNum(&buffer[0]);

			auto p = input;
			input = output;
			output = p;
		}

		checkCudaErrors(cudaDeviceSynchronize());

		incr.readByNum(&buffer[0]);
#endif

		incrementBlocks << <blockPerGrid, blocksize >> > (
			dst.ptr(),
			dst.maxNum(),
			incr.ptr());

		dst.readByNum(&buffer[0]);

		int xxx = 0;
	}
}

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

			if (innerBlockPerGrid <= 1) {
				cudaMemcpy(tmp.ptr(), work.ptr(), work.bytes(), cudaMemcpyDeviceToDevice);
