From 0fc0a0bed02e208692f3345aedf3b9ffed67db53 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Wed, 4 Dec 2024 14:58:38 +0100 Subject: [PATCH 01/10] Introduce new field --- include/rgl/api/core.h | 1 + src/RGLFields.hpp | 8 ++++++++ 2 files changed, 9 insertions(+) diff --git a/include/rgl/api/core.h b/include/rgl/api/core.h index 1b13f83a..2697c4df 100644 --- a/include/rgl/api/core.h +++ b/include/rgl/api/core.h @@ -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, diff --git a/src/RGLFields.hpp b/src/RGLFields.hpp index d12d8098..f979bcdb 100644 --- a/src/RGLFields.hpp +++ b/src/RGLFields.hpp @@ -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 @@ -73,6 +74,7 @@ inline const std::set& getAllRealFields() ENTITY_ID_I32, INTENSITY_F32, INTENSITY_U8, + REFLECTIVITY_F32, LASER_RETRO_F32, RING_ID_U16, AZIMUTH_F32, @@ -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 @@ -156,6 +159,7 @@ inline std::size_t getFieldSize(rgl_field_t type) case IS_GROUND_I32: return Field::size; case INTENSITY_F32: return Field::size; case INTENSITY_U8: return Field::size; + case REFLECTIVITY_F32: return Field::size; case LASER_RETRO_F32: return Field::size; case RING_ID_U16: return Field::size; case AZIMUTH_F32: return Field::size; @@ -215,6 +219,7 @@ inline std::shared_ptr createArray(rgl_field_t type, Args&&... args) case ENTITY_ID_I32: return Subclass::type>::create(std::forward(args)...); case INTENSITY_F32: return Subclass::type>::create(std::forward(args)...); case INTENSITY_U8: return Subclass::type>::create(std::forward(args)...); + case REFLECTIVITY_F32: return Subclass::type>::create(std::forward(args)...); case LASER_RETRO_F32: return Subclass::type>::create(std::forward(args)...); case RING_ID_U16: return Subclass::type>::create(std::forward(args)...); case AZIMUTH_F32: return Subclass::type>::create(std::forward(args)...); @@ -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"; @@ -291,6 +297,7 @@ inline std::vector 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}; @@ -342,6 +349,7 @@ inline std::vector 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"}; From afa83a98f3e1e37ede84c845d7816eb9943b8a6f Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Wed, 4 Dec 2024 17:37:17 +0100 Subject: [PATCH 02/10] First working version of reflectivity calculation with hardcoded alpha. --- src/gpu/MultiReturn.hpp | 1 + src/gpu/RaytraceRequestContext.hpp | 1 + src/gpu/nodeKernels.cu | 6 +++ src/gpu/optixPrograms.cu | 10 ++-- src/graph/NodesCore.hpp | 6 ++- src/graph/RaytraceNode.cpp | 1 + test/include/helpers/fieldGenerators.hpp | 3 ++ test/include/helpers/testPointCloud.hpp | 1 + test/src/helpers/pointsTest.cpp | 20 ++++--- test/src/scene/textureTest.cpp | 67 +++++++++++++++++++++++- 10 files changed, 104 insertions(+), 12 deletions(-) diff --git a/src/gpu/MultiReturn.hpp b/src/gpu/MultiReturn.hpp index 028f303a..c924826a 100644 --- a/src/gpu/MultiReturn.hpp +++ b/src/gpu/MultiReturn.hpp @@ -27,6 +27,7 @@ struct MultiReturnSamplesPointers Field::type* isHit; Field::type* distance; Field::type* intensity; + Field::type* reflectivity; Field::type* laserRetro; Field::type* entityId; Field::type* absVelocity; diff --git a/src/gpu/RaytraceRequestContext.hpp b/src/gpu/RaytraceRequestContext.hpp index 0e32a667..29c28e19 100644 --- a/src/gpu/RaytraceRequestContext.hpp +++ b/src/gpu/RaytraceRequestContext.hpp @@ -58,6 +58,7 @@ struct RaytraceRequestContext Field::type* distance; Field::type* intensityF32; Field::type* intensityU8; + Field::type* reflectivityF32; Field::type* laserRetro; Field::type* timestampF64; Field::type* timestampU32; diff --git a/src/gpu/nodeKernels.cu b/src/gpu/nodeKernels.cu index 3917397c..bb4e78e8 100644 --- a/src/gpu/nodeKernels.cu +++ b/src/gpu/nodeKernels.cu @@ -127,6 +127,9 @@ __device__ void saveReturnAsHit(const RaytraceRequestContext* ctx, int beamIdx, static_cast(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; } diff --git a/src/gpu/optixPrograms.cu b/src/gpu/optixPrograms.cu index 16b3cc95..4a9cab5e 100644 --- a/src/gpu/optixPrograms.cu +++ b/src/gpu/optixPrograms.cu @@ -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; diff --git a/src/graph/NodesCore.hpp b/src/graph/NodesCore.hpp index fca18c9e..168c6812 100644 --- a/src/graph/NodesCore.hpp +++ b/src/graph/NodesCore.hpp @@ -158,7 +158,8 @@ struct RaytraceNode : IPointsNode explicit MultiReturnSamples(StreamBoundObjectsManager& arrayMgr) : isHit(DeviceAsyncArray::type>::create(arrayMgr)), distance(DeviceAsyncArray::type>::create(arrayMgr)), - intensity(DeviceAsyncArray::type>::create(arrayMgr)) + intensity(DeviceAsyncArray::type>::create(arrayMgr)), + reflectivity(DeviceAsyncArray::type>::create(arrayMgr)) {} void adjustToFields(const std::unordered_map& 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::type>::Ptr isHit; DeviceAsyncArray::type>::Ptr distance; DeviceAsyncArray::type>::Ptr intensity; + DeviceAsyncArray::type>::Ptr reflectivity; // Additional field data for multi-return samples. DeviceAsyncArray::type>::Ptr laserRetro; diff --git a/src/graph/RaytraceNode.cpp b/src/graph/RaytraceNode.cpp index 9425aa1b..66a90f61 100644 --- a/src/graph/RaytraceNode.cpp +++ b/src/graph/RaytraceNode.cpp @@ -119,6 +119,7 @@ void RaytraceNode::enqueueExecImpl() .distance = getPtrTo(), .intensityF32 = getPtrTo(), .intensityU8 = getPtrTo(), + .reflectivityF32 = getPtrTo(), .laserRetro = getPtrTo(), .timestampF64 = getPtrTo(), .timestampU32 = getPtrTo(), diff --git a/test/include/helpers/fieldGenerators.hpp b/test/include/helpers/fieldGenerators.hpp index 8de8a174..d9415d54 100644 --- a/test/include/helpers/fieldGenerators.hpp +++ b/test/include/helpers/fieldGenerators.hpp @@ -36,6 +36,9 @@ static std::function::type(int)> genIntensityF32 = [](int i static std::function::type(int)> genIntensityU8 = [](int i) { return i % std::numeric_limits::type>::max(); }; +static std::function::type(int)> genReflectivityF32 = [](int i) { + return static_cast(i) / (static_cast(i + 1)); +}; static std::function::type(int)> genLaserRetro = [](int i) { return static_cast(i) / (static_cast(i + 1)); }; diff --git a/test/include/helpers/testPointCloud.hpp b/test/include/helpers/testPointCloud.hpp index 19420033..1127e358 100644 --- a/test/include/helpers/testPointCloud.hpp +++ b/test/include/helpers/testPointCloud.hpp @@ -289,6 +289,7 @@ class TestPointCloud {ENTITY_ID_I32, [&](std::size_t count) {setFieldValues(generateFieldValues(count, genEntityId));}}, {INTENSITY_F32, [&](std::size_t count) {setFieldValues(generateFieldValues(count, genIntensityF32));}}, {INTENSITY_U8, [&](std::size_t count) {setFieldValues(generateFieldValues(count, genIntensityU8));}}, + {REFLECTIVITY_F32, [&](std::size_t count) {setFieldValues(generateFieldValues(count, genReflectivityF32));}}, {LASER_RETRO_F32, [&](std::size_t count) {setFieldValues(generateFieldValues(count, genLaserRetro));}}, {RING_ID_U16, [&](std::size_t count) {setFieldValues(generateFieldValues(count, genRingId));}}, {AZIMUTH_F32, [&](std::size_t count) {setFieldValues(generateFieldValues(count, genAzimuth));}}, diff --git a/test/src/helpers/pointsTest.cpp b/test/src/helpers/pointsTest.cpp index 77fdaa6e..3bbf34c6 100644 --- a/test/src/helpers/pointsTest.cpp +++ b/test/src/helpers/pointsTest.cpp @@ -5,18 +5,20 @@ class PointCloudTest : public ::testing::TestWithParam { protected: - std::vector 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 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 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 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::type> pointCoord; std::vector::type> isHit; std::vector::type> rayIdx; std::vector::type> entityId; std::vector::type> intensity; + std::vector::type> reflectivity; std::vector::type> ringId; std::vector::type> azimuth; std::vector::type> distance; @@ -31,6 +33,7 @@ class PointCloudTest : public ::testing::TestWithParam Field::type rayIdx; Field::type entityId; Field::type intensity; + Field::type reflectivity; Field::type ringId; Field::type azimuth; Field::type distance; @@ -45,8 +48,8 @@ class PointCloudTest : public ::testing::TestWithParam 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 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 pointCloud.setFieldValues(rayIdx); pointCloud.setFieldValues(entityId); pointCloud.setFieldValues(intensity); + pointCloud.setFieldValues(reflectivity); pointCloud.setFieldValues(ringId); pointCloud.setFieldValues(azimuth); pointCloud.setFieldValues(distance); @@ -92,6 +97,7 @@ class PointCloudTest : public ::testing::TestWithParam EXPECT_EQ(pointCloud.getFieldValues(), rayIdx); EXPECT_EQ(pointCloud.getFieldValues(), entityId); EXPECT_EQ(pointCloud.getFieldValues(), intensity); + EXPECT_EQ(pointCloud.getFieldValues(), reflectivity); EXPECT_EQ(pointCloud.getFieldValues(), ringId); EXPECT_EQ(pointCloud.getFieldValues(), azimuth); EXPECT_EQ(pointCloud.getFieldValues(), distance); diff --git a/test/src/scene/textureTest.cpp b/test/src/scene/textureTest.cpp index 91e51b9e..9da3b4b9 100644 --- a/test/src/scene/textureTest.cpp +++ b/test/src/scene/textureTest.cpp @@ -59,7 +59,7 @@ TEST_P(TextureTest, rgl_texture_reading) std::vector rays = {// Ray must be incident perpendicular to the surface to receive all intensity Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}).toRGL()}; - std::vector yieldFields = {INTENSITY_F32}; + std::vector 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(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 rays = {// Ray must be incident perpendicular to the surface to receive all intensity + Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}).toRGL()}; + + std::vector 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::type> outIntensity; + std::vector<::Field::type> outReflectivity; + std::vector<::Field::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) From 918c20d7c5c6da661cfeb801bd357dfdd224bc13 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Wed, 4 Dec 2024 18:01:21 +0100 Subject: [PATCH 03/10] Introduce alpha of reflectivity. --- include/rgl/api/core.h | 9 +++++ src/api/apiCore.cpp | 21 ++++++++++ src/gpu/RaytraceRequestContext.hpp | 1 + src/gpu/optixPrograms.cu | 2 +- src/graph/NodesCore.hpp | 2 + src/graph/RaytraceNode.cpp | 1 + src/tape/TapeCore.hpp | 3 ++ test/CMakeLists.txt | 2 + test/src/scene/textureTest.cpp | 65 ------------------------------ 9 files changed, 40 insertions(+), 66 deletions(-) diff --git a/include/rgl/api/core.h b/include/rgl/api/core.h index 2697c4df..a3c75c90 100644 --- a/include/rgl/api/core.h +++ b/include/rgl/api/core.h @@ -893,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 const for every hit point. + * Default reflectivity alpha is set to 1.0f. + * @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 diff --git a/src/api/apiCore.cpp b/src/api/apiCore.cpp index 38fc516e..750db89e 100644 --- a/src/api/apiCore.cpp +++ b/src/api/apiCore.cpp @@ -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()); } +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(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(); + rgl_node_t node = state.nodes.at(nodeId); + rgl_node_raytrace_configure_reflectivity_alpha(node, yamlNode[1].as()); +} + RGL_API rgl_status_t rgl_node_raytrace_configure_return_mode(rgl_node_t node, rgl_return_mode_t return_mode) { auto status = rglSafeCall([&]() { diff --git a/src/gpu/RaytraceRequestContext.hpp b/src/gpu/RaytraceRequestContext.hpp index 29c28e19..0cc2ca6a 100644 --- a/src/gpu/RaytraceRequestContext.hpp +++ b/src/gpu/RaytraceRequestContext.hpp @@ -29,6 +29,7 @@ struct RaytraceRequestContext float farNonHitDistance; float defaultIntensity; + float reflectivityAlpha; const Mat3x4f* raysWorld; size_t rayCount; diff --git a/src/gpu/optixPrograms.cu b/src/gpu/optixPrograms.cu index 4a9cab5e..c8a6ee4c 100644 --- a/src/gpu/optixPrograms.cu +++ b/src/gpu/optixPrograms.cu @@ -166,7 +166,7 @@ extern "C" __global__ void __closesthit__() } intensity *= cosIncidentAngle; - float reflectivityAlpha = 0.1f; + float reflectivityAlpha = ctx.reflectivityAlpha; float reflectivity = reflectivityAlpha * intensity * distance * distance; Vec3f absPointVelocity{NAN}; diff --git a/src/graph/NodesCore.hpp b/src/graph/NodesCore.hpp index 168c6812..44c48ebc 100644 --- a/src/graph/NodesCore.hpp +++ b/src/graph/NodesCore.hpp @@ -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) { @@ -248,6 +249,7 @@ struct RaytraceNode : IPointsNode float nearNonHitDistance{std::numeric_limits::infinity()}; float farNonHitDistance{std::numeric_limits::infinity()}; float defaultIntensity = 0.0f; + float reflectivityAlpha = 0.1f; MultiReturnSamples mrSampleData = MultiReturnSamples{arrayMgr}; rgl_return_mode_t returnMode = RGL_RETURN_FIRST; diff --git a/src/graph/RaytraceNode.cpp b/src/graph/RaytraceNode.cpp index 66a90f61..e6d9f936 100644 --- a/src/graph/RaytraceNode.cpp +++ b/src/graph/RaytraceNode.cpp @@ -98,6 +98,7 @@ void RaytraceNode::enqueueExecImpl() .nearNonHitDistance = nearNonHitDistance, .farNonHitDistance = farNonHitDistance, .defaultIntensity = defaultIntensity, + .reflectivityAlpha = reflectivityAlpha, .raysWorld = raysPtr, .rayCount = raysNode->getRayCount(), .rayOriginToWorld = raysNode->getCumulativeRayTransfrom(), diff --git a/src/tape/TapeCore.hpp b/src/tape/TapeCore.hpp index 2017cb17..1bdda151 100644 --- a/src/tape/TapeCore.hpp +++ b/src/tape/TapeCore.hpp @@ -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); @@ -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("tape_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), diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index a492542b..e4390c7d 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -48,11 +48,13 @@ 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 src/scene/incidentAngleTest.cpp src/graph/multiReturnTest.cpp + src/scene/reflectivityTest.cpp ) if (RGL_BUILD_ROS2_EXTENSION) diff --git a/test/src/scene/textureTest.cpp b/test/src/scene/textureTest.cpp index 9da3b4b9..e9036d0e 100644 --- a/test/src/scene/textureTest.cpp +++ b/test/src/scene/textureTest.cpp @@ -87,71 +87,6 @@ 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(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 rays = {// Ray must be incident perpendicular to the surface to receive all intensity - Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}).toRGL()}; - - std::vector 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::type> outIntensity; - std::vector<::Field::type> outReflectivity; - std::vector<::Field::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) From 985c465f2c11f40d34cab022e931ab4070930cd1 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Wed, 4 Dec 2024 18:05:46 +0100 Subject: [PATCH 04/10] Cleanup --- include/rgl/api/core.h | 4 ++-- src/RGLFields.hpp | 2 +- test/CMakeLists.txt | 1 - 3 files changed, 3 insertions(+), 4 deletions(-) diff --git a/include/rgl/api/core.h b/include/rgl/api/core.h index a3c75c90..66ece031 100644 --- a/include/rgl/api/core.h +++ b/include/rgl/api/core.h @@ -895,8 +895,8 @@ RGL_API rgl_status_t rgl_node_raytrace_configure_default_intensity(rgl_node_t no /** * Modifies RaytraceNode to set reflectivity alpha. - * Reflectivity alpha is used to calculate reflectivity of the hit point. This value is const for every hit point. - * Default reflectivity alpha is set to 1.0f. + * 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. * @param node RaytraceNode to modify. * @param reflectivity_alpha Reflectivity alpha to set. */ diff --git a/src/RGLFields.hpp b/src/RGLFields.hpp index f979bcdb..14330cf9 100644 --- a/src/RGLFields.hpp +++ b/src/RGLFields.hpp @@ -74,7 +74,7 @@ inline const std::set& getAllRealFields() ENTITY_ID_I32, INTENSITY_F32, INTENSITY_U8, - REFLECTIVITY_F32, + REFLECTIVITY_F32, LASER_RETRO_F32, RING_ID_U16, AZIMUTH_F32, diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index e4390c7d..c3c626fb 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -54,7 +54,6 @@ set(RGL_TEST_FILES src/synchronization/testKernel.cu src/scene/incidentAngleTest.cpp src/graph/multiReturnTest.cpp - src/scene/reflectivityTest.cpp ) if (RGL_BUILD_ROS2_EXTENSION) From 0cedd337f33890989c5671d4414855bda59a2aa7 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Wed, 4 Dec 2024 18:20:00 +0100 Subject: [PATCH 05/10] Fix tape call --- src/tape/TapeCore.hpp | 2 +- test/src/TapeTest.cpp | 3 +++ 2 files changed, 4 insertions(+), 1 deletion(-) diff --git a/src/tape/TapeCore.hpp b/src/tape/TapeCore.hpp index 1bdda151..fa378fde 100644 --- a/src/tape/TapeCore.hpp +++ b/src/tape/TapeCore.hpp @@ -123,7 +123,7 @@ 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("tape_node_raytrace_configure_reflectivity_alpha", + 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), diff --git a/test/src/TapeTest.cpp b/test/src/TapeTest.cpp index 4e7654e1..4b4e2c8a 100644 --- a/test/src/TapeTest.cpp +++ b/test/src/TapeTest.cpp @@ -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)); From 085115589a60aa808925f1559867a7b451ae8f68 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Wed, 4 Dec 2024 18:25:41 +0100 Subject: [PATCH 06/10] Introduce separate test. --- test/src/scene/reflectivityTest.cpp | 82 +++++++++++++++++++++++++++++ test/src/scene/textureTest.cpp | 2 +- 2 files changed, 83 insertions(+), 1 deletion(-) create mode 100644 test/src/scene/reflectivityTest.cpp diff --git a/test/src/scene/reflectivityTest.cpp b/test/src/scene/reflectivityTest.cpp new file mode 100644 index 00000000..9b93366f --- /dev/null +++ b/test/src/scene/reflectivityTest.cpp @@ -0,0 +1,82 @@ +#include +#include +#include +#include +#include + +#include +struct ReflectivityTest : public RGLTestWithParam> +{}; + +INSTANTIATE_TEST_SUITE_P(Parametrized, ReflectivityTest, + testing::Combine(testing::Values(0.0f, 0.1f, 10.0f), testing::Values(u_char(0), u_char(127), u_char(255)))); + +TEST_P(ReflectivityTest, read_value) +{ + auto [alpha, value] = GetParam(); + + int textureWidth = 100; + int textureHeight = 100; + + rgl_texture_t texture = nullptr; + rgl_entity_t entity = nullptr; + rgl_mesh_t mesh = makeCubeMesh(); + auto textureRawData = generateStaticColorTexture(textureWidth, textureHeight, value); + + EXPECT_RGL_SUCCESS(rgl_texture_create(&texture, textureRawData.data(), textureWidth, textureHeight)); + 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 rays = {// Ray must be incident perpendicular to the surface to receive all intensity + Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}).toRGL()}; + + std::vector 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_node_raytrace_configure_reflectivity_alpha(raytraceNode, alpha)); + + EXPECT_RGL_SUCCESS(rgl_graph_run(raytraceNode)); + + std::vector<::Field::type> outIntensity; + std::vector<::Field::type> outReflectivity; + std::vector<::Field::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 = alpha * distance * distance * intensity; + EXPECT_NEAR(reflectivityValue, outReflectivity.at(i), EPSILON_F); + } +} \ No newline at end of file diff --git a/test/src/scene/textureTest.cpp b/test/src/scene/textureTest.cpp index e9036d0e..91e51b9e 100644 --- a/test/src/scene/textureTest.cpp +++ b/test/src/scene/textureTest.cpp @@ -59,7 +59,7 @@ TEST_P(TextureTest, rgl_texture_reading) std::vector rays = {// Ray must be incident perpendicular to the surface to receive all intensity Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}).toRGL()}; - std::vector yieldFields = {INTENSITY_F32, REFLECTIVITY_F32}; + std::vector yieldFields = {INTENSITY_F32}; EXPECT_RGL_SUCCESS(rgl_node_rays_from_mat3x4f(&useRaysNode, rays.data(), rays.size())); EXPECT_RGL_SUCCESS(rgl_node_raytrace(&raytraceNode, nullptr)); From f422484025efefe9ed5a98687a2b049a0f6bff60 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Mon, 16 Dec 2024 11:40:28 +0100 Subject: [PATCH 07/10] PR fixes v1 --- include/rgl/api/core.h | 8 +++++++- src/gpu/MultiReturn.hpp | 1 - src/gpu/nodeKernels.cu | 3 ++- src/gpu/optixPrograms.cu | 8 +++----- src/graph/NodesCore.hpp | 6 +----- test/src/scene/reflectivityTest.cpp | 25 +++++++++++++++++++------ 6 files changed, 32 insertions(+), 19 deletions(-) diff --git a/include/rgl/api/core.h b/include/rgl/api/core.h index 66ece031..074adfb6 100644 --- a/include/rgl/api/core.h +++ b/include/rgl/api/core.h @@ -358,7 +358,6 @@ 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, @@ -442,6 +441,13 @@ typedef enum : int32_t */ RGL_FIELD_LASER_RETRO_F32, + /** + * Reflectivity of the hit point. + * Reflectivity is a scalar value that describes how much light is reflected by the hit point. + * Reflectivity is a property of the material of the hit point. + */ + RGL_FIELD_REFLECTIVITY_F32, + // Dummy fields RGL_FIELD_PADDING_8 = 1024, RGL_FIELD_PADDING_16, diff --git a/src/gpu/MultiReturn.hpp b/src/gpu/MultiReturn.hpp index c924826a..028f303a 100644 --- a/src/gpu/MultiReturn.hpp +++ b/src/gpu/MultiReturn.hpp @@ -27,7 +27,6 @@ struct MultiReturnSamplesPointers Field::type* isHit; Field::type* distance; Field::type* intensity; - Field::type* reflectivity; Field::type* laserRetro; Field::type* entityId; Field::type* absVelocity; diff --git a/src/gpu/nodeKernels.cu b/src/gpu/nodeKernels.cu index bb4e78e8..17d8b2fe 100644 --- a/src/gpu/nodeKernels.cu +++ b/src/gpu/nodeKernels.cu @@ -128,7 +128,8 @@ __device__ void saveReturnAsHit(const RaytraceRequestContext* ctx, int beamIdx, UINT8_MAX; } if (ctx->reflectivityF32 != nullptr) { - ctx->reflectivityF32[returnPointIdx] = ctx->mrSamples.reflectivity[sampleIdx]; + const auto distance2 = ctx->mrSamples.distance[sampleIdx] * ctx->mrSamples.distance[sampleIdx]; + ctx->reflectivityF32[returnPointIdx] = ctx->reflectivityAlpha * ctx->mrSamples.intensity[sampleIdx] * distance2; } if (ctx->entityId != nullptr) { ctx->entityId[returnPointIdx] = ctx->mrSamples.entityId[sampleIdx]; diff --git a/src/gpu/optixPrograms.cu b/src/gpu/optixPrograms.cu index c8a6ee4c..b19e160c 100644 --- a/src/gpu/optixPrograms.cu +++ b/src/gpu/optixPrograms.cu @@ -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 reflectivity, float laserRetro, int objectID, +__device__ void saveSampleAsHit(int sampleIdx, float distance, float intensity, float laserRetro, int objectID, const Vec3f& absVelocity, const Vec3f& relVelocity, float radialSpeed, const Vec3f& normal, float incidentAngle); __device__ void saveNonHitBeamSamples(int beamIdx, float nonHitDistance); @@ -167,7 +167,6 @@ extern "C" __global__ void __closesthit__() intensity *= cosIncidentAngle; float reflectivityAlpha = ctx.reflectivityAlpha; - float reflectivity = reflectivityAlpha * intensity * distance * distance; Vec3f absPointVelocity{NAN}; Vec3f relPointVelocity{NAN}; @@ -214,7 +213,7 @@ extern "C" __global__ void __closesthit__() radialSpeed = hitRays.normalized().dot(relPointVelocity); } - saveSampleAsHit(mrSampleIdx, distance, intensity, reflectivity, laserRetro, entityId, absPointVelocity, relPointVelocity, radialSpeed, + saveSampleAsHit(mrSampleIdx, distance, intensity, laserRetro, entityId, absPointVelocity, relPointVelocity, radialSpeed, wNormal, incidentAngle); } @@ -256,14 +255,13 @@ __device__ void saveSampleAsNonHit(int sampleIdx, float nonHitDistance) ctx.mrSamples.distance[sampleIdx] = nonHitDistance; } -__device__ void saveSampleAsHit(int sampleIdx, float distance, float intensity,float reflectivity, float laserRetro, int objectID, +__device__ void saveSampleAsHit(int sampleIdx, float distance, float intensity, 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; diff --git a/src/graph/NodesCore.hpp b/src/graph/NodesCore.hpp index 44c48ebc..aac211e5 100644 --- a/src/graph/NodesCore.hpp +++ b/src/graph/NodesCore.hpp @@ -159,8 +159,7 @@ struct RaytraceNode : IPointsNode explicit MultiReturnSamples(StreamBoundObjectsManager& arrayMgr) : isHit(DeviceAsyncArray::type>::create(arrayMgr)), distance(DeviceAsyncArray::type>::create(arrayMgr)), - intensity(DeviceAsyncArray::type>::create(arrayMgr)), - reflectivity(DeviceAsyncArray::type>::create(arrayMgr)) + intensity(DeviceAsyncArray::type>::create(arrayMgr)) {} void adjustToFields(const std::unordered_map& inFieldData, @@ -185,7 +184,6 @@ 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); @@ -203,7 +201,6 @@ 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, @@ -227,7 +224,6 @@ struct RaytraceNode : IPointsNode DeviceAsyncArray::type>::Ptr isHit; DeviceAsyncArray::type>::Ptr distance; DeviceAsyncArray::type>::Ptr intensity; - DeviceAsyncArray::type>::Ptr reflectivity; // Additional field data for multi-return samples. DeviceAsyncArray::type>::Ptr laserRetro; diff --git a/test/src/scene/reflectivityTest.cpp b/test/src/scene/reflectivityTest.cpp index 9b93366f..a83b6ff4 100644 --- a/test/src/scene/reflectivityTest.cpp +++ b/test/src/scene/reflectivityTest.cpp @@ -5,15 +5,19 @@ #include #include -struct ReflectivityTest : public RGLTestWithParam> +struct ReflectivityTest : public RGLTestWithParam> {}; INSTANTIATE_TEST_SUITE_P(Parametrized, ReflectivityTest, - testing::Combine(testing::Values(0.0f, 0.1f, 10.0f), testing::Values(u_char(0), u_char(127), u_char(255)))); + testing::Combine( + testing::Values(0.012f, 0.12f, 1.23f), + testing::Values(u_char(0), u_char(127), u_char(255)), + testing::Values(0.012, 0.123, 1.23) + )); TEST_P(ReflectivityTest, read_value) { - auto [alpha, value] = GetParam(); + auto [alpha, value, distance] = GetParam(); int textureWidth = 100; int textureHeight = 100; @@ -29,6 +33,10 @@ TEST_P(ReflectivityTest, read_value) EXPECT_RGL_SUCCESS(rgl_entity_create(&entity, nullptr, mesh)); EXPECT_RGL_SUCCESS(rgl_entity_set_intensity_texture(entity, texture)); + // Scale the cube in order to test different distances + rgl_mat3x4f cubeTransform = Mat3x4f::TRS({0, 0, 0}, {0, 0, 0}, {distance, distance, distance}).toRGL(); + EXPECT_RGL_SUCCESS(rgl_entity_set_transform(entity, &cubeTransform)); + // Create RGL graph pipeline. rgl_node_t useRaysNode = nullptr, raytraceNode = nullptr, compactNode = nullptr, yieldNode = nullptr; @@ -74,9 +82,14 @@ TEST_P(ReflectivityTest, read_value) for (int i = 0; i < outCount; ++i) { EXPECT_NEAR(((float) value), outIntensity.at(i), EPSILON_F); - float distance = outDistance.at(i); + float outDdistance = outDistance.at(i); float intensity = outIntensity.at(i); - float reflectivityValue = alpha * distance * distance * intensity; - EXPECT_NEAR(reflectivityValue, outReflectivity.at(i), EPSILON_F); + float reflectivityValue = alpha * outDdistance * outDdistance * intensity; + printf("Distance: %f, Intensity: %f, Reflectivity: %f, Alpha: %f\n", outDdistance, intensity, reflectivityValue, alpha); + + // Reflectivity test is conducted with greater epsilon. + // This is due to lack of distance impact on intensity. + // As long as distance is not included into intensity calculations, reflectivity value will grow relatively fast with the distance. + EXPECT_NEAR(reflectivityValue, outReflectivity.at(i), 1e-3f); } } \ No newline at end of file From 97ea7775dd720991320e49304ea6d66f57a5ed17 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Mon, 16 Dec 2024 12:12:00 +0100 Subject: [PATCH 08/10] minor formating fixes --- include/rgl/api/core.h | 8 ++++---- test/src/scene/reflectivityTest.cpp | 3 +-- 2 files changed, 5 insertions(+), 6 deletions(-) diff --git a/include/rgl/api/core.h b/include/rgl/api/core.h index 074adfb6..e4414d69 100644 --- a/include/rgl/api/core.h +++ b/include/rgl/api/core.h @@ -442,10 +442,10 @@ typedef enum : int32_t RGL_FIELD_LASER_RETRO_F32, /** - * Reflectivity of the hit point. - * Reflectivity is a scalar value that describes how much light is reflected by the hit point. - * Reflectivity is a property of the material of the hit point. - */ + * Reflectivity of the hit point. + * Reflectivity is a scalar value that describes how much light is reflected by the hit point. + * Reflectivity is a property of the material of the hit point. + */ RGL_FIELD_REFLECTIVITY_F32, // Dummy fields diff --git a/test/src/scene/reflectivityTest.cpp b/test/src/scene/reflectivityTest.cpp index a83b6ff4..efc1fc10 100644 --- a/test/src/scene/reflectivityTest.cpp +++ b/test/src/scene/reflectivityTest.cpp @@ -12,8 +12,7 @@ INSTANTIATE_TEST_SUITE_P(Parametrized, ReflectivityTest, testing::Combine( testing::Values(0.012f, 0.12f, 1.23f), testing::Values(u_char(0), u_char(127), u_char(255)), - testing::Values(0.012, 0.123, 1.23) - )); + testing::Values(0.012, 0.123, 1.23))); TEST_P(ReflectivityTest, read_value) { From 63c43113d416e5fec552950630d3105978232fa8 Mon Sep 17 00:00:00 2001 From: Piotr Mrozik <33581747+PiotrMrozik@users.noreply.github.com> Date: Mon, 16 Dec 2024 13:28:31 +0100 Subject: [PATCH 09/10] Update test/src/scene/reflectivityTest.cpp Co-authored-by: Piotr Rybicki --- test/src/scene/reflectivityTest.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/test/src/scene/reflectivityTest.cpp b/test/src/scene/reflectivityTest.cpp index efc1fc10..c537fecf 100644 --- a/test/src/scene/reflectivityTest.cpp +++ b/test/src/scene/reflectivityTest.cpp @@ -81,10 +81,9 @@ TEST_P(ReflectivityTest, read_value) for (int i = 0; i < outCount; ++i) { EXPECT_NEAR(((float) value), outIntensity.at(i), EPSILON_F); - float outDdistance = outDistance.at(i); + float outDistance = outDistance.at(i); float intensity = outIntensity.at(i); - float reflectivityValue = alpha * outDdistance * outDdistance * intensity; - printf("Distance: %f, Intensity: %f, Reflectivity: %f, Alpha: %f\n", outDdistance, intensity, reflectivityValue, alpha); + float reflectivityValue = alpha * outDistance * outDistance * intensity; // Reflectivity test is conducted with greater epsilon. // This is due to lack of distance impact on intensity. From c1f04f9135edca994cdc3adea4675be7fe36bd18 Mon Sep 17 00:00:00 2001 From: PiotrMrozik Date: Mon, 16 Dec 2024 13:36:42 +0100 Subject: [PATCH 10/10] fix naming conflict --- src/gpu/optixPrograms.cu | 2 -- test/src/scene/reflectivityTest.cpp | 4 ++-- 2 files changed, 2 insertions(+), 4 deletions(-) diff --git a/src/gpu/optixPrograms.cu b/src/gpu/optixPrograms.cu index b19e160c..16b3cc95 100644 --- a/src/gpu/optixPrograms.cu +++ b/src/gpu/optixPrograms.cu @@ -166,8 +166,6 @@ extern "C" __global__ void __closesthit__() } intensity *= cosIncidentAngle; - float reflectivityAlpha = ctx.reflectivityAlpha; - Vec3f absPointVelocity{NAN}; Vec3f relPointVelocity{NAN}; float radialSpeed{NAN}; diff --git a/test/src/scene/reflectivityTest.cpp b/test/src/scene/reflectivityTest.cpp index c537fecf..1f6b6eee 100644 --- a/test/src/scene/reflectivityTest.cpp +++ b/test/src/scene/reflectivityTest.cpp @@ -81,9 +81,9 @@ TEST_P(ReflectivityTest, read_value) for (int i = 0; i < outCount; ++i) { EXPECT_NEAR(((float) value), outIntensity.at(i), EPSILON_F); - float outDistance = outDistance.at(i); + float outDistanceValue = outDistance.at(i); float intensity = outIntensity.at(i); - float reflectivityValue = alpha * outDistance * outDistance * intensity; + float reflectivityValue = alpha * outDistanceValue * outDistanceValue * intensity; // Reflectivity test is conducted with greater epsilon. // This is due to lack of distance impact on intensity.