 		dst.init(x.size());
 
 		scan(src, dst);
-	}
-}
+
+		std::vector<int> buffer(x.size());
+		dst.readByNum(&buffer[0]);
+
+		int xxx = 0;
+#else
+		int b[] = { 1, 0, 1, 0, 1, 0, 1, 0 };
+		int v[] = { 0, 1, 2, 3, 4, 5, 6, 7 };
+
+		AT_ASSERT(AT_COUNTOF(b) == AT_COUNTOF(v));
+
+		int num = AT_COUNTOF(b);
+
+		std::vector<int> buffer(num);
+
+		idaten::TypedCudaMemory<int> bools;
+		bools.init(num);
+		bools.writeByNum(b, num);
+
+		idaten::TypedCudaMemory<int> indices;
+		indices.init(num);
+
+		scan(bools, indices);
+
+		indices.readByNum(&buffer[0]);
+
+		idaten::TypedCudaMemory<int> values;
+		values.init(num);
+		values.writeByNum(v, num);
+
+		idaten::TypedCudaMemory<int> dst;
+		dst.init(num);
+
+		idaten::TypedCudaMemory<int> count;
+		count.init(1);
+
+		int blockPerGrid = (num - 1) / blocksize + 1;
+
+		scatter << <blockPerGrid, blocksize >> > (
+			dst.ptr(),
+			count.ptr(),
+			dst.maxNum(),
+			bools.ptr(),
+			indices.ptr(),
+			values.ptr());
+
+		dst.readByNum(&buffer[0]);
+
+		int _count = -1;
+		count.readByNum(&_count);
+
+		int xxx = 0;
+#endif
+	}
+}


// NOTE
// https://developer.nvidia.com/gpugems/GPUGems3/gpugems3_ch39.html

// ubNPʂŌvZ exclusiveScan ̑al𑫂̂vZ.
__global__ void computeBlockCount(

	int index = threadIdx.x;
	int offset = 1;

	auto n = blockIdx.x * blockDim.x + threadIdx.x;
	if (n * 2 >= num) {
		return;
	}

	// Copy input data to shared memory
	temp[2 * index] = src[2 * index + (blockIdx.x * blockDim.x * 2)];
	temp[2 * index + 1] = src[2 * index + 1 + (blockIdx.x * blockDim.x * 2)];
