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.

Sign up for GitHub

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

Architectural changes to support multi-return feature #285

Merged
merged 4 commits into from
May 23, 2024
Merged

Architectural changes to support multi-return feature #285

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
d289d82
Architectural changes to support multi-return feature
prybicki May 15, 2024
bc8cfa2
Review fixes
prybicki May 15, 2024
b5be808
Review fixes
prybicki May 23, 2024
778547e
Review fixes
prybicki May 23, 2024
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
21 changes: 21 additions & 0 deletions src/api/apiCore.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1304,6 +1304,27 @@ void TapeCore::tape_node_gaussian_noise_distance(const YAML::Node& yamlNode, Pla
state.nodes.insert({nodeId, node});
}

RGL_API rgl_status_t rgl_node_multi_return_switch(rgl_node_t* node, rgl_return_type_t return_type)
{
auto status = rglSafeCall([&]() {
RGL_API_LOG("rgl_node_multi_return_switch(node={}, return_type={})", repr(node), return_type);
CHECK_ARG(node != nullptr);

createOrUpdateNode<MultiReturnSwitchNode>(node, return_type);
});
TAPE_HOOK(node, return_type);
return status;
}

void TapeCore::tape_node_multi_return_switch(const YAML::Node& yamlNode, PlaybackState& state)
{
auto nodeId = yamlNode[0].as<TapeAPIObjectID>();
auto return_type = (rgl_return_type_t) yamlNode[1].as<int>();
prybicki marked this conversation as resolved.
Show resolved Hide resolved
rgl_node_t node = state.nodes.contains(nodeId) ? state.nodes.at(nodeId) : nullptr;
rgl_node_multi_return_switch(&node, return_type);
state.nodes.insert({nodeId, node});
}

rgl_status_t rgl_node_is_alive(rgl_node_t node, bool* out_alive)
{
auto status = rglSafeCall([&]() {
Expand Down
29 changes: 18 additions & 11 deletions src/gpu/nodeKernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -229,25 +229,32 @@ __global__ void kProcessBeamSamplesFirstLast(size_t beamCount, int samplesPerBea
LIMIT(beamCount);

const auto beamIdx = tid;
int firstIdx = 0;
int lastIdx = 0;
int firstIdx = -1;
int lastIdx = -1;
for (int sampleIdx = 0; sampleIdx < samplesPerBeam; ++sampleIdx) {
if (beamSamples.isHit[beamIdx * samplesPerBeam + sampleIdx] == 0) {
continue;
}
if (beamSamples.distance[beamIdx * samplesPerBeam + sampleIdx] <
beamSamples.distance[beamIdx * samplesPerBeam + firstIdx]) {
auto currentFirstDistance = firstIdx >= 0 ? beamSamples.distance[beamIdx * samplesPerBeam + firstIdx] : FLT_MAX;
auto currentLastDistance = lastIdx >= 0 ? beamSamples.distance[beamIdx * samplesPerBeam + lastIdx] : -FLT_MAX;
if (beamSamples.distance[beamIdx * samplesPerBeam + sampleIdx] < currentFirstDistance) {
firstIdx = sampleIdx;
}
if (beamSamples.distance[beamIdx * samplesPerBeam + sampleIdx] >
beamSamples.distance[beamIdx * samplesPerBeam + lastIdx]) {
if (beamSamples.distance[beamIdx * samplesPerBeam + sampleIdx] > currentLastDistance) {
PawelLiberadzki marked this conversation as resolved.
Show resolved Hide resolved
lastIdx = sampleIdx;
}
}
first.xyz[beamIdx] = beamSamples.xyz[beamIdx * samplesPerBeam + firstIdx];
first.distance[beamIdx] = beamSamples.distance[beamIdx * samplesPerBeam + firstIdx];
last.xyz[beamIdx] = beamSamples.xyz[beamIdx * samplesPerBeam + lastIdx];
last.distance[beamIdx] = beamSamples.distance[beamIdx * samplesPerBeam + lastIdx];
first.isHit[beamIdx] = firstIdx >= 0;
if (first.isHit[beamIdx]) {
first.xyz[beamIdx] = beamSamples.xyz[beamIdx * samplesPerBeam + firstIdx];
first.distance[beamIdx] = beamSamples.distance[beamIdx * samplesPerBeam + firstIdx];
PawelLiberadzki marked this conversation as resolved.
Show resolved Hide resolved
}

last.isHit[beamIdx] = lastIdx >= 0;
if (last.isHit[beamIdx]) {
last.xyz[beamIdx] = beamSamples.xyz[beamIdx * samplesPerBeam + lastIdx];
last.distance[beamIdx] = beamSamples.distance[beamIdx * samplesPerBeam + lastIdx];
}
Copy link
Contributor

Choose a reason for hiding this comment

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

Is it not like if beamSamples.isHit[beamIdx * samplesPerBeam + sampleIdx] != 0 (line 235 above) then you will always have first and last? If this is true, then two if-statements here may technically be reduced to one.

}


Expand Down Expand Up @@ -327,4 +334,4 @@ void gpuProcessBeamSamplesFirstLast(cudaStream_t stream, size_t beamCount, int s
MultiReturnPointers first, MultiReturnPointers last)
{
run(kProcessBeamSamplesFirstLast, stream, beamCount, samplesPerBeam, beamSamples, first, last);
}
}
29 changes: 29 additions & 0 deletions src/graph/NodesCore.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -124,6 +124,8 @@ struct RaytraceNode : IPointsNode
return std::const_pointer_cast<const IAnyArray>(fieldData.at(field));
}

IAnyArray::ConstPtr getFieldDataMultiReturn(rgl_field_t field, rgl_return_type_t);

// RaytraceNode specific
void setVelocity(const Vec3f& linearVelocity, const Vec3f& angularVelocity);
void enableRayDistortion(bool enabled) { doApplyDistortion = enabled; }
Expand Down Expand Up @@ -189,6 +191,33 @@ struct RaytraceNode : IPointsNode
void setFields(const std::set<rgl_field_t>& fields);
};

struct MultiReturnSwitchNode : IPointsNodeSingleInput
{
using Ptr = std::shared_ptr<MultiReturnSwitchNode>;
void setParameters(rgl_return_type_t returnType) { this->returnType = returnType; }

// Node
void validateImpl() override
{
IPointsNodeSingleInput::validateImpl();
rtxInput = getExactlyOneInputOfType<RaytraceNode>();
}

void enqueueExecImpl() override {}

// Point cloud description
Copy link
Contributor

Choose a reason for hiding this comment

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

Probably unsed docs from copy/paste?


// Data getters
IAnyArray::ConstPtr getFieldData(rgl_field_t field) override
{
return rtxInput->getFieldDataMultiReturn(field, returnType);
}

private:
rgl_return_type_t returnType;
RaytraceNode::Ptr rtxInput;
};

struct TransformPointsNode : IPointsNodeSingleInput
{
using Ptr = std::shared_ptr<TransformPointsNode>;
Expand Down
21 changes: 21 additions & 0 deletions src/graph/RaytraceNode.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -137,6 +137,27 @@ void RaytraceNode::enqueueExecImpl()
mrSamples.getPointers(), mrFirst.getPointers(), mrLast.getPointers());
}

IAnyArray::ConstPtr RaytraceNode::getFieldDataMultiReturn(rgl_field_t field, rgl_return_type_t type)
{
if (type == RGL_RETURN_TYPE_FIRST) {
switch (field) {
case XYZ_VEC3_F32: return mrFirst.xyz;
case DISTANCE_F32: return mrFirst.distance;
case IS_HIT_I32: return mrFirst.isHit;
default: throw InvalidPipeline(fmt::format("Multi-Return not supported for this field ({})", toString(field)));
}
}
if (type == RGL_RETURN_TYPE_LAST) {
switch (field) {
case XYZ_VEC3_F32: return mrLast.xyz;
case DISTANCE_F32: return mrLast.distance;
case IS_HIT_I32: return mrLast.isHit;
default: throw InvalidPipeline(fmt::format("Multi-Return not supported for this field ({})", toString(field)));
}
}
throw InvalidPipeline(fmt::format("Unknown multi-return type ({})", type));
}

void RaytraceNode::setFields(const std::set<rgl_field_t>& fields)
{
auto keyViewer = std::views::keys(fieldData);
Expand Down
2 changes: 2 additions & 0 deletions src/tape/TapeCore.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,6 +68,7 @@ class TapeCore
static void tape_node_gaussian_noise_angular_ray(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_gaussian_noise_angular_hitpoint(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_gaussian_noise_distance(const YAML::Node& yamlNode, PlaybackState& state);
static void tape_node_multi_return_switch(const YAML::Node& yamlNode, PlaybackState& state);

// Called once in the translation unit
static inline bool autoExtendTapeFunctions = std::invoke([]() {
Expand Down Expand Up @@ -123,6 +124,7 @@ class TapeCore
TAPE_CALL_MAPPING("rgl_node_gaussian_noise_angular_ray", TapeCore::tape_node_gaussian_noise_angular_ray),
TAPE_CALL_MAPPING("rgl_node_gaussian_noise_angular_hitpoint", TapeCore::tape_node_gaussian_noise_angular_hitpoint),
TAPE_CALL_MAPPING("rgl_node_gaussian_noise_distance", TapeCore::tape_node_gaussian_noise_distance),
TAPE_CALL_MAPPING("rgl_node_multi_return_switch", TapeCore::tape_node_multi_return_switch),
};
TapePlayer::extendTapeFunctions(tapeFunctions);
return true;
Expand Down
Loading