Warp optimise correspondence search

Does a pixel by pixel correspondence search organised to improve memory caching and access by organising around a warp

components/operators/src/correspondence.cu

@@ -98,13 +98,8 @@ __global__ void corresponding_point_kernel(
         float4x4 pose2,  // Inverse
         Camera cam1,
         Camera cam2, ftl::cuda::MvMLSParams params) {
-
-    // Each warp picks point in p1
-    //const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
-	//const int x = (blockIdx.x*blockDim.x + threadIdx.x); // / WARP_SIZE;
-	//const int y = blockIdx.y*blockDim.y + threadIdx.y;
 	
-	const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
+	//const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
 	const int x = (blockIdx.x*8 + (threadIdx.x%4) + 4*(threadIdx.x / 16)); // / WARP_SIZE;
 	const int y = blockIdx.y*8 + threadIdx.x/4 + 4*threadIdx.y;
 
@@ -127,19 +122,13 @@ __global__ void corresponding_point_kernel(
         //float bestdepth = 0.0f;
         short2 bestScreen = make_short2(-1,-1);
 		float bestdepth = 0.0f;
-		float bestdepth2 = 0.0f;
+		//float bestdepth2 = 0.0f;
         float bestweight = 0.0f;
         float bestcolour = 0.0f;
-        float bestdweight = 0.0f;
+        //float bestdweight = 0.0f;
         float totalcolour = 0.0f;
         int count = 0;
-        float contrib = 0.0f;
-        
-        //const float step_interval = params.range / (COR_STEPS / 2);
-        
-        //const float3 rayStep_world = pose1.getFloat3x3() * cam1.screenToCam(x,y,step_interval);
-        //const float3 rayStart_2 = pose2 * world1;
-		//const float3 rayStep_2 = pose2.getFloat3x3() * rayStep_world;
+        //float contrib = 0.0f;
 		
 		const float3 camPosOrigin = pose2 * world1;
         const float2 lineOrigin = cam2.camToScreen<float2>(camPosOrigin);
@@ -158,15 +147,7 @@ __global__ void corresponding_point_kernel(
         // Project to p2 using cam2
         // Each thread takes a possible correspondence and calculates a weighting
         //const int lane = tid % WARP_SIZE;
-        for (int i=0; i<COR_STEPS; ++i) {
-            //const int j = i - (COR_STEPS/2);
-            //const float depth_adjust = (float)j * step_interval;
-
-            // Calculate adjusted depth 3D point in camera 2 space
-            //const float3 worldPos = world1 + j * rayStep_world; //(pose1 * cam1.screenToCam(x, y, depth_adjust));
-            //const float3 camPos = rayStart_2 + j * rayStep_2; //pose2 * worldPos;
-			//const float2 screen = cam2.camToScreen<float2>(camPos);
-			
+        for (int i=0; i<COR_STEPS; ++i) {			
 			float weight = (linePos.x >= cam2.width || linePos.y >= cam2.height) ? 0.0f : 1.0f;
 
 			// Generate a colour correspondence value
@@ -187,7 +168,7 @@ __global__ void corresponding_point_kernel(
             //weight *= weightFunction<FUNCTION>(params, dweight, cweight);
 
             ++count;
-            contrib += weight;
+            //contrib += weight;
             bestcolour = max(cweight, bestcolour);
             //bestdweight = max(dweight, bestdweight);
             totalcolour += cweight;
@@ -200,10 +181,10 @@ __global__ void corresponding_point_kernel(
                 //bestScreen = make_short2(screen.x+0.5f, screen.y+0.5f);
 			//}
 			
-			depthPos += depthM*float(INTERVAL);
-			depthPos2 += depthM2*float(INTERVAL);
-            linePos.x += float(INTERVAL);
-            linePos.y += lineM*float(INTERVAL);
+			depthPos += depthM;
+			depthPos2 += depthM2;
+            linePos.x += 1.0f;
+            linePos.y += lineM;
         }
 
         const float avgcolour = totalcolour/(float)count;