 		return;
 	}
 
-	dst[index] = src0[(index + 1) * stride - 1] + src1[(index + 1) * stride - 1];
-
 	if (index == 0) {
 		dst[index] = 0;
+	}
+	else {
+		dst[index] = src0[index * stride - 1] + src1[index * stride - 1];
 	}
 }
 

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

		idaten::TypedCudaMemory<int> src;
		src.init(AT_COUNTOF(x));
		src.writeByNum(x, AT_COUNTOF(x));

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

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

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


		offset >>= 1;
		__syncthreads();

		if (index < d && offset > 0) {
			int ai = offset * (2 * index + 1) - 1;
			int bi = offset * (2 * index + 2) - 1;
			int t = temp[ai];
