Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Feature/reflectivity #341

Open
wants to merge 6 commits into
base: develop
Choose a base branch
from
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
10 changes: 10 additions & 0 deletions include/rgl/api/core.h
Original file line number Diff line number Diff line change
Expand Up @@ -358,6 +358,7 @@ typedef enum : int32_t
* Same as RGL_FIELD_INTENSITY_F32, but uint8_t type.
*/
RGL_FIELD_INTENSITY_U8,
RGL_FIELD_REFLECTIVITY_F32,
RGL_FIELD_IS_HIT_I32,
RGL_FIELD_IS_GROUND_I32,
RGL_FIELD_RAY_IDX_U32,
Expand Down Expand Up @@ -892,6 +893,15 @@ RGL_API rgl_status_t rgl_node_raytrace_configure_beam_divergence(rgl_node_t node
*/
RGL_API rgl_status_t rgl_node_raytrace_configure_default_intensity(rgl_node_t node, float default_intensity);

/**
* Modifies RaytraceNode to set reflectivity alpha.
* Reflectivity alpha is used to calculate reflectivity of the hit point. This value is constant for every hit point.
* Default reflectivity alpha is set to 0.1f.
Copy link
Collaborator

Choose a reason for hiding this comment

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

What's the reason behind this 0.1f choice?

* @param node RaytraceNode to modify.
* @param reflectivity_alpha Reflectivity alpha to set.
*/
RGL_API rgl_status_t rgl_node_raytrace_configure_reflectivity_alpha(rgl_node_t node, float reflectivity_alpha);

/**
* Modifies RaytraceNode to set return mode.
* Point return types (RGL_FIELD_RETURN_TYPE_U8) will be set to corresponding rgl_return_type_t values, e.g. return mode
Expand Down
8 changes: 8 additions & 0 deletions src/RGLFields.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,7 @@ typedef unsigned char TextureTexelFormat;
#define ENTITY_ID_I32 RGL_FIELD_ENTITY_ID_I32
#define INTENSITY_F32 RGL_FIELD_INTENSITY_F32
#define INTENSITY_U8 RGL_FIELD_INTENSITY_U8
#define REFLECTIVITY_F32 RGL_FIELD_REFLECTIVITY_F32
#define LASER_RETRO_F32 RGL_FIELD_LASER_RETRO_F32
#define RING_ID_U16 RGL_FIELD_RING_ID_U16
#define AZIMUTH_F32 RGL_FIELD_AZIMUTH_F32
Expand Down Expand Up @@ -73,6 +74,7 @@ inline const std::set<rgl_field_t>& getAllRealFields()
ENTITY_ID_I32,
INTENSITY_F32,
INTENSITY_U8,
REFLECTIVITY_F32,
LASER_RETRO_F32,
RING_ID_U16,
AZIMUTH_F32,
Expand Down Expand Up @@ -122,6 +124,7 @@ FIELD(RAY_IDX_U32, uint32_t); // PCL uses uint32_t
FIELD(ENTITY_ID_I32, int32_t);
FIELD(INTENSITY_F32, float);
FIELD(INTENSITY_U8, uint8_t);
FIELD(REFLECTIVITY_F32, float);
FIELD(LASER_RETRO_F32, float);
FIELD(IS_HIT_I32, int32_t); // Signed may be faster
FIELD(IS_GROUND_I32, int32_t); // Signed may be faster
Expand Down Expand Up @@ -156,6 +159,7 @@ inline std::size_t getFieldSize(rgl_field_t type)
case IS_GROUND_I32: return Field<IS_GROUND_I32>::size;
case INTENSITY_F32: return Field<INTENSITY_F32>::size;
case INTENSITY_U8: return Field<INTENSITY_U8>::size;
case REFLECTIVITY_F32: return Field<REFLECTIVITY_F32>::size;
case LASER_RETRO_F32: return Field<LASER_RETRO_F32>::size;
case RING_ID_U16: return Field<RING_ID_U16>::size;
case AZIMUTH_F32: return Field<AZIMUTH_F32>::size;
Expand Down Expand Up @@ -215,6 +219,7 @@ inline std::shared_ptr<IAnyArray> createArray(rgl_field_t type, Args&&... args)
case ENTITY_ID_I32: return Subclass<Field<ENTITY_ID_I32>::type>::create(std::forward<Args>(args)...);
case INTENSITY_F32: return Subclass<Field<INTENSITY_F32>::type>::create(std::forward<Args>(args)...);
case INTENSITY_U8: return Subclass<Field<INTENSITY_U8>::type>::create(std::forward<Args>(args)...);
case REFLECTIVITY_F32: return Subclass<Field<REFLECTIVITY_F32>::type>::create(std::forward<Args>(args)...);
case LASER_RETRO_F32: return Subclass<Field<LASER_RETRO_F32>::type>::create(std::forward<Args>(args)...);
case RING_ID_U16: return Subclass<Field<RING_ID_U16>::type>::create(std::forward<Args>(args)...);
case AZIMUTH_F32: return Subclass<Field<AZIMUTH_F32>::type>::create(std::forward<Args>(args)...);
Expand Down Expand Up @@ -251,6 +256,7 @@ inline std::string toString(rgl_field_t type)
case ENTITY_ID_I32: return "ENTITY_ID_I32";
case INTENSITY_F32: return "INTENSITY_F32";
case INTENSITY_U8: return "INTENSITY_U8";
case REFLECTIVITY_F32: return "REFLECTIVITY_F32";
case LASER_RETRO_F32: return "LASER_RETRO_F32";
case RING_ID_U16: return "RING_ID_U16";
case AZIMUTH_F32: return "AZIMUTH_F32";
Expand Down Expand Up @@ -291,6 +297,7 @@ inline std::vector<uint8_t> toRos2Fields(rgl_field_t type)
case ENTITY_ID_I32: return {sensor_msgs::msg::PointField::INT32};
case INTENSITY_F32: return {sensor_msgs::msg::PointField::FLOAT32};
case INTENSITY_U8: return {sensor_msgs::msg::PointField::UINT8};
case REFLECTIVITY_F32: return {sensor_msgs::msg::PointField::FLOAT32};
case LASER_RETRO_F32: return {sensor_msgs::msg::PointField::FLOAT32};
case RING_ID_U16: return {sensor_msgs::msg::PointField::UINT16};
case AZIMUTH_F32: return {sensor_msgs::msg::PointField::FLOAT32};
Expand Down Expand Up @@ -342,6 +349,7 @@ inline std::vector<std::string> toRos2Names(rgl_field_t type)
case INTENSITY_F32:
case INTENSITY_U8:
return {"intensity"};
case REFLECTIVITY_F32: return {"reflectivity"};
case LASER_RETRO_F32: return {"laser_retro"};
case RING_ID_U16: return {"channel"};
case AZIMUTH_F32: return {"azimuth"};
Expand Down
21 changes: 21 additions & 0 deletions src/api/apiCore.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1066,6 +1066,27 @@ void TapeCore::tape_node_raytrace_configure_default_intensity(const YAML::Node&
rgl_node_raytrace_configure_default_intensity(node, yamlNode[1].as<float>());
}

RGL_API rgl_status_t rgl_node_raytrace_configure_reflectivity_alpha(rgl_node_t node, float reflectivity_alpha)
{
auto status = rglSafeCall([&]() {
RGL_API_LOG("rgl_node_raytrace_configure_reflectivity_alpha(node={}, default_intensity={})", repr(node),
reflectivity_alpha);
CHECK_ARG(node != nullptr);
CHECK_ARG(reflectivity_alpha >= 0.0f);
RaytraceNode::Ptr raytraceNode = Node::validatePtr<RaytraceNode>(node);
raytraceNode->setReflectivityAlpha(reflectivity_alpha);
});
TAPE_HOOK(node, reflectivity_alpha);
return status;
}

void TapeCore::tape_node_raytrace_configure_reflectivity_alpha(const YAML::Node& yamlNode, PlaybackState& state)
{
auto nodeId = yamlNode[0].as<TapeAPIObjectID>();
rgl_node_t node = state.nodes.at(nodeId);
rgl_node_raytrace_configure_reflectivity_alpha(node, yamlNode[1].as<float>());
}

RGL_API rgl_status_t rgl_node_raytrace_configure_return_mode(rgl_node_t node, rgl_return_mode_t return_mode)
{
auto status = rglSafeCall([&]() {
Expand Down
1 change: 1 addition & 0 deletions src/gpu/MultiReturn.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@ struct MultiReturnSamplesPointers
Field<IS_HIT_I32>::type* isHit;
Field<DISTANCE_F32>::type* distance;
Field<INTENSITY_F32>::type* intensity;
Field<REFLECTIVITY_F32>::type* reflectivity;
Copy link
Collaborator

Choose a reason for hiding this comment

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

Remove

Field<LASER_RETRO_F32>::type* laserRetro;
Field<ENTITY_ID_I32>::type* entityId;
Field<ABSOLUTE_VELOCITY_VEC3_F32>::type* absVelocity;
Expand Down
2 changes: 2 additions & 0 deletions src/gpu/RaytraceRequestContext.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -29,6 +29,7 @@ struct RaytraceRequestContext
float farNonHitDistance;

float defaultIntensity;
float reflectivityAlpha;

const Mat3x4f* raysWorld;
size_t rayCount;
Expand Down Expand Up @@ -58,6 +59,7 @@ struct RaytraceRequestContext
Field<DISTANCE_F32>::type* distance;
Field<INTENSITY_F32>::type* intensityF32;
Field<INTENSITY_U8>::type* intensityU8;
Field<REFLECTIVITY_F32>::type* reflectivityF32;
Field<LASER_RETRO_F32>::type* laserRetro;
Field<TIME_STAMP_F64>::type* timestampF64;
Field<TIME_STAMP_U32>::type* timestampU32;
Expand Down
6 changes: 6 additions & 0 deletions src/gpu/nodeKernels.cu
Original file line number Diff line number Diff line change
Expand Up @@ -127,6 +127,9 @@ __device__ void saveReturnAsHit(const RaytraceRequestContext* ctx, int beamIdx,
static_cast<uint8_t>(std::round(ctx->mrSamples.intensity[sampleIdx])) :
UINT8_MAX;
}
if (ctx->reflectivityF32 != nullptr) {
ctx->reflectivityF32[returnPointIdx] = ctx->mrSamples.reflectivity[sampleIdx];
}
Comment on lines +130 to +132
Copy link
Collaborator

Choose a reason for hiding this comment

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

Here you can use selected sample to compute reflectivity. This way reflectivity can be removed from multi-return samples to save memory.
This suggestion was NOT COMPILED.

Suggested change
if (ctx->reflectivityF32 != nullptr) {
ctx->reflectivityF32[returnPointIdx] = ctx->mrSamples.reflectivity[sampleIdx];
}
if (ctx->reflectivityF32 != nullptr) {
auto distance2 = ctx->mrSamples.intensity[sampleIdx] * ctx->mrSamples.intensity[sampleIdx];
ctx->reflectivityF32[returnPointIdx] = reflectivityAlpha * ctx->mrSamples.intensity[sampleIdx] * distance2;
}

if (ctx->entityId != nullptr) {
ctx->entityId[returnPointIdx] = ctx->mrSamples.entityId[sampleIdx];
}
Expand Down Expand Up @@ -183,6 +186,9 @@ __device__ void saveReturnAsNonHit(const RaytraceRequestContext* ctx, int firstS
if (ctx->intensityU8 != nullptr) {
ctx->intensityU8[returnPointIdx] = 0;
}
if (ctx->reflectivityF32 != nullptr) {
ctx->reflectivityF32[returnPointIdx] = 0;
}
if (ctx->entityId != nullptr) {
ctx->entityId[returnPointIdx] = RGL_ENTITY_INVALID_ID;
}
Expand Down
10 changes: 7 additions & 3 deletions src/gpu/optixPrograms.cu
Original file line number Diff line number Diff line change
Expand Up @@ -34,7 +34,7 @@ extern "C" static __constant__ RaytraceRequestContext ctx;

// Helper functions
__device__ void saveSampleAsNonHit(int sampleIdx, float nonHitDistance);
__device__ void saveSampleAsHit(int sampleIdx, float distance, float intensity, float laserRetro, int objectID,
__device__ void saveSampleAsHit(int sampleIdx, float distance, float intensity, float reflectivity, float laserRetro, int objectID,
const Vec3f& absVelocity, const Vec3f& relVelocity, float radialSpeed, const Vec3f& normal,
float incidentAngle);
__device__ void saveNonHitBeamSamples(int beamIdx, float nonHitDistance);
Expand Down Expand Up @@ -166,6 +166,9 @@ extern "C" __global__ void __closesthit__()
}
intensity *= cosIncidentAngle;

float reflectivityAlpha = ctx.reflectivityAlpha;
float reflectivity = reflectivityAlpha * intensity * distance * distance;

Vec3f absPointVelocity{NAN};
Vec3f relPointVelocity{NAN};
float radialSpeed{NAN};
Expand Down Expand Up @@ -211,7 +214,7 @@ extern "C" __global__ void __closesthit__()
radialSpeed = hitRays.normalized().dot(relPointVelocity);
}

saveSampleAsHit(mrSampleIdx, distance, intensity, laserRetro, entityId, absPointVelocity, relPointVelocity, radialSpeed,
saveSampleAsHit(mrSampleIdx, distance, intensity, reflectivity, laserRetro, entityId, absPointVelocity, relPointVelocity, radialSpeed,
wNormal, incidentAngle);
}

Expand Down Expand Up @@ -253,13 +256,14 @@ __device__ void saveSampleAsNonHit(int sampleIdx, float nonHitDistance)
ctx.mrSamples.distance[sampleIdx] = nonHitDistance;
}

__device__ void saveSampleAsHit(int sampleIdx, float distance, float intensity, float laserRetro, int objectID,
__device__ void saveSampleAsHit(int sampleIdx, float distance, float intensity,float reflectivity, float laserRetro, int objectID,
const Vec3f& absVelocity, const Vec3f& relVelocity, float radialSpeed, const Vec3f& normal,
float incidentAngle)
{
ctx.mrSamples.isHit[sampleIdx] = true;
ctx.mrSamples.distance[sampleIdx] = distance;
ctx.mrSamples.intensity[sampleIdx] = intensity;
ctx.mrSamples.reflectivity[sampleIdx] = reflectivity;

if (ctx.mrSamples.laserRetro != nullptr) {
ctx.mrSamples.laserRetro[sampleIdx] = laserRetro;
Expand Down
8 changes: 7 additions & 1 deletion src/graph/NodesCore.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -139,6 +139,7 @@ struct RaytraceNode : IPointsNode
void setNonHitDistanceValues(float nearDistance, float farDistance);
void setNonHitsMask(const int8_t* maskRaw, size_t maskPointCount);
void setDefaultIntensity(float intensity) { defaultIntensity = intensity; }
void setReflectivityAlpha(float reflectivity_alpha) { reflectivityAlpha = reflectivity_alpha; }
void setReturnMode(rgl_return_mode_t mode)
{
if (mode == RGL_RETURN_UNKNOWN) {
Expand All @@ -158,7 +159,8 @@ struct RaytraceNode : IPointsNode
explicit MultiReturnSamples(StreamBoundObjectsManager& arrayMgr)
: isHit(DeviceAsyncArray<Field<IS_HIT_I32>::type>::create(arrayMgr)),
distance(DeviceAsyncArray<Field<DISTANCE_F32>::type>::create(arrayMgr)),
intensity(DeviceAsyncArray<Field<INTENSITY_F32>::type>::create(arrayMgr))
intensity(DeviceAsyncArray<Field<INTENSITY_F32>::type>::create(arrayMgr)),
reflectivity(DeviceAsyncArray<Field<REFLECTIVITY_F32>::type>::create(arrayMgr))
{}

void adjustToFields(const std::unordered_map<rgl_field_t, IAnyArray::Ptr>& inFieldData,
Expand All @@ -183,6 +185,7 @@ struct RaytraceNode : IPointsNode
isHit->resize(size, false, false);
distance->resize(size, false, false);
intensity->resize(size, false, false);
reflectivity->resize(size, false, false);
resizeField(laserRetro, size);
resizeField(entityId, size);
resizeField(absVelocity, size);
Expand All @@ -200,6 +203,7 @@ struct RaytraceNode : IPointsNode
.isHit = isHit->getWritePtr(),
.distance = distance->getWritePtr(),
.intensity = intensity->getWritePtr(),
.reflectivity = reflectivity->getWritePtr(),
.laserRetro = laserRetro ? laserRetro->getWritePtr() : nullptr,
.entityId = entityId ? entityId->getWritePtr() : nullptr,
.absVelocity = absVelocity ? absVelocity->getWritePtr() : nullptr,
Expand All @@ -223,6 +227,7 @@ struct RaytraceNode : IPointsNode
DeviceAsyncArray<Field<IS_HIT_I32>::type>::Ptr isHit;
DeviceAsyncArray<Field<DISTANCE_F32>::type>::Ptr distance;
DeviceAsyncArray<Field<INTENSITY_F32>::type>::Ptr intensity;
DeviceAsyncArray<Field<REFLECTIVITY_F32>::type>::Ptr reflectivity;

// Additional field data for multi-return samples.
DeviceAsyncArray<Field<LASER_RETRO_F32>::type>::Ptr laserRetro;
Expand All @@ -244,6 +249,7 @@ struct RaytraceNode : IPointsNode
float nearNonHitDistance{std::numeric_limits<float>::infinity()};
float farNonHitDistance{std::numeric_limits<float>::infinity()};
float defaultIntensity = 0.0f;
float reflectivityAlpha = 0.1f;

MultiReturnSamples mrSampleData = MultiReturnSamples{arrayMgr};
rgl_return_mode_t returnMode = RGL_RETURN_FIRST;
Expand Down
2 changes: 2 additions & 0 deletions src/graph/RaytraceNode.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -98,6 +98,7 @@ void RaytraceNode::enqueueExecImpl()
.nearNonHitDistance = nearNonHitDistance,
.farNonHitDistance = farNonHitDistance,
.defaultIntensity = defaultIntensity,
.reflectivityAlpha = reflectivityAlpha,
.raysWorld = raysPtr,
.rayCount = raysNode->getRayCount(),
.rayOriginToWorld = raysNode->getCumulativeRayTransfrom(),
Expand All @@ -119,6 +120,7 @@ void RaytraceNode::enqueueExecImpl()
.distance = getPtrTo<DISTANCE_F32>(),
.intensityF32 = getPtrTo<INTENSITY_F32>(),
.intensityU8 = getPtrTo<INTENSITY_U8>(),
.reflectivityF32 = getPtrTo<REFLECTIVITY_F32>(),
.laserRetro = getPtrTo<LASER_RETRO_F32>(),
.timestampF64 = getPtrTo<TIME_STAMP_F64>(),
.timestampU32 = getPtrTo<TIME_STAMP_U32>(),
Expand Down
3 changes: 3 additions & 0 deletions src/tape/TapeCore.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -60,6 +60,7 @@ class TapeCore
static void tape_node_raytrace_configure_mask(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_raytrace_configure_beam_divergence(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_raytrace_configure_default_intensity(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_raytrace_configure_reflectivity_alpha(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_raytrace_configure_return_mode(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_points_format(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_points_yield(const YAML::Node& yamlNode, PlaybackState& state);
Expand Down Expand Up @@ -122,6 +123,8 @@ class TapeCore
TapeCore::tape_node_raytrace_configure_beam_divergence),
TAPE_CALL_MAPPING("rgl_node_raytrace_configure_default_intensity",
TapeCore::tape_node_raytrace_configure_default_intensity),
TAPE_CALL_MAPPING("rgl_node_raytrace_configure_reflectivity_alpha",
TapeCore::tape_node_raytrace_configure_reflectivity_alpha),
TAPE_CALL_MAPPING("rgl_node_raytrace_configure_return_mode",
TapeCore::tape_node_raytrace_configure_return_mode),
TAPE_CALL_MAPPING("rgl_node_points_format", TapeCore::tape_node_points_format),
Expand Down
1 change: 1 addition & 0 deletions test/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -48,6 +48,7 @@ set(RGL_TEST_FILES
src/scene/entityVelocityTest.cpp
src/scene/meshAPITest.cpp
src/scene/textureTest.cpp
src/scene/reflectivityTest.cpp
src/synchronization/graphAndCopyStream.cpp
src/synchronization/graphThreadSynchronization.cpp
src/synchronization/testKernel.cu
Expand Down
3 changes: 3 additions & 0 deletions test/include/helpers/fieldGenerators.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,9 @@ static std::function<Field<INTENSITY_F32>::type(int)> genIntensityF32 = [](int i
static std::function<Field<INTENSITY_U8>::type(int)> genIntensityU8 = [](int i) {
return i % std::numeric_limits<Field<INTENSITY_U8>::type>::max();
};
static std::function<Field<REFLECTIVITY_F32>::type(int)> genReflectivityF32 = [](int i) {
return static_cast<float>(i) / (static_cast<float>(i + 1));
};
static std::function<Field<LASER_RETRO_F32>::type(int)> genLaserRetro = [](int i) {
return static_cast<float>(i) / (static_cast<float>(i + 1));
};
Expand Down
1 change: 1 addition & 0 deletions test/include/helpers/testPointCloud.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -289,6 +289,7 @@ class TestPointCloud
{ENTITY_ID_I32, [&](std::size_t count) {setFieldValues<ENTITY_ID_I32>(generateFieldValues(count, genEntityId));}},
{INTENSITY_F32, [&](std::size_t count) {setFieldValues<INTENSITY_F32>(generateFieldValues(count, genIntensityF32));}},
{INTENSITY_U8, [&](std::size_t count) {setFieldValues<INTENSITY_U8>(generateFieldValues(count, genIntensityU8));}},
{REFLECTIVITY_F32, [&](std::size_t count) {setFieldValues<REFLECTIVITY_F32>(generateFieldValues(count, genReflectivityF32));}},
{LASER_RETRO_F32, [&](std::size_t count) {setFieldValues<LASER_RETRO_F32>(generateFieldValues(count, genLaserRetro));}},
{RING_ID_U16, [&](std::size_t count) {setFieldValues<RING_ID_U16>(generateFieldValues(count, genRingId));}},
{AZIMUTH_F32, [&](std::size_t count) {setFieldValues<AZIMUTH_F32>(generateFieldValues(count, genAzimuth));}},
Expand Down
3 changes: 3 additions & 0 deletions test/src/TapeTest.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -252,6 +252,9 @@ TEST_F(TapeTest, RecordPlayAllCalls)
float defaultIntensity = 1.1f;
EXPECT_RGL_SUCCESS(rgl_node_raytrace_configure_default_intensity(raytrace, defaultIntensity));

float reflectivityAlpha = 0.1;
EXPECT_RGL_SUCCESS(rgl_node_raytrace_configure_reflectivity_alpha(raytrace, reflectivityAlpha));

rgl_return_mode_t returnMode = RGL_RETURN_FIRST;
EXPECT_RGL_SUCCESS(rgl_node_raytrace_configure_return_mode(raytrace, returnMode));

Expand Down
Loading
Loading