Auto 835 add laser retro support to rgl plugin

include/rgl/api/core.h

@@ -190,6 +190,7 @@ typedef enum
 	RGL_FIELD_RING_ID_U16,
 	RGL_FIELD_RETURN_TYPE_U8,
 	RGL_FIELD_TIME_STAMP_F64,
+	RGL_FIELD_LASER_RETRO_F32,
 	// Dummy fields
 	RGL_FIELD_PADDING_8 = 1024,
 	RGL_FIELD_PADDING_16,
@@ -308,6 +309,15 @@ rgl_entity_destroy(rgl_entity_t entity);
 RGL_API rgl_status_t
 rgl_entity_set_pose(rgl_entity_t entity, const rgl_mat3x4f *local_to_world_tf);
 
+
+/**
+ * Set laser retro value for the given entity.
+ * @param entity Entity to modify
+ * @param int laser retrovalue to set. If not set,it will be set to 100.0.
+ */
+RGL_API rgl_status_t
+rgl_entity_set_laser_retro(rgl_entity_t entity, float retro);
+
 /******************************** NODES ********************************/
 
 /**

src/RGLFields.hpp

@@ -33,6 +33,7 @@
 #define DISTANCE_F32 RGL_FIELD_DISTANCE_F32
 #define RETURN_TYPE_U8 RGL_FIELD_RETURN_TYPE_U8
 #define TIME_STAMP_F64 RGL_FIELD_TIME_STAMP_F64
+#define LASER_RETRO_F32 RGL_FIELD_LASER_RETRO_F32
 #define PADDING_8 RGL_FIELD_PADDING_8
 #define PADDING_16 RGL_FIELD_PADDING_16
 #define PADDING_32 RGL_FIELD_PADDING_32
@@ -57,6 +58,7 @@ FIELD(AZIMUTH_F32, float);
 FIELD(RING_ID_U16, uint16_t);
 FIELD(RETURN_TYPE_U8, uint8_t);
 FIELD(TIME_STAMP_F64, double);
+FIELD(LASER_RETRO_F32,float);
 FIELD(PADDING_8, uint8_t);
 FIELD(PADDING_16, uint16_t);
 FIELD(PADDING_32, uint32_t);
@@ -73,6 +75,7 @@ inline std::size_t getFieldSize(rgl_field_t type)
 		case DISTANCE_F32: return Field<DISTANCE_F32>::size;
 		case RETURN_TYPE_U8: return Field<RETURN_TYPE_U8>::size;
 		case TIME_STAMP_F64: return Field<TIME_STAMP_F64>::size;
+		case LASER_RETRO_F32: return Field<LASER_RETRO_F32>::size;
 		case PADDING_8: return Field<PADDING_8>::size;
 		case PADDING_16: return Field<PADDING_16>::size;
 		case PADDING_32: return Field<PADDING_32>::size;
@@ -113,6 +116,7 @@ inline VArray::Ptr createVArray(rgl_field_t type, std::size_t initialSize)
 		case DISTANCE_F32: return VArray::create<Field<DISTANCE_F32>::type>(initialSize);
 		case RETURN_TYPE_U8: return VArray::create<Field<RETURN_TYPE_U8>::type>(initialSize);
 		case TIME_STAMP_F64: return VArray::create<Field<TIME_STAMP_F64>::type>(initialSize);
+		case LASER_RETRO_F32: return VArray::create<Field<LASER_RETRO_F32>::type>(initialSize);
 		case IS_HIT_I32: return VArray::create<Field<IS_HIT_I32>::type>(initialSize);
 	}
 	throw std::invalid_argument(fmt::format("createVArray: unknown RGL field {}", type));
@@ -130,6 +134,7 @@ inline std::string toString(rgl_field_t type)
 		case DISTANCE_F32: return "DISTANCE_F32";
 		case RETURN_TYPE_U8: return "RETURN_TYPE_U8";
 		case TIME_STAMP_F64: return "TIME_STAMP_F64";
+		case LASER_RETRO_F32: return "LASER_RETRO_F32";
 		case PADDING_8: return "PADDING_8";
 		case PADDING_16: return "PADDING_16";
 		case PADDING_32: return "PADDING_32";

src/Tape.hpp

@@ -169,6 +169,7 @@ class TapePlay
 	void tape_entity_create(const YAML::Node& yamlNode);
 	void tape_entity_destroy(const YAML::Node& yamlNode);
 	void tape_entity_set_pose(const YAML::Node& yamlNode);
+	void tape_entity_set_laser_retro(const YAML::Node& yamlNode);
 	void tape_graph_run(const YAML::Node& yamlNode);
 	void tape_graph_destroy(const YAML::Node& yamlNode);
 	void tape_graph_get_result_size(const YAML::Node& yamlNode);

src/api/api.cpp

@@ -427,6 +427,25 @@ void TapePlay::tape_entity_set_pose(const YAML::Node& yamlNode)
 		reinterpret_cast<const rgl_mat3x4f*>(fileMmap + yamlNode[1].as<size_t>()));
 }
 
+RGL_API rgl_status_t
+rgl_entity_set_laser_retro(rgl_entity_t entity, float retro)
+{
+	auto status = rglSafeCall([&]() {
+		RGL_API_LOG("rgl_entity_set_laser_retro(entity={}, retro={})", (void *) entity, retro);
+		CHECK_ARG(entity != nullptr);
+		CHECK_ARG(retro > 0);
+		Entity::validatePtr(entity)->setLaserRetro(retro);
+	});
+	TAPE_HOOK(entity, retro);
+	return status;
+}
+
+void TapePlay::tape_entity_set_laser_retro(const YAML::Node& yamlNode)
+{
+	rgl_entity_set_laser_retro(tapeEntities[yamlNode[0].as<size_t>()],
+					  yamlNode[1].as<Field<LASER_RETRO_F32>::type>());
+}
+
 RGL_API rgl_status_t
 rgl_graph_run(rgl_node_t node)
 {

src/gpu/RaytraceRequestContext.hpp

@@ -38,5 +38,6 @@ struct RaytraceRequestContext
 	Field<RING_ID_U16>::type* ringIdx;
 	Field<DISTANCE_F32>::type* distance;
 	Field<INTENSITY_F32>::type* intensity;
+	Field<LASER_RETRO_F32>::type* laserRetro;
 };
 static_assert(std::is_trivially_copyable<RaytraceRequestContext>::value);

src/gpu/ShaderBindingTableTypes.h

@@ -8,6 +8,7 @@ struct TriangleMeshSBTData {
 	const Vec3i *index;
 	size_t vertex_count;
 	size_t index_count;
+    float laser_retro;
 };
 
 

src/gpu/optixPrograms.cu

@@ -23,6 +23,8 @@
 #include <gpu/RaytraceRequestContext.hpp>
 #include <gpu/ShaderBindingTableTypes.h>
 
+#define DEFAULT_LASER_RETRO 100.0
+
 extern "C" static __constant__ RaytraceRequestContext ctx;
 
 struct Vec3fPayload
@@ -54,7 +56,7 @@ Vec3f decodePayloadVec3f(const Vec3fPayload& src)
 
 template<bool isFinite>
 __forceinline__ __device__
-void saveRayResult(const Vec3f* xyz=nullptr, const Vec3f* origin=nullptr)
+void saveRayResult(const Vec3f* xyz=nullptr, const Vec3f* origin=nullptr, const float retro = DEFAULT_LASER_RETRO)
 {
 	const int rayIdx = optixGetLaunchIndex().x;
 	if (ctx.xyz != nullptr) {
@@ -81,6 +83,9 @@ void saveRayResult(const Vec3f* xyz=nullptr, const Vec3f* origin=nullptr)
 	if (ctx.intensity != nullptr) {
 		ctx.intensity[rayIdx] = 100;
 	}
+	if (ctx.laserRetro != nullptr) {
+		ctx.laserRetro[rayIdx] = isFinite? retro : 0.0;
+	}
 }
 
 extern "C" __global__ void __raygen__()
@@ -121,12 +126,14 @@ extern "C" __global__ void __closesthit__()
 	Vec3f hitObject = Vec3f((1 - u - v) * A + u * B + v * C);
 	Vec3f hitWorld = optixTransformPointFromObjectToWorldSpace(hitObject);
 
+	float retro = sbtData.laser_retro;
+
 	Vec3f origin = decodePayloadVec3f({
 		optixGetPayload_0(),
 		optixGetPayload_1(),
 		optixGetPayload_2()
 	});
-	saveRayResult<true>(&hitWorld, &origin);
+	saveRayResult<true>(&hitWorld, &origin, retro);
 }
 
 extern "C" __global__ void __miss__()

src/graph/RaytraceNode.cpp

@@ -59,6 +59,7 @@ void RaytraceNode::schedule(cudaStream_t stream)
 		.ringIdx = getPtrTo<RING_ID_U16>(),
 		.distance = getPtrTo<DISTANCE_F32>(),
 		.intensity = getPtrTo<INTENSITY_F32>(),
+		.laserRetro = getPtrTo<LASER_RETRO_F32>(),
 	};
 
 	CUdeviceptr pipelineArgsPtr = requestCtx->getCUdeviceptr();

src/scene/Entity.cpp

@@ -19,7 +19,8 @@ API_OBJECT_INSTANCE(Entity);
 Entity::Entity(std::shared_ptr<Mesh> mesh, std::optional<std::string> name)
 : mesh(std::move(mesh))
 , transform(Mat3x4f::identity())
-, humanReadableName(std::move(name)) { }
+, humanReadableName(std::move(name))
+, laser_retro(DEFAULT_LASER_RETRO) { }
 
 void Entity::setTransform(Mat3x4f newTransform)
 {
@@ -29,6 +30,14 @@ void Entity::setTransform(Mat3x4f newTransform)
 	}
 }
 
+void Entity::setLaserRetro(float retro)
+{
+	laser_retro = retro;
+	if (auto activeScene = scene.lock()) {
+		activeScene->requestASRebuild();
+	}
+}
+
 OptixInstance Entity::getIAS(int idx)
 {
 	// NOTE: this assumes a single SBT record per GAS

src/scene/Entity.hpp

@@ -20,6 +20,7 @@
 #include <utility>
 #include <math/Mat3x4f.hpp>
 
+constexpr float DEFAULT_LASER_RETRO = 100.0;
 
 struct Entity : APIObject<Entity>
 {
@@ -28,11 +29,13 @@ struct Entity : APIObject<Entity>
 	// TODO(prybicki): low-prio optimization: do not rebuild whole IAS if only transform changed
 	void setTransform(Mat3x4f newTransform);
 	OptixInstance getIAS(int idx);
-
+	void setLaserRetro(float retro);
+	const float getLaserRetro() { return laser_retro;}
 	std::shared_ptr<Mesh> mesh;
 	std::weak_ptr<Scene> scene;
 private:
 	Mat3x4f transform;
+	float laser_retro;
 	std::optional<std::string> humanReadableName;
 	friend struct APIObject<Entity>;
 	friend struct Scene;

src/scene/Scene.cpp

@@ -85,6 +85,7 @@ OptixShaderBindingTable Scene::buildSBT()
 			.index = mesh->dIndices.readDevice(),
 			.vertex_count = mesh->dVertices.getElemCount(),
 			.index_count = mesh->dIndices.getElemCount(),
+			.laser_retro = entity->getLaserRetro(),
 		};
 	}
 	dHitgroupRecords.copyFromHost(hHitgroupRecords);