Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -72,8 +72,8 @@ public class HeightMapExtractor
private final FloatPointer parametersHostPointer;
private final FloatPointer parametersDevicePointer;

private int centerIndex;
private int cellsPerAxis;
private final int centerIndex;
private final int cellsPerAxis;
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Cleanup


private final RigidBodyTransform previousSensorToWorld = new RigidBodyTransform();
private int previousCellX;
Expand Down Expand Up @@ -160,7 +160,6 @@ public void reset(double footHeight, float loweredValue)
resetOffset -= loweredValue;

globalMeanMap.setTo(new Scalar(resetOffset));
emptyGlobalHeightMap.setTo(new Scalar(resetOffset));
}

public void update(GpuMat latestDepthImageGPU,
Expand Down Expand Up @@ -315,19 +314,19 @@ public void update(GpuMat latestDepthImageGPU,
}
}

float zeroValueForEmptySpaces = 0.0f;
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ok so the idea behind the next bit of code is to run a kernel, that's going to store the entire height map area. But the only non-zero values will be squares that have data from the local maps. At the beginning we set the entire map to be 0.0 (zeroValueForEmptySpaces) and the values get overwritten if there is data in the cell on the localMeanMap (local height map).

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

private static final float ZERO_FOR_EMPTY_SPACES = 0.0f; at top of file?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm ok to leave it if you are. I like it being close to the code that's using it.

// ---------- Run the registration kernel for an empty global height map ----------
{
int emptyRegistrationGridSizeXY = (cellsPerAxis + BLOCK_SIZE_XY - 1) / BLOCK_SIZE_XY;
dim3 registerKernelGridDim = new dim3(emptyRegistrationGridSizeXY, emptyRegistrationGridSizeXY, 1);

// Need to reset the empty global map before using it so when its filled it starts with all "zero" values
emptyGlobalHeightMap.setTo(new Scalar(resetOffset));
emptyGlobalHeightMap.setTo(new Scalar(zeroValueForEmptySpaces));

emptyRegisterKernel.withPointer(localMeanMap.data()).withLong(localMeanMap.step());
emptyRegisterKernel.withPointer(emptyGlobalHeightMap.data()).withLong(emptyGlobalHeightMap.step());
emptyRegisterKernel.withPointer(zUpCameraToWorldAlignedGroundDevicePointer);
emptyRegisterKernel.withPointer(parametersDevicePointer);
emptyRegisterKernel.withFloat(resetOffset);

emptyRegisterKernel.run(stream, registerKernelGridDim, blockSize, 0);

Expand All @@ -344,7 +343,7 @@ public void update(GpuMat latestDepthImageGPU,
planOffsetKernel.withPointer(globalMeanMap.data()).withLong(globalMeanMap.step());
planOffsetKernel.withPointer(emptyGlobalHeightMap.data()).withLong(emptyGlobalHeightMap.step());
planOffsetKernel.withFloat(zDriftInMeters);
planOffsetKernel.withFloat(resetOffset);
planOffsetKernel.withFloat(zeroValueForEmptySpaces);
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Not sure why we were comparing the value to this resetOffset guy. We want to compare it to the zero value in the emptyGlobalHeightMap that we just computed in the kernel before this.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This zerValueForEmptySpaces is always 0.0, right? I don't think you need to pass it to the kernel then, just hard code the value in the kernel. Less memory to transfer between CPU and GPU when launching the kernel that way.

Can like #define ZERO_FOR_EMPTY_SPACES 0.0f or const float ZERO_FOR_EMPTY_SPACES = 0.0f;

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That's an interesting thought. I'm happy to try that as a future improvement to see what the memory transfer would be. Since its currently one float (32 bits). I'm not too worried about it and I find it easier to understand if the field lives in Java so its visible to the user.

planOffsetKernel.withPointer(parametersDevicePointer);

planOffsetKernel.run(stream, planOffsetKernelGridDim, blockSize, 0);
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -277,7 +277,7 @@ extern "C"
__global__ void heightMapEmptyRegistrationKernel(float *localMap, size_t pitchLocal,
float *globalMap, size_t pitchGlobal,
float *zUpCameraToWorldAlignedGround,
float *params, float resetOffset)
float *params)
{
int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
Expand All @@ -297,6 +297,7 @@ __global__ void heightMapEmptyRegistrationKernel(float *localMap, size_t pitchLo

float *localHeight = (float *)((char *)localMap + localIndex.x * pitchLocal) + localIndex.y;

// This is a way of checking if we have any read data in the spot, its very unlikely that any real data is 0.0
if (*localHeight == 0.0f)
return;

Expand All @@ -305,9 +306,17 @@ __global__ void heightMapEmptyRegistrationKernel(float *localMap, size_t pitchLo
}

extern "C"
/**
@brief Compute Plan Offset KERNEL: This kernel is not very intuitive. For background we've got a status message
being published that gives us an offset in Z in meters that the robot thinks its drifted by.
Our goal is to update the height map accordingly. However, it doesn't make sense to update any values that are
in the live view of the camera, because: a) they will just get overwritten the next image, and b) that data doesn't drift
because its live data. So we pass in this matrix of values to skip because those values are outside the live view
of the camera.
*/
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let me know if this makes sense!

__global__ void planOffsetKernel(float *matrixToModify, size_t pitchMatrixToModify,
float *matrixValuesToSkip, size_t pitchMatrixValuesToSkip,
float offsetInZ, float resetOffset, float *params)
float offsetInZ, float zeroValueForEmptySpaces, float *params)
{
int indexX = blockIdx.x * blockDim.x + threadIdx.x;
int indexY = blockIdx.y * blockDim.y + threadIdx.y;
Expand All @@ -318,9 +327,8 @@ __global__ void planOffsetKernel(float *matrixToModify, size_t pitchMatrixToModi
return;

float *skipRow = (float *)((char *)matrixValuesToSkip + indexX * pitchMatrixValuesToSkip);
// This is less then or equal to due to a round error that can give +- 1 offsets
// This skips the cells that have real data in them coming from the values to skip
if (fabsf(skipRow[indexY] - resetOffset) >= 2.0f)
if (skipRow[indexY] != zeroValueForEmptySpaces)
return;

float *matrixRow = (float *)((char *)matrixToModify + indexX * pitchMatrixToModify);
Expand Down
Loading