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

Merged
merged 10 commits into from
Dec 16, 2024
Merged
Changes from 1 commit
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
Prev Previous commit
Next Next commit
First working version of reflectivity calculation with hardcoded alpha.
PiotrMrozik committed Dec 4, 2024
commit afa83a98f3e1e37ede84c845d7816eb9943b8a6f
1 change: 1 addition & 0 deletions src/gpu/MultiReturn.hpp
Original file line number Diff line number Diff line change
@@ -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;
Field<LASER_RETRO_F32>::type* laserRetro;
Field<ENTITY_ID_I32>::type* entityId;
Field<ABSOLUTE_VELOCITY_VEC3_F32>::type* absVelocity;
1 change: 1 addition & 0 deletions src/gpu/RaytraceRequestContext.hpp
Original file line number Diff line number Diff line change
@@ -58,6 +58,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;
6 changes: 6 additions & 0 deletions src/gpu/nodeKernels.cu
Original file line number Diff line number Diff line change
@@ -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];
}
if (ctx->entityId != nullptr) {
ctx->entityId[returnPointIdx] = ctx->mrSamples.entityId[sampleIdx];
}
@@ -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;
}
10 changes: 7 additions & 3 deletions src/gpu/optixPrograms.cu
Original file line number Diff line number Diff line change
@@ -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);
@@ -166,6 +166,9 @@ extern "C" __global__ void __closesthit__()
}
intensity *= cosIncidentAngle;

float reflectivityAlpha = 0.1f;
float reflectivity = reflectivityAlpha * intensity * distance * distance;

Vec3f absPointVelocity{NAN};
Vec3f relPointVelocity{NAN};
float radialSpeed{NAN};
@@ -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);
}

@@ -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;
6 changes: 5 additions & 1 deletion src/graph/NodesCore.hpp
Original file line number Diff line number Diff line change
@@ -158,7 +158,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,
@@ -183,6 +184,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);
@@ -200,6 +202,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,
@@ -223,6 +226,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;
1 change: 1 addition & 0 deletions src/graph/RaytraceNode.cpp
Original file line number Diff line number Diff line change
@@ -119,6 +119,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>(),
3 changes: 3 additions & 0 deletions test/include/helpers/fieldGenerators.hpp
Original file line number Diff line number Diff line change
@@ -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));
};
1 change: 1 addition & 0 deletions test/include/helpers/testPointCloud.hpp
Original file line number Diff line number Diff line change
@@ -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));}},
20 changes: 13 additions & 7 deletions test/src/helpers/pointsTest.cpp
Original file line number Diff line number Diff line change
@@ -5,18 +5,20 @@
class PointCloudTest : public ::testing::TestWithParam<int>
{
protected:
std::vector<rgl_field_t> allNotDummyFields = {XYZ_VEC3_F32, IS_HIT_I32, RAY_IDX_U32, ENTITY_ID_I32, INTENSITY_F32,
RING_ID_U16, AZIMUTH_F32, DISTANCE_F32, RETURN_TYPE_U8, TIME_STAMP_F64};
std::vector<rgl_field_t> allNotDummyFields = {XYZ_VEC3_F32, IS_HIT_I32, RAY_IDX_U32, ENTITY_ID_I32,
INTENSITY_F32, REFLECTIVITY_F32, RING_ID_U16, AZIMUTH_F32,
DISTANCE_F32, RETURN_TYPE_U8, TIME_STAMP_F64};

std::vector<rgl_field_t> fieldsWithPaddings = {PADDING_32, XYZ_VEC3_F32, PADDING_16, IS_HIT_I32, PADDING_8,
RAY_IDX_U32, ENTITY_ID_I32, INTENSITY_F32, RING_ID_U16, AZIMUTH_F32,
DISTANCE_F32, RETURN_TYPE_U8, TIME_STAMP_F64, PADDING_16, PADDING_32};
std::vector<rgl_field_t> fieldsWithPaddings = {
PADDING_32, XYZ_VEC3_F32, PADDING_16, IS_HIT_I32, PADDING_8, RAY_IDX_U32, ENTITY_ID_I32, INTENSITY_F32,
REFLECTIVITY_F32, RING_ID_U16, AZIMUTH_F32, DISTANCE_F32, RETURN_TYPE_U8, TIME_STAMP_F64, PADDING_16, PADDING_32};

std::vector<Field<XYZ_VEC3_F32>::type> pointCoord;
std::vector<Field<IS_HIT_I32>::type> isHit;
std::vector<Field<RAY_IDX_U32>::type> rayIdx;
std::vector<Field<ENTITY_ID_I32>::type> entityId;
std::vector<Field<INTENSITY_F32>::type> intensity;
std::vector<Field<REFLECTIVITY_F32>::type> reflectivity;
std::vector<Field<RING_ID_U16>::type> ringId;
std::vector<Field<AZIMUTH_F32>::type> azimuth;
std::vector<Field<DISTANCE_F32>::type> distance;
@@ -31,6 +33,7 @@ class PointCloudTest : public ::testing::TestWithParam<int>
Field<RAY_IDX_U32>::type rayIdx;
Field<ENTITY_ID_I32>::type entityId;
Field<INTENSITY_F32>::type intensity;
Field<REFLECTIVITY_F32>::type reflectivity;
Field<RING_ID_U16>::type ringId;
Field<AZIMUTH_F32>::type azimuth;
Field<DISTANCE_F32>::type distance;
@@ -45,8 +48,8 @@ class PointCloudTest : public ::testing::TestWithParam<int>
points.reserve(pointsCount);
for (int i = 0; i < pointsCount; i++) {
points.emplace_back(TestPointStruct{genCoord(i), genHalfHit(i), genRayIdx(i), genEntityId(i), genIntensityF32(i),
genRingId(i), genAzimuth(i), genDistance(i), genReturnType(i),
genTimeStampF64(i)});
genReflectivityF32(i), genRingId(i), genAzimuth(i), genDistance(i),
genReturnType(i), genTimeStampF64(i)});
}
return points;
}
@@ -58,6 +61,7 @@ class PointCloudTest : public ::testing::TestWithParam<int>
rayIdx = generateFieldValues(pointsCount, genRayIdx);
entityId = generateFieldValues(pointsCount, genEntityId);
intensity = generateFieldValues(pointsCount, genIntensityF32);
reflectivity = generateFieldValues(pointsCount, genReflectivityF32);
ringId = generateFieldValues(pointsCount, genRingId);
azimuth = generateFieldValues(pointsCount, genAzimuth);
distance = generateFieldValues(pointsCount, genDistance);
@@ -72,6 +76,7 @@ class PointCloudTest : public ::testing::TestWithParam<int>
pointCloud.setFieldValues<RAY_IDX_U32>(rayIdx);
pointCloud.setFieldValues<ENTITY_ID_I32>(entityId);
pointCloud.setFieldValues<INTENSITY_F32>(intensity);
pointCloud.setFieldValues<REFLECTIVITY_F32>(reflectivity);
pointCloud.setFieldValues<RING_ID_U16>(ringId);
pointCloud.setFieldValues<AZIMUTH_F32>(azimuth);
pointCloud.setFieldValues<DISTANCE_F32>(distance);
@@ -92,6 +97,7 @@ class PointCloudTest : public ::testing::TestWithParam<int>
EXPECT_EQ(pointCloud.getFieldValues<RAY_IDX_U32>(), rayIdx);
EXPECT_EQ(pointCloud.getFieldValues<ENTITY_ID_I32>(), entityId);
EXPECT_EQ(pointCloud.getFieldValues<INTENSITY_F32>(), intensity);
EXPECT_EQ(pointCloud.getFieldValues<REFLECTIVITY_F32>(), reflectivity);
EXPECT_EQ(pointCloud.getFieldValues<RING_ID_U16>(), ringId);
EXPECT_EQ(pointCloud.getFieldValues<AZIMUTH_F32>(), azimuth);
EXPECT_EQ(pointCloud.getFieldValues<DISTANCE_F32>(), distance);
67 changes: 66 additions & 1 deletion test/src/scene/textureTest.cpp
Original file line number Diff line number Diff line change
@@ -59,7 +59,7 @@ TEST_P(TextureTest, rgl_texture_reading)
std::vector<rgl_mat3x4f> rays = {// Ray must be incident perpendicular to the surface to receive all intensity
Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}).toRGL()};

std::vector<rgl_field_t> yieldFields = {INTENSITY_F32};
std::vector<rgl_field_t> yieldFields = {INTENSITY_F32, REFLECTIVITY_F32};

EXPECT_RGL_SUCCESS(rgl_node_rays_from_mat3x4f(&useRaysNode, rays.data(), rays.size()));
EXPECT_RGL_SUCCESS(rgl_node_raytrace(&raytraceNode, nullptr));
@@ -87,6 +87,71 @@ TEST_P(TextureTest, rgl_texture_reading)
}
}

TEST_P(TextureTest, rgl_reflectivity)
{
auto [width, height, value] = GetParam();

rgl_texture_t texture = nullptr;
rgl_entity_t entity = nullptr;
rgl_mesh_t mesh = makeCubeMesh();
auto textureRawData = generateStaticColorTexture<TextureTexelFormat>(width, height, value);

EXPECT_RGL_SUCCESS(rgl_texture_create(&texture, textureRawData.data(), width, height));
EXPECT_RGL_SUCCESS(rgl_mesh_set_texture_coords(mesh, cubeUVs, ARRAY_SIZE(cubeUVs)));

EXPECT_RGL_SUCCESS(rgl_entity_create(&entity, nullptr, mesh));
EXPECT_RGL_SUCCESS(rgl_entity_set_intensity_texture(entity, texture));

// Create RGL graph pipeline.
rgl_node_t useRaysNode = nullptr, raytraceNode = nullptr, compactNode = nullptr, yieldNode = nullptr;

std::vector<rgl_mat3x4f> rays = {// Ray must be incident perpendicular to the surface to receive all intensity
Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}).toRGL()};

std::vector<rgl_field_t> yieldFields = {INTENSITY_F32, REFLECTIVITY_F32, DISTANCE_F32};

EXPECT_RGL_SUCCESS(rgl_node_rays_from_mat3x4f(&useRaysNode, rays.data(), rays.size()));
EXPECT_RGL_SUCCESS(rgl_node_raytrace(&raytraceNode, nullptr));
EXPECT_RGL_SUCCESS(rgl_node_points_compact_by_field(&compactNode, RGL_FIELD_IS_HIT_I32));
EXPECT_RGL_SUCCESS(rgl_node_points_yield(&yieldNode, yieldFields.data(), yieldFields.size()));

EXPECT_RGL_SUCCESS(rgl_graph_node_add_child(useRaysNode, raytraceNode));
EXPECT_RGL_SUCCESS(rgl_graph_node_add_child(raytraceNode, compactNode));
EXPECT_RGL_SUCCESS(rgl_graph_node_add_child(compactNode, yieldNode));

EXPECT_RGL_SUCCESS(rgl_graph_run(raytraceNode));

std::vector<::Field<INTENSITY_F32>::type> outIntensity;
std::vector<::Field<REFLECTIVITY_F32>::type> outReflectivity;
std::vector<::Field<DISTANCE_F32>::type> outDistance;

int32_t outCount, outSizeOf;
EXPECT_RGL_SUCCESS(rgl_graph_get_result_size(yieldNode, INTENSITY_F32, &outCount, &outSizeOf));
EXPECT_EQ(outSizeOf, getFieldSize(INTENSITY_F32));

EXPECT_RGL_SUCCESS(rgl_graph_get_result_size(yieldNode, REFLECTIVITY_F32, &outCount, &outSizeOf));
EXPECT_EQ(outSizeOf, getFieldSize(REFLECTIVITY_F32));

EXPECT_RGL_SUCCESS(rgl_graph_get_result_size(yieldNode, DISTANCE_F32, &outCount, &outSizeOf));
EXPECT_EQ(outSizeOf, getFieldSize(DISTANCE_F32));

outIntensity.resize(outCount);
outReflectivity.resize(outCount);
outDistance.resize(outCount);

EXPECT_RGL_SUCCESS(rgl_graph_get_result_data(yieldNode, INTENSITY_F32, outIntensity.data()));
EXPECT_RGL_SUCCESS(rgl_graph_get_result_data(yieldNode, REFLECTIVITY_F32, outReflectivity.data()));
EXPECT_RGL_SUCCESS(rgl_graph_get_result_data(yieldNode, DISTANCE_F32, outDistance.data()));

for (int i = 0; i < outCount; ++i) {
EXPECT_NEAR(((float) value), outIntensity.at(i), EPSILON_F);
float distance = outDistance.at(i);
float intensity = outIntensity.at(i);
float reflectivityValue = 0.1f * distance * distance * intensity;
EXPECT_NEAR(reflectivityValue, outReflectivity.at(i), EPSILON_F);
}
}

// Use-case test. Create texture, mesh and entity. Set texture to entity and run graph pipeline.
// As a result, we should get the cube with assigned gradient-checkerboard texture.
TEST_P(TextureTest, rgl_texture_use_case)