Skip to content
Merged
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
68 changes: 64 additions & 4 deletions tsd/src/tsd/algorithms/cpu/outline.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -14,17 +14,27 @@ namespace tsd::algorithms::cpu {

namespace {

inline uint32_t shadePixel(uint32_t c_in)
inline uint32_t shadePixel(uint32_t c_in, uint32_t outlineColor)
{
auto c_in_f = helium::cvt_color_to_float4(c_in);
auto c_h = tsd::math::float4(1.f, 0.5f, 0.f, 1.f);
auto c_h = helium::cvt_color_to_float4(outlineColor);
auto c_out = tsd::math::lerp(c_in_f, c_h, 0.8f);
return helium::cvt_color_to_uint32(c_out);
}

inline bool isValidPrimitive(uint32_t objectId, uint32_t primitiveId)
{
return objectId != ~0u && primitiveId != ~0u;
}

inline uint64_t primitiveKey(uint32_t objectId, uint32_t primitiveId)
{
return (uint64_t(objectId) << 32) | primitiveId;
}

} // namespace

void outline(const uint32_t *objectId,
void outlineObject(const uint32_t *objectId,
uint32_t *color,
uint32_t outlineId,
uint32_t width,
Expand All @@ -46,7 +56,57 @@ void outline(const uint32_t *objectId,
}

if (cnt > 1 && cnt < 8)
color[i] = shadePixel(color[i]);
color[i] = shadePixel(
color[i], helium::cvt_color_to_uint32({1.f, 0.5f, 0.f, 1.f}));
});
}

void outlinePrimitives(const uint32_t *objectId,
const uint32_t *primitiveId,
uint32_t *color,
uint32_t outlineColor,
uint32_t thickness,
uint32_t width,
uint32_t height)
{
const uint32_t radius = std::max(1u, thickness);
detail::parallel_for(0u, width * height, [=](uint32_t i) {
uint32_t y = i / width;
uint32_t x = i % width;

bool sawValid = false;
bool sawInvalid = false;
bool sawDifferentValid = false;
uint64_t firstValidKey = 0;

for (unsigned fy = std::max(0u, y - radius);
fy <= std::min(height - 1, y + radius) && !sawDifferentValid;
fy++) {
for (unsigned fx = std::max(0u, x - radius);
fx <= std::min(width - 1, x + radius);
Comment thread
jeffamstutz marked this conversation as resolved.
fx++) {
const size_t fi = fx + size_t(width) * fy;
const uint32_t neighborObjectId = objectId[fi];
const uint32_t neighborPrimitiveId = primitiveId[fi];

if (!isValidPrimitive(neighborObjectId, neighborPrimitiveId)) {
sawInvalid = true;
continue;
}

const uint64_t key = primitiveKey(neighborObjectId, neighborPrimitiveId);
if (!sawValid) {
sawValid = true;
firstValidKey = key;
} else if (key != firstValidKey) {
sawDifferentValid = true;
break;
}
}
}

if (sawValid && (sawInvalid || sawDifferentValid))
color[i] = shadePixel(color[i], outlineColor);
});
}

Expand Down
13 changes: 12 additions & 1 deletion tsd/src/tsd/algorithms/cpu/outline.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -9,10 +9,21 @@ namespace tsd::algorithms::cpu {

// Highlight edges of a selected object by blending an orange tint
// on pixels where the 3x3 neighborhood straddles the object boundary.
void outline(const uint32_t *objectId,
void outlineObject(const uint32_t *objectId,
uint32_t *color,
uint32_t outlineId,
uint32_t width,
uint32_t height);

// Highlight edges between neighboring primitive identities, where identity is
// defined by the pair (objectId, primitiveId). Background transitions are
// treated as silhouettes.
void outlinePrimitives(const uint32_t *objectId,
const uint32_t *primitiveId,
uint32_t *color,
uint32_t outlineColor,
uint32_t thickness,
uint32_t width,
uint32_t height);

} // namespace tsd::algorithms::cpu
105 changes: 97 additions & 8 deletions tsd/src/tsd/algorithms/cuda/outline.cu
Original file line number Diff line number Diff line change
Expand Up @@ -17,17 +17,30 @@ namespace tsd::algorithms::cuda {

namespace {

__forceinline__ __device__ uint32_t shadePixel(uint32_t c_in)
__forceinline__ __device__ uint32_t shadePixel(
uint32_t c_in, uint32_t outlineColor)
{
auto c_in_f = helium::cvt_color_to_float4(c_in);
auto c_h = tsd::math::float4(1.f, 0.5f, 0.f, 1.f);
auto c_h = helium::cvt_color_to_float4(outlineColor);
auto c_out = tsd::math::lerp(c_in_f, c_h, 0.8f);
return helium::cvt_color_to_uint32(c_out);
}

__forceinline__ __device__ bool isValidPrimitive(
uint32_t objectId, uint32_t primitiveId)
{
return objectId != ~0u && primitiveId != ~0u;
}

__forceinline__ __device__ uint64_t primitiveKey(
uint32_t objectId, uint32_t primitiveId)
{
return (uint64_t(objectId) << 32) | primitiveId;
}

} // namespace

void outline(cudaStream_t stream,
void outlineObject(cudaStream_t stream,
const uint32_t *objectId,
uint32_t *color,
uint32_t outlineId,
Expand All @@ -42,8 +55,11 @@ void outline(cudaStream_t stream,
uint32_t x = i % width;

int cnt = 0;
for (unsigned fy = max(0u, y - 1); fy <= min(height - 1, y + 1); fy++) {
for (unsigned fx = max(0u, x - 1); fx <= min(width - 1, x + 1);
for (unsigned fy = std::max(0u, y - 1);
fy <= std::min(height - 1, y + 1);
fy++) {
for (unsigned fx = std::max(0u, x - 1);
fx <= std::min(width - 1, x + 1);
fx++) {
size_t fi = fx + size_t(width) * fy;
if (objectId[fi] == outlineId)
Expand All @@ -52,17 +68,90 @@ void outline(cudaStream_t stream,
}

if (cnt > 1 && cnt < 8)
color[i] = shadePixel(color[i]);
color[i] = shadePixel(
color[i], helium::cvt_color_to_uint32({1.f, 0.5f, 0.f, 1.f}));
});
}

void outline(const uint32_t *objectId,
void outlineObject(const uint32_t *objectId,
uint32_t *color,
uint32_t outlineId,
uint32_t width,
uint32_t height)
{
outline(cudaStream_t{0}, objectId, color, outlineId, width, height);
outlineObject(cudaStream_t{0}, objectId, color, outlineId, width, height);
}

void outlinePrimitives(cudaStream_t stream,
const uint32_t *objectId,
const uint32_t *primitiveId,
uint32_t *color,
uint32_t outlineColor,
uint32_t thickness,
uint32_t width,
uint32_t height)
{
const uint32_t radius = std::max(1u, thickness);
thrust::for_each(thrust::cuda::par.on(stream),
thrust::make_counting_iterator(0u),
thrust::make_counting_iterator(width * height),
[=] __device__(uint32_t i) {
uint32_t y = i / width;
uint32_t x = i % width;

bool sawValid = false;
bool sawInvalid = false;
bool sawDifferentValid = false;
uint64_t firstValidKey = 0;

for (unsigned fy = std::max(0u, y - radius);
fy <= std::min(height - 1, y + radius) && !sawDifferentValid;
fy++) {
for (unsigned fx = std::max(0u, x - radius);
fx <= std::min(width - 1, x + radius);
fx++) {
const size_t fi = fx + size_t(width) * fy;
const uint32_t neighborObjectId = objectId[fi];
const uint32_t neighborPrimitiveId = primitiveId[fi];

if (!isValidPrimitive(neighborObjectId, neighborPrimitiveId)) {
sawInvalid = true;
continue;
}

const uint64_t key =
primitiveKey(neighborObjectId, neighborPrimitiveId);
if (!sawValid) {
sawValid = true;
firstValidKey = key;
} else if (key != firstValidKey) {
sawDifferentValid = true;
break;
}
}
}

if (sawValid && (sawInvalid || sawDifferentValid))
color[i] = shadePixel(color[i], outlineColor);
});
}

void outlinePrimitives(const uint32_t *objectId,
const uint32_t *primitiveId,
uint32_t *color,
uint32_t outlineColor,
uint32_t thickness,
uint32_t width,
uint32_t height)
{
outlinePrimitives(cudaStream_t{0},
objectId,
primitiveId,
color,
outlineColor,
thickness,
width,
height);
}

} // namespace tsd::algorithms::cuda
21 changes: 19 additions & 2 deletions tsd/src/tsd/algorithms/cuda/outline.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -8,17 +8,34 @@

namespace tsd::algorithms::cuda {

void outline(cudaStream_t stream,
void outlineObject(cudaStream_t stream,
const uint32_t *objectId,
uint32_t *color,
uint32_t outlineId,
uint32_t width,
uint32_t height);

void outline(const uint32_t *objectId,
void outlineObject(const uint32_t *objectId,
uint32_t *color,
uint32_t outlineId,
uint32_t width,
uint32_t height);

void outlinePrimitives(cudaStream_t stream,
const uint32_t *objectId,
const uint32_t *primitiveId,
uint32_t *color,
uint32_t outlineColor,
uint32_t thickness,
uint32_t width,
uint32_t height);

void outlinePrimitives(const uint32_t *objectId,
const uint32_t *primitiveId,
uint32_t *color,
uint32_t outlineColor,
uint32_t thickness,
uint32_t width,
uint32_t height);

} // namespace tsd::algorithms::cuda
2 changes: 2 additions & 0 deletions tsd/src/tsd/rendering/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -19,6 +19,7 @@ PRIVATE
pipeline/passes/CopyToSDLTexturePass.cpp
pipeline/passes/MultiDeviceSceneRenderPass.cpp
pipeline/passes/OutlineRenderPass.cpp
pipeline/passes/PrimitiveOutlineRenderPass.cpp
pipeline/passes/OutputTransformPass.cpp
pipeline/passes/PickPass.cpp
pipeline/passes/ImagePass.cpp
Expand Down Expand Up @@ -47,6 +48,7 @@ if (TSD_USE_CUDA)
pipeline/passes/ClearBuffersPass.cpp
pipeline/passes/ImagePass.cpp
pipeline/passes/OutlineRenderPass.cpp
pipeline/passes/PrimitiveOutlineRenderPass.cpp
pipeline/passes/OutputTransformPass.cpp
pipeline/passes/ToneMapPass.cpp
pipeline/passes/VisualizeAOVPass.cpp
Expand Down
1 change: 1 addition & 0 deletions tsd/src/tsd/rendering/pipeline/ImagePipeline.h
Original file line number Diff line number Diff line change
Expand Up @@ -13,6 +13,7 @@
#include "passes/CopyToColorBufferPass.hpp"
#include "passes/MultiDeviceSceneRenderPass.h"
#include "passes/OutlineRenderPass.h"
#include "passes/PrimitiveOutlineRenderPass.h"
#include "passes/PickPass.h"
#include "passes/SaveToFilePass.h"
#include "passes/VisualizeAOVPass.h"
Expand Down
4 changes: 2 additions & 2 deletions tsd/src/tsd/rendering/pipeline/passes/OutlineRenderPass.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -30,12 +30,12 @@ void OutlineRenderPass::render(ImageBuffers &b, int stageId)

#ifdef TSD_ALGORITHMS_HAS_CUDA
if (b.stream) {
tsd::algorithms::cuda::outline(
tsd::algorithms::cuda::outlineObject(
b.stream, b.objectId, b.color, m_outlineId, size.x, size.y);
return;
}
#endif
tsd::algorithms::cpu::outline(
tsd::algorithms::cpu::outlineObject(
b.objectId, b.color, m_outlineId, size.x, size.y);
}

Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,60 @@
// Copyright 2026 NVIDIA Corporation
// SPDX-License-Identifier: Apache-2.0

#include "PrimitiveOutlineRenderPass.h"
// tsd_algorithms
#include "tsd/algorithms/cpu/outline.hpp"
#ifdef TSD_ALGORITHMS_HAS_CUDA
#include "tsd/algorithms/cuda/outline.hpp"
#endif
// helium
#include <helium/helium_math.h>

namespace tsd::rendering {

PrimitiveOutlineRenderPass::PrimitiveOutlineRenderPass() = default;

PrimitiveOutlineRenderPass::~PrimitiveOutlineRenderPass() = default;

void PrimitiveOutlineRenderPass::setOutlineColor(
const tsd::math::float4 &color)
{
m_outlineColor = color;
}

void PrimitiveOutlineRenderPass::setThickness(uint32_t thickness)
{
m_thickness = thickness;
}

void PrimitiveOutlineRenderPass::render(ImageBuffers &b, int stageId)
{
if (!b.objectId || !b.primitiveId || stageId == 0)
return;

const auto size = getDimensions();
const auto outlineColor = helium::cvt_color_to_uint32(m_outlineColor);

#ifdef TSD_ALGORITHMS_HAS_CUDA
if (b.stream) {
tsd::algorithms::cuda::outlinePrimitives(b.stream,
b.objectId,
b.primitiveId,
b.color,
outlineColor,
m_thickness,
size.x,
size.y);
return;
}
#endif
tsd::algorithms::cpu::outlinePrimitives(b.objectId,
b.primitiveId,
b.color,
outlineColor,
m_thickness,
size.x,
size.y);
}

} // namespace tsd::rendering
Loading
Loading