Skip to content

Commit

Permalink
Still improving OS compatibility and other corrections
Browse files Browse the repository at this point in the history
  • Loading branch information
MarianoJT88 committed Jan 17, 2015
1 parent 153e739 commit e0f177b
Show file tree
Hide file tree
Showing 10 changed files with 146 additions and 127 deletions.
3 changes: 1 addition & 2 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -34,6 +34,7 @@ ENDIF (BUILD_EVALUATOR)
SET(DEFAULT_BUILD_RT_VISUALIZATION ON)
SET(BUILD_RT_VISUALIZATION ${DEFAULT_BUILD_RT_VISUALIZATION} CACHE BOOL "Build the scene flow estimator with real-time visualization that requires MRPT and OpenNI2")

#Build real-time scene flow with visualization
IF (BUILD_RT_VISUALIZATION)
FIND_PACKAGE(MRPT REQUIRED base gui opengl)

Expand All @@ -50,8 +51,6 @@ IF (BUILD_RT_VISUALIZATION)
SET(OpenNI_lib "${OpenNI2_libdir}/libOpenNI2.so")
ENDIF (UNIX)

message( "esto vale $ENV{OPENNI2_INCLUDE}" )

ADD_EXECUTABLE(Scene-Flow-Visualization
main_scene_flow_visualization.cpp
scene_flow_visualization.cpp
Expand Down
10 changes: 10 additions & 0 deletions README.txt
Original file line number Diff line number Diff line change
@@ -0,0 +1,10 @@
Number of iterations at every level

Number of threads and blocks - Threads should be higher than 25 for the gaussian mask computation

Dependencies
OpenCV
MRPT
OpenNI2

Only two camera modes
7 changes: 2 additions & 5 deletions main_scene_flow_impair.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -21,9 +21,6 @@
** **
*****************************************************************************/

#include <stdio.h>
#include <string.h>
//#include <stdlib.h>
#include "scene_flow_impair.h"

using namespace std;
Expand All @@ -36,9 +33,9 @@ using namespace std;
int main(int num_arg, char *argv[])
{
//==============================================================================
// Read function arguments
// Read arguments
//==============================================================================
unsigned int rows = 240; //Default values
unsigned int rows = 240; //Default value

if (num_arg <= 1); //No arguments
else if ( string(argv[1]) == "--help")
Expand Down
4 changes: 0 additions & 4 deletions main_scene_flow_visualization.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -21,9 +21,6 @@
** **
*****************************************************************************/

#include <stdio.h>
#include <string.h>
//#include <stdlib.h>
#include "scene_flow_visualization.h"


Expand Down Expand Up @@ -61,7 +58,6 @@ int main(int num_arg, char *argv[])

if ( string(argv[i]) == "--fps")
fps = stoi(argv[i+1]);
//fps = strtof(argv[i+1], NULL);

if ( string(argv[i]) == "--rows")
rows = stoi(argv[i+1]);
Expand Down
29 changes: 15 additions & 14 deletions pdflow_cudalib.cu
Original file line number Diff line number Diff line change
Expand Up @@ -42,7 +42,7 @@ __host__ void CSF_cuda::allocateDevMemory()

for (unsigned int i = 0; i<pyr_levels; i++)
{
s = powf(2,i);
s = static_cast<unsigned int>(powf(2,i));
cudaMalloc((void**)&colour_dev[i], width*height*sizeof(float)/(s*s) );
cudaMalloc((void**)&colour_old_dev[i], width*height*sizeof(float)/(s*s) );
cudaMalloc((void**)&depth_dev[i], width*height*sizeof(float)/(s*s) );
Expand Down Expand Up @@ -189,7 +189,7 @@ __host__ void CSF_cuda::readParameters(unsigned int rows_host, unsigned int cols
fovv = fovv_host;
f_dist = f_dist_host;

//Allocate gaussian mask
//Allocate and copy gaussian mask
cudaError_t err = cudaMalloc((void**)&g_mask_dev, 5*5*sizeof(float));
//printf("%s", cudaGetErrorString(err));
cudaMemcpy(g_mask_dev, g_mask, 5*5*sizeof(float), cudaMemcpyHostToDevice);
Expand Down Expand Up @@ -802,7 +802,7 @@ __device__ void CSF_cuda::computeStepSizes(unsigned int index)
//=============================================================================
__device__ void CSF_cuda::updateDualVariables(unsigned int index)
{
//Create aux variables
//Create aux variables to avoid repetitive global memory access
float module_p;
float pd = pd_dev[index], puu = puu_dev[index], puv = puv_dev[index];
float pvu = pvu_dev[index], pvv = pvv_dev[index], pwu = pwu_dev[index], pwv = pwv_dev[index];
Expand Down Expand Up @@ -1035,7 +1035,8 @@ __device__ void CSF_cuda::filterSolution(unsigned int index)
const unsigned int v = index%rows_i;
const unsigned int u = index/rows_i;

//Median filter
// Weighted median filter
//----------------------------------------------------------------------------------------
fieldAndPresence up[9], vp[9], wp[9];
float pres_cum_u[9], pres_cum_v[9], pres_cum_w[9], pres_med;
int indr, indc, ind_loop;
Expand All @@ -1060,7 +1061,7 @@ __device__ void CSF_cuda::filterSolution(unsigned int index)
continue;
}

//Compute presence of every element
//Compute weights
const float pres = 1.f/(1.f + kd*powf(depth_old - depth_old_dev[level_image][ind_loop],2.f) + kddt*powf(ddt_dev[ind_loop],2.f));

up[v_index].field = du_new_dev[ind_loop]; up[v_index].pres = pres;
Expand All @@ -1069,12 +1070,12 @@ __device__ void CSF_cuda::filterSolution(unsigned int index)
v_index++;
}

//Sort vectors (both the solution and the presence)
//Sort vectors (both the solution and the weights)
bubbleSortDev(up, point_count);
bubbleSortDev(vp, point_count);
bubbleSortDev(wp, point_count);

//Compute cumulative presence
//Compute cumulative weight
pres_cum_u[0] = up[0].pres; pres_cum_v[0] = vp[0].pres; pres_cum_w[0] = wp[0].pres;
for (unsigned int i=1; i<point_count; i++)
{
Expand Down Expand Up @@ -1138,13 +1139,12 @@ __device__ void CSF_cuda::filterSolution(unsigned int index)

__device__ void CSF_cuda::computeMotionField(unsigned int index)
{
//Fill the matrices dx,dy,dz with the final solution
const float x_incr = 2.f*f_dist*tanf(0.5f*fovh)/(cols_i-1); //In meters
const float y_incr = 2.f*f_dist*tanf(0.5f*fovv)/(rows_i-1); //In meters
const float x_incr = 2.f*f_dist*tanf(0.5f*fovh)/(cols_i-1);
const float y_incr = 2.f*f_dist*tanf(0.5f*fovv)/(rows_i-1);
const float fx = f_dist/x_incr;
const float fy = f_dist/y_incr;


//Fill the matrices dx,dy,dz with the scene flow estimate
if (depth_old_dev[level_image][index] > 0)
{
dx_dev[index] = dw_l_dev[index];
Expand Down Expand Up @@ -1240,7 +1240,7 @@ void MotionFieldBridge(CSF_cuda *csf)

void DebugBridge(CSF_cuda *csf_device)
{
printf("Entro en el debug bridge");
printf("Executing debug kernel");
DebugKernel <<<1,1>>>(csf_device);
}

Expand All @@ -1255,8 +1255,9 @@ void BridgeBack(CSF_cuda *csf_host, CSF_cuda *csf_device)
//=============================================================================
__global__ void DebugKernel(CSF_cuda *csf)
{
printf("\n dx: ");
for (unsigned int i = 0; i< (csf->rows_i)*(csf->cols_i); i++) //(csf->rows)*(csf->cols)
//Add here the code you want to use for debugging
printf("\n dx: ");
for (unsigned int i = 0; i< (csf->rows_i)*(csf->cols_i); i++)
printf(" %f", csf->dx_dev[i]);

}
Expand Down
4 changes: 2 additions & 2 deletions pdflow_cudalib.h
Original file line number Diff line number Diff line change
Expand Up @@ -25,8 +25,8 @@
#include <cuda_runtime.h>
#include <cuda.h>

#define N_blocks 256 //128
#define N_threads 128 //256
#define N_blocks 256
#define N_threads 128

//Warning!!!!!! Number of threads should be higher than 25 - See definition of "computePyramidLevel()"

Expand Down
Loading

0 comments on commit e0f177b

Please sign in to comment.