From 5846e47c12e6c5a4b35139a17e947299ca56cf96 Mon Sep 17 00:00:00 2001 From: Tobias Reiter Date: Wed, 7 Jan 2026 19:06:17 +0100 Subject: [PATCH 1/2] Fix material handling in GPU tracer --- include/viennaray/gpu/raygTrace.hpp | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/include/viennaray/gpu/raygTrace.hpp b/include/viennaray/gpu/raygTrace.hpp index 5fb345b..65a1a4c 100644 --- a/include/viennaray/gpu/raygTrace.hpp +++ b/include/viennaray/gpu/raygTrace.hpp @@ -243,14 +243,20 @@ template class Trace { template void setMaterialIds(const std::vector &materialIds, - const bool mapToConsecutive = true) { + const bool mapToConsecutive = true, + const std::set &pUniqueMaterialIds = {}) { assert(materialIds.size() == launchParams_.numElements); - if (mapToConsecutive) { - uniqueMaterialIds_.clear(); + uniqueMaterialIds_.clear(); + if (!pUniqueMaterialIds.empty()) { + uniqueMaterialIds_ = pUniqueMaterialIds; + } else { for (auto &matId : materialIds) { uniqueMaterialIds_.insert(static_cast(matId)); } + } + + if (mapToConsecutive) { std::unordered_map materialIdMap; int currentId = 0; for (auto &uniqueMaterialId : uniqueMaterialIds_) { @@ -266,6 +272,7 @@ template class Trace { materialIdsBuffer_.allocUpload(materialIdsMapped); } else { std::vector materialIdsMapped(launchParams_.numElements); +#pragma omp parallel for for (int i = 0; i < launchParams_.numElements; i++) { materialIdsMapped[i] = static_cast(materialIds[i]); } From 56c0163b40bec2427366e7e3f5af29cf340b9c00 Mon Sep 17 00:00:00 2001 From: Tobias Reiter Date: Wed, 7 Jan 2026 20:02:12 +0100 Subject: [PATCH 2/2] Add material map on GPU --- include/viennaray/gpu/raygLaunchParams.hpp | 1 + include/viennaray/gpu/raygTrace.hpp | 7 +++++++ 2 files changed, 8 insertions(+) diff --git a/include/viennaray/gpu/raygLaunchParams.hpp b/include/viennaray/gpu/raygLaunchParams.hpp index 0f0f4d8..666df4b 100644 --- a/include/viennaray/gpu/raygLaunchParams.hpp +++ b/include/viennaray/gpu/raygLaunchParams.hpp @@ -40,6 +40,7 @@ struct LaunchParams { float sticking = 1.f; float cosineExponent = 1.f; int *materialIds; + int *materialMap; float *materialSticking; void *customData; diff --git a/include/viennaray/gpu/raygTrace.hpp b/include/viennaray/gpu/raygTrace.hpp index 65a1a4c..311274c 100644 --- a/include/viennaray/gpu/raygTrace.hpp +++ b/include/viennaray/gpu/raygTrace.hpp @@ -97,6 +97,7 @@ template class Trace { if (materialIdsBuffer_.sizeInBytes != 0) { launchParams_.materialIds = (int *)materialIdsBuffer_.dPointer(); + launchParams_.materialMap = (int *)materialMapBuffer_.dPointer(); } launchParams_.seed = config_.rngSeed + config_.runNumber++; @@ -256,6 +257,10 @@ template class Trace { } } + std::vector materialMap(uniqueMaterialIds_.begin(), + uniqueMaterialIds_.end()); + materialMapBuffer_.allocUpload(materialMap); + if (mapToConsecutive) { std::unordered_map materialIdMap; int currentId = 0; @@ -354,6 +359,7 @@ template class Trace { buffer.free(); } materialIdsBuffer_.free(); + materialMapBuffer_.free(); for (auto &buffer : materialStickingBuffer_) { buffer.free(); } @@ -756,6 +762,7 @@ template class Trace { std::vector> particles_; CudaBuffer dataPerParticleBuffer_; // same for all particles std::vector materialStickingBuffer_; // different for particles + CudaBuffer materialMapBuffer_; // sbt data CudaBuffer cellDataBuffer_;