Skip to content

Commit

Permalink
Fix Gaussian noise: transforms updates & distance noise calculation (#…
Browse files Browse the repository at this point in the history
…217)

* Add tests

* Update lookAtOriginTransform every gaussian noise nodes exec

* Fix distance Gaussian noise calculation

* Refactor test

* Review change

* Review changes

* Fix bug with distance calculation

* Grammar correction
  • Loading branch information
msz-rai authored Nov 19, 2023
1 parent 7803c2f commit 2e91c81
Show file tree
Hide file tree
Showing 8 changed files with 162 additions and 61 deletions.
23 changes: 10 additions & 13 deletions src/gpu/gaussianNoiseKernels.cu
Original file line number Diff line number Diff line change
Expand Up @@ -48,24 +48,21 @@ __global__ void kAddGaussianNoiseAngularHitpoint(size_t pointCount, float mean,

__global__ void kAddGaussianNoiseDistance(size_t pointCount, float mean, float stDevBase, float stDevRisePerMeter,
Mat3x4f lookAtOriginTransform, curandStatePhilox4_32_10_t* randomStates,
const Field<XYZ_VEC3_F32>::type* inPoints, Field<XYZ_VEC3_F32>::type* outPoints,
const Field<XYZ_VEC3_F32>::type* inPoints,
const Field<DISTANCE_F32>::type* inDistances, Field<XYZ_VEC3_F32>::type* outPoints,
Field<DISTANCE_F32>::type* outDistances)
{
LIMIT(pointCount);

Field<XYZ_VEC3_F32>::type originPoint = lookAtOriginTransform * inPoints[tid];

float distance = originPoint.length();
float distanceInducedStDev = distance * stDevRisePerMeter;
float distanceInducedStDev = inDistances[tid] * stDevRisePerMeter;
float totalStDev = distanceInducedStDev + stDevBase;

float distanceError = mean + curand_normal(&randomStates[tid]) * totalStDev;

if (outDistances != nullptr) {
outDistances[tid] = distance + distanceError;
}
Field<XYZ_VEC3_F32>::type pointInRayOriginTransform = lookAtOriginTransform * inPoints[tid];

outPoints[tid] = inPoints[tid] + inPoints[tid].normalize() * distanceError;
outPoints[tid] = lookAtOriginTransform.inverse() *
(pointInRayOriginTransform + pointInRayOriginTransform.normalize() * distanceError);
outDistances[tid] = inDistances[tid] + distanceError;
}

void gpuAddGaussianNoiseAngularRay(cudaStream_t stream, size_t rayCount, float mean, float stDev, rgl_axis_t rotationAxis,
Expand All @@ -87,9 +84,9 @@ void gpuAddGaussianNoiseAngularHitpoint(cudaStream_t stream, size_t pointCount,

void gpuAddGaussianNoiseDistance(cudaStream_t stream, size_t pointCount, float mean, float stDevBase, float stDevRisePerMeter,
Mat3x4f lookAtOriginTransform, curandStatePhilox4_32_10_t* randomStates,
const Field<XYZ_VEC3_F32>::type* inPoints, Field<XYZ_VEC3_F32>::type* outPoints,
Field<DISTANCE_F32>::type* outDistances)
const Field<XYZ_VEC3_F32>::type* inPoints, const Field<DISTANCE_F32>::type* inDistances,
Field<XYZ_VEC3_F32>::type* outPoints, Field<DISTANCE_F32>::type* outDistances)
{
run(kAddGaussianNoiseDistance, stream, pointCount, mean, stDevBase, stDevRisePerMeter, lookAtOriginTransform, randomStates,
inPoints, outPoints, outDistances);
inPoints, inDistances, outPoints, outDistances);
}
4 changes: 2 additions & 2 deletions src/gpu/gaussianNoiseKernels.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -27,5 +27,5 @@ void gpuAddGaussianNoiseAngularHitpoint(cudaStream_t stream, size_t pointCount,
Field<XYZ_VEC3_F32>::type* outPoints, Field<DISTANCE_F32>::type* outDistances);
void gpuAddGaussianNoiseDistance(cudaStream_t stream, size_t pointCount, float mean, float stDevBase, float stDevRisePerMeter,
Mat3x4f lookAtOriginTransform, curandStatePhilox4_32_10_t* randomStates,
const Field<XYZ_VEC3_F32>::type* inPoints, Field<XYZ_VEC3_F32>::type* outPoints,
Field<DISTANCE_F32>::type* outDistances);
const Field<XYZ_VEC3_F32>::type* inPoints, const Field<DISTANCE_F32>::type* inDistances,
Field<XYZ_VEC3_F32>::type* outPoints, Field<DISTANCE_F32>::type* outDistances);
5 changes: 2 additions & 3 deletions src/graph/GaussianNoiseAngularHitpointNode.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,6 @@ void GaussianNoiseAngularHitpointNode::setParameters(float mean, float stDev, rg
void GaussianNoiseAngularHitpointNode::validateImpl()
{
IPointsNodeSingleInput::validateImpl();
lookAtOriginTransform = input->getLookAtOriginTransform();

// This node will modify field DISTANCE_F32 if present.
// In the future: only one field should be modified.
Expand Down Expand Up @@ -59,8 +58,8 @@ void GaussianNoiseAngularHitpointNode::enqueueExecImpl()
const auto* inXyzPtr = input->getFieldDataTyped<XYZ_VEC3_F32>()->asSubclass<DeviceAsyncArray>()->getReadPtr();
auto* outXyzPtr = outXyz->getWritePtr();
auto* randPtr = randomizationStates->getWritePtr();
gpuAddGaussianNoiseAngularHitpoint(getStreamHandle(), pointCount, mean, stDev, rotationAxis, lookAtOriginTransform, randPtr,
inXyzPtr, outXyzPtr, outDistancePtr);
gpuAddGaussianNoiseAngularHitpoint(getStreamHandle(), pointCount, mean, stDev, rotationAxis,
input->getLookAtOriginTransform(), randPtr, inXyzPtr, outXyzPtr, outDistancePtr);
}

IAnyArray::ConstPtr GaussianNoiseAngularHitpointNode::getFieldData(rgl_field_t field)
Expand Down
13 changes: 4 additions & 9 deletions src/graph/GaussianNoiseAngularRayNode.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -23,25 +23,20 @@ void GaussianNoiseAngularRaysNode::setParameters(float mean, float stDev, rgl_ax
this->rotationAxis = rotationAxis;
}

void GaussianNoiseAngularRaysNode::validateImpl()
void GaussianNoiseAngularRaysNode::enqueueExecImpl()
{
IRaysNodeSingleInput::validateImpl();
lookAtOriginTransform = input->getCumulativeRayTransfrom().inverse();

// In case rays have changed
auto rayCount = input->getRayCount();
rays->resize(rayCount, false, false);

if (randomizationStates->getCount() < rayCount) {
randomizationStates->resize(rayCount, false, false);
gpuSetupRandomNumberGenerator(getStreamHandle(), rayCount, randomDevice(), randomizationStates->getWritePtr());
}
}

void GaussianNoiseAngularRaysNode::enqueueExecImpl()
{
const auto* inRaysPtr = input->getRays()->asSubclass<DeviceAsyncArray>()->getReadPtr();
auto* outRaysPtr = rays->getWritePtr();
auto* randPtr = randomizationStates->getWritePtr();
gpuAddGaussianNoiseAngularRay(getStreamHandle(), getRayCount(), mean, stDev, rotationAxis, lookAtOriginTransform, randPtr,
inRaysPtr, outRaysPtr);
gpuAddGaussianNoiseAngularRay(getStreamHandle(), getRayCount(), mean, stDev, rotationAxis,
input->getCumulativeRayTransfrom().inverse(), randPtr, inRaysPtr, outRaysPtr);
}
32 changes: 6 additions & 26 deletions src/graph/GaussianNoiseDistanceNode.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -23,52 +23,32 @@ void GaussianNoiseDistanceNode::setParameters(float mean, float stDevBase, float
this->stDevRisePerMeter = stDevRisePerMeter;
}

void GaussianNoiseDistanceNode::validateImpl()
{
IPointsNodeSingleInput::validateImpl();
lookAtOriginTransform = input->getLookAtOriginTransform();

// This node will modifty field DISTANCE_F32 if present.
// In the future: only one field should be modified.
// Other fields that depend on the main field (for now, it's XYZ_VEC3_F32) should be calculated somewhere else (e.g., in data getters nodes).
if (input->hasField(DISTANCE_F32)) {
if (outDistance == nullptr) {
outDistance = DeviceAsyncArray<Field<DISTANCE_F32>::type>::create(arrayMgr);
}
} else {
outDistance.reset();
}
}

void GaussianNoiseDistanceNode::enqueueExecImpl()
{
auto pointCount = input->getPointCount();
outXyz->resize(pointCount, false, false);

Field<DISTANCE_F32>::type* outDistancePtr = nullptr;
if (outDistance != nullptr) {
outDistance->resize(pointCount, false, false);
outDistancePtr = outDistance->getWritePtr();
}
outDistance->resize(pointCount, false, false);

if (randomizationStates->getCount() < pointCount) {
randomizationStates->resize(pointCount, false, false);
gpuSetupRandomNumberGenerator(getStreamHandle(), pointCount, randomDevice(), randomizationStates->getWritePtr());
}

const auto* inXyzPtr = input->getFieldDataTyped<XYZ_VEC3_F32>()->asSubclass<DeviceAsyncArray>()->getReadPtr();
const auto* inDistancePtr = input->getFieldDataTyped<DISTANCE_F32>()->asSubclass<DeviceAsyncArray>()->getReadPtr();
auto* outXyzPtr = outXyz->getWritePtr();
auto* outDistancePtr = outDistance->getWritePtr();
auto* randPtr = randomizationStates->getWritePtr();
gpuAddGaussianNoiseDistance(getStreamHandle(), pointCount, mean, stDevBase, stDevRisePerMeter, lookAtOriginTransform,
randPtr, inXyzPtr, outXyzPtr, outDistancePtr);
gpuAddGaussianNoiseDistance(getStreamHandle(), pointCount, mean, stDevBase, stDevRisePerMeter,
input->getLookAtOriginTransform(), randPtr, inXyzPtr, inDistancePtr, outXyzPtr, outDistancePtr);
}

IAnyArray::ConstPtr GaussianNoiseDistanceNode::getFieldData(rgl_field_t field)
{
if (field == XYZ_VEC3_F32) {
return outXyz;
}
if (field == DISTANCE_F32 && outDistance != nullptr) {
if (field == DISTANCE_F32) {
return outDistance;
}
return input->getFieldData(field);
Expand Down
10 changes: 3 additions & 7 deletions src/graph/NodesCore.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -386,7 +386,6 @@ struct GaussianNoiseAngularRaysNode : IRaysNodeSingleInput
void setParameters(float mean, float stSev, rgl_axis_t rotationAxis);

// Node
void validateImpl() override;
void enqueueExecImpl() override;

// Data getters
Expand All @@ -397,7 +396,6 @@ struct GaussianNoiseAngularRaysNode : IRaysNodeSingleInput
float stDev;
rgl_axis_t rotationAxis;
std::random_device randomDevice;
Mat3x4f lookAtOriginTransform;

DeviceAsyncArray<curandStatePhilox4_32_10_t>::Ptr randomizationStates =
DeviceAsyncArray<curandStatePhilox4_32_10_t>::create(arrayMgr);
Expand Down Expand Up @@ -425,7 +423,6 @@ struct GaussianNoiseAngularHitpointNode : IPointsNodeSingleInput
float stDev;
rgl_axis_t rotationAxis;
std::random_device randomDevice;
Mat3x4f lookAtOriginTransform;

DeviceAsyncArray<curandStatePhilox4_32_10_t>::Ptr randomizationStates =
DeviceAsyncArray<curandStatePhilox4_32_10_t>::create(arrayMgr);
Expand All @@ -440,11 +437,10 @@ struct GaussianNoiseDistanceNode : IPointsNodeSingleInput
void setParameters(float mean, float stDevBase, float stDevRisePerMeter);

// Node
void validateImpl() override;
void enqueueExecImpl() override;

// Node requirements
std::vector<rgl_field_t> getRequiredFieldList() const override { return {XYZ_VEC3_F32}; };
std::vector<rgl_field_t> getRequiredFieldList() const override { return {XYZ_VEC3_F32, DISTANCE_F32}; };

// Data getters
IAnyArray::ConstPtr getFieldData(rgl_field_t field) override;
Expand All @@ -454,10 +450,10 @@ struct GaussianNoiseDistanceNode : IPointsNodeSingleInput
float stDevBase;
float stDevRisePerMeter;
std::random_device randomDevice;
Mat3x4f lookAtOriginTransform;

DeviceAsyncArray<curandStatePhilox4_32_10_t>::Ptr randomizationStates =
DeviceAsyncArray<curandStatePhilox4_32_10_t>::create(arrayMgr);
DeviceAsyncArray<Field<XYZ_VEC3_F32>::type>::Ptr outXyz = DeviceAsyncArray<Field<XYZ_VEC3_F32>::type>::create(arrayMgr);
DeviceAsyncArray<Field<DISTANCE_F32>::type>::Ptr outDistance = nullptr;
DeviceAsyncArray<Field<DISTANCE_F32>::type>::Ptr outDistance = DeviceAsyncArray<Field<DISTANCE_F32>::type>::create(
arrayMgr);
};
2 changes: 1 addition & 1 deletion test/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -8,6 +8,7 @@ set(RGL_TEST_FILES
src/externalLibraryTest.cpp
src/features/loggingTests.cpp
src/gaussianStressTest.cpp
src/gaussianPoseIndependentTest.cpp
src/graph/addChildTest.cpp
src/graph/graphCaseTest.cpp
src/graph/nodes/CompactPointsNodeTest.cpp
Expand Down Expand Up @@ -40,7 +41,6 @@ set(RGL_TEST_FILES
src/VelocityDistortionTest.cpp
)


# Only Linux
if ((NOT WIN32))
# list(APPEND RGL_TEST_FILES
Expand Down
134 changes: 134 additions & 0 deletions test/src/gaussianPoseIndependentTest.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,134 @@
#include <RGLFields.hpp>
#include <helpers/commonHelpers.hpp>
#include <helpers/lidarHelpers.hpp>
#include <helpers/sceneHelpers.hpp>

// The aim of this test is to verify that gaussian noise is not dependent on the ray's pose.
// The test will compare point clouds captured inside cubes.
// Cubes are spawned in different locations. Raytrace will always be performed from the same pose in relation to the cube.
// Gaussian noise is configured with zero standard deviation so there is no randomization.
// The test checks that the noise is always applied correctly in ray coordinates.
// There was a bug in the past where the point cloud transform was not updated correctly in the Gaussian noise nodes.

constexpr float MEAN = 0.1f;
constexpr float EPSILON_DIFF = 1e-4f;

const std::vector<rgl_mat3x4f> lidarRays = makeLidar3dRays(360, 180, 2, 1);
const std::vector<Mat3x4f> cubePoses = {
Mat3x4f::identity(),
Mat3x4f::TRS({100, 200, 300}, {43, 123, 99}),
Mat3x4f::TRS({5, 5, 5}, {45, 45, 45}),
Mat3x4f::TRS({123, 321, 10.5f}, {90, 0, 35.6f}),
Mat3x4f::identity(),
};

struct RaytraceResult
{
Field<XYZ_VEC3_F32>::type xyz;
Field<DISTANCE_F32>::type distance;

static std::vector<rgl_field_t> getFields() { return {XYZ_VEC3_F32, DISTANCE_F32}; }
};

struct GaussianPoseIndependentTest : public RGLTest
{
GaussianPoseIndependentTest()
{
rgl_mat3x4f identity = Mat3x4f::identity().toRGL();

// Setup scene
cube = makeEntity(makeCubeMesh());
EXPECT_RGL_SUCCESS(rgl_entity_set_pose(cube, &identity));

// Create nodes
EXPECT_RGL_SUCCESS(rgl_node_rays_from_mat3x4f(&useRays, lidarRays.data(), lidarRays.size()));
EXPECT_RGL_SUCCESS(rgl_node_rays_transform(&lidarPose, &identity));
EXPECT_RGL_SUCCESS(rgl_node_raytrace(&raytrace, nullptr));
EXPECT_RGL_SUCCESS(rgl_node_points_transform(&toLidarFrame, &identity));
auto formatFields = RaytraceResult::getFields();
EXPECT_RGL_SUCCESS(rgl_node_points_format(&format, formatFields.data(), formatFields.size()));
}

void runTest()
{
bool isFirstRaytrace = true;

std::vector<RaytraceResult> firstRaytraceResult;
std::vector<RaytraceResult> lastRaytraceResult;
firstRaytraceResult.resize(lidarRays.size());
lastRaytraceResult.resize(lidarRays.size());

for (auto&& cubeTf : cubePoses) {
rgl_mat3x4f lidarTf = cubeTf.toRGL();
rgl_mat3x4f lidarTfInv = cubeTf.inverse().toRGL();
ASSERT_RGL_SUCCESS(rgl_entity_set_pose(cube, &lidarTf));
ASSERT_RGL_SUCCESS(rgl_node_rays_transform(&lidarPose, &lidarTf));
ASSERT_RGL_SUCCESS(rgl_node_points_transform(&toLidarFrame, &lidarTfInv));
ASSERT_RGL_SUCCESS(rgl_graph_run(raytrace));

if (isFirstRaytrace) {
ASSERT_RGL_SUCCESS(rgl_graph_get_result_data(format, RGL_FIELD_DYNAMIC_FORMAT, firstRaytraceResult.data()));
isFirstRaytrace = false;
continue;
}

ASSERT_RGL_SUCCESS(rgl_graph_get_result_data(format, RGL_FIELD_DYNAMIC_FORMAT, lastRaytraceResult.data()));
for (int i = 0; i < lidarRays.size(); ++i) {
ASSERT_NEAR(firstRaytraceResult[i].xyz.x(), lastRaytraceResult[i].xyz.x(), EPSILON_DIFF);
ASSERT_NEAR(firstRaytraceResult[i].xyz.y(), lastRaytraceResult[i].xyz.y(), EPSILON_DIFF);
ASSERT_NEAR(firstRaytraceResult[i].xyz.z(), lastRaytraceResult[i].xyz.z(), EPSILON_DIFF);
ASSERT_NEAR(firstRaytraceResult[i].distance, lastRaytraceResult[i].distance, EPSILON_DIFF);
}
}
}

rgl_entity_t cube = nullptr;
rgl_node_t useRays = nullptr;
rgl_node_t lidarPose = nullptr;
rgl_node_t raytrace = nullptr;
rgl_node_t toLidarFrame = nullptr;
rgl_node_t noise = nullptr;
rgl_node_t format = nullptr;
};

TEST_F(GaussianPoseIndependentTest, GaussianNoiseDistance)
{
ASSERT_RGL_SUCCESS(rgl_node_gaussian_noise_distance(&noise, MEAN, 0.0f, 0.0f));

// Connect nodes into graph
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(useRays, lidarPose));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(lidarPose, raytrace));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(raytrace, noise));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(noise, toLidarFrame));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(toLidarFrame, format));

runTest();
}

TEST_F(GaussianPoseIndependentTest, GaussianNoiseAngularHitpoint)
{
ASSERT_RGL_SUCCESS(rgl_node_gaussian_noise_angular_hitpoint(&noise, MEAN, 0.0f, RGL_AXIS_Y));

// Connect nodes into graph
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(useRays, lidarPose));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(lidarPose, raytrace));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(raytrace, noise));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(noise, toLidarFrame));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(toLidarFrame, format));

runTest();
}

TEST_F(GaussianPoseIndependentTest, GaussianNoiseAngularRay)
{
ASSERT_RGL_SUCCESS(rgl_node_gaussian_noise_angular_ray(&noise, MEAN, 0.0f, RGL_AXIS_Y));

// Connect nodes into graph
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(useRays, lidarPose));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(lidarPose, noise));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(noise, raytrace));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(raytrace, toLidarFrame));
ASSERT_RGL_SUCCESS(rgl_graph_node_add_child(toLidarFrame, format));

runTest();
}

0 comments on commit 2e91c81

Please sign in to comment.