Tune CT a little

lib/libsgm/src/census_transform.cu

@@ -118,41 +118,23 @@ __global__ void census_transform_kernel(
 	const int x = (blockIdx.x*blockDim.x + threadIdx.x);
 	const int y = blockIdx.y*blockDim.y + threadIdx.y;
 
-	dest[x+y*width] = 0;
+	feature_type res = 0;
 
 	if (x >= RADIUS_X && y >= RADIUS_Y && x < width-RADIUS_X && y < height-RADIUS_Y) {
-		short center = src[y*pitch+x];
-		uint8_t i = 0; // bit counter for *out
-		// possible BUG in operator(), gets called more than once per pixel;
-		// local variable for sub-bitstring to avoid data race (no read
-		// dependency to out; writes are identical)
-		feature_type res = 0;
+		const T center = src[y*pitch+x];
 
-		for (int wy = -RADIUS_Y; wy <= RADIUS_Y; wy++) {
-			for (int wx = -RADIUS_X; wx <= RADIUS_X; wx++) {
-				const int y_ = y + wy;
-				const int x_ = x + wx;
+		#pragma unroll
+		for (int wy = -RADIUS_Y; wy <= RADIUS_Y; ++wy) {
+			const int i = (y + wy) * pitch + x;
 
-				if (y == 0 && x == 0) {
-					continue;
+			#pragma unroll
+			for (int wx = -RADIUS_X; wx <= RADIUS_X; ++wx) {
+				res = (res << 1) | (center < (src[i+wx]) ? 1 : 0);
 			}
-
-				// zero if first value, otherwise shift to left
-				res = (res << 1);
-				res |= (center < (src[y_*pitch+x_]) ? 1 : 0);
-
-				// if all bits set, continue to next element
-				/*if (++i % 64 == 0) {
-					*out = res;
-					out++;
-				}*/
 		}
 	}
-		//if ((i - 1)%64 != 0) {
-			// write remaining bits
+
 	dest[x+y*width] = res;
-		//}
-	}
 }
 
 template <typename T>