- Fixed problem for CUDA 9.0 compatibility: the Census kernel was com…

…piled with too many registers and it was not able to run, because of the 1024 threads per block requirement (not enough threads)

- Fixed an indexing bug, thanks to @ankit1089sny

cost_aggregation.h

@@ -46,7 +46,7 @@ __device__ __forceinline__ void CostAggregationGenericIndexesIncrement(int *inde
 template<int add_index, bool recompute, bool join_dispcomputation>
 __device__ __forceinline__ void CostAggregationDiagonalGenericIndexesIncrement(int *index, int *index_im, int *col, const int cols, const int initial_row, const int i, const int dis) {
 	*col += add_index;
-	if(add_index > 0 && *col > cols) {
+	if(add_index > 0 && *col >= cols) {
 		*col = 0;
 	} else if(*col < 0) {
 		*col = cols-1;

costs.cu

@@ -21,7 +21,9 @@
 #include "costs.h"
 #include <stdio.h>
 
-__global__ void CenterSymmetricCensusKernelSM2(const uint8_t *im, const uint8_t *im2, cost_t *transform, cost_t *transform2, const uint32_t rows, const uint32_t cols) {
+__global__ void 
+__launch_bounds__(1024, 2)
+CenterSymmetricCensusKernelSM2(const uint8_t *im, const uint8_t *im2, cost_t *transform, cost_t *transform2, const uint32_t rows, const uint32_t cols) {
 	const int idx = blockIdx.x*blockDim.x+threadIdx.x;
 	const int idy = blockIdx.y*blockDim.y+threadIdx.y;