diff --git a/src/core/device/thrust.h b/src/core/device/thrust.h index 81a3423..29029c1 100644 --- a/src/core/device/thrust.h +++ b/src/core/device/thrust.h @@ -33,7 +33,7 @@ class thrust_cached_resource final : public: void release() { - Log(Info, "thrust_cached_resource::release()"); + Log(Debug, "thrust_cached_resource::release()"); // Deallocate all outstanding blocks in both lists. for (typename free_blocks_container::iterator i = free_blocks.begin(); i != free_blocks.end(); ++i) m_upstream->do_deallocate(i->second, i->first.first, i->first.second); @@ -44,17 +44,17 @@ class thrust_cached_resource final : } void_ptr do_allocate(std::size_t bytes, std::size_t alignment) override { - Log(Info, "thrust_cached_resource::do_allocate(): num_bytes == %zu", bytes); + Log(Debug, "thrust_cached_resource::do_allocate(): num_bytes == %zu", bytes); void_ptr result = nullptr; typename free_blocks_container::iterator free_block = free_blocks.find({bytes, alignment}); if (free_block != free_blocks.end()) { - Log(Info, "thrust_cached_resource::do_allocate(): found a free block of %zd bytes", bytes); + Log(Debug, "thrust_cached_resource::do_allocate(): found a free block of %zd bytes", bytes); result = free_block->second; free_blocks.erase(free_block); } else { - Log(Info, "thrust_cached_resource::do_allocate(): allocating new block of %zd bytes", bytes); + Log(Debug, "thrust_cached_resource::do_allocate(): allocating new block of %zd bytes", bytes); result = m_upstream->do_allocate(bytes, alignment); } @@ -63,7 +63,7 @@ class thrust_cached_resource final : } void do_deallocate(void_ptr ptr, std::size_t bytes, std::size_t alignment) override { - Log(Info, "thrust_cached_resource::do_deallocate(): ptr == %p", reinterpret_cast(ptr.get())); + Log(Debug, "thrust_cached_resource::do_deallocate(): ptr == %p", reinterpret_cast(ptr.get())); //typename allocated_blocks_container::iterator iter = allocated_blocks.find(ptr); typename allocated_blocks_container::iterator iter = std::find_if(allocated_blocks.begin(), diff --git a/src/render/spectrum.h b/src/render/spectrum.h index e5cb542..a95b336 100644 --- a/src/render/spectrum.h +++ b/src/render/spectrum.h @@ -466,8 +466,10 @@ KRR_CALLABLE float RGBColorSpace::lum(const SpectrumType& s, const SampledWavele SampledSpectrum Ys = CIE_Y->sample(lambda); SampledSpectrum pdf = lambda.pdf(); return SampledSpectrum(Ys * s).safeDiv(pdf).mean() / CIE_Y_integral; + } else { + static_assert(false, "SpectrumType must be either RGB or SampledSpectrum"); + return 0; } - else static_assert(false, "SpectrumType must be either RGB or SampledSpectrum"); } KRR_CALLABLE float luminance(RGB color) { diff --git a/src/render/wavefront/integrator.cpp b/src/render/wavefront/integrator.cpp index 87d8f05..ba619bc 100644 --- a/src/render/wavefront/integrator.cpp +++ b/src/render/wavefront/integrator.cpp @@ -140,6 +140,7 @@ void WavefrontPathTracer::generateScatterRays(int depth) { } { PROFILE("Sort indices"); + // par_nosync only appears on recently cuda versions... thrust::sort(thrust::cuda::par_nosync(*alloc).on(KRR_DEFAULT_STREAM), keys, keys + maxQueueSize, [] KRR_DEVICE(const ScatterRayKeyIndex &a, const ScatterRayKeyIndex &b) { @@ -151,14 +152,12 @@ void WavefrontPathTracer::generateScatterRays(int depth) { // sorted to auxiliary buffer GPUParallelFor(maxQueueSize, [=] KRR_DEVICE (int index) { if (index >= queue->size()) return; - ScatterRayQueue::GetSetIndirector w = queue->operator[](keys[index].index); - auxBuffer->operator[](index) = w.operator krr::ScatterRayWorkItem(); + auxBuffer->operator[](index) = queue->operator[](keys[index].index); }, KRR_DEFAULT_STREAM); // blit back GPUParallelFor(maxQueueSize, [=] KRR_DEVICE (int index) { if (index >= queue->size()) return; - queue->operator[](index) = - auxBuffer->operator[](index).operator krr::ScatterRayWorkItem(); + queue->operator[](index) = auxBuffer->operator[](index); }, KRR_DEFAULT_STREAM); } } diff --git a/src/util/soac.cpp b/src/util/soac.cpp index eeddf9e..323f450 100644 --- a/src/util/soac.cpp +++ b/src/util/soac.cpp @@ -423,19 +423,31 @@ int main(int argc, char* argv[]) { printf(" }\n"); printf(" KRR_CALLABLE void operator=(const %s &a) {\n", workItemName.c_str()); - printf(" DCHECK(soa != nullptr);\n"); + printf(" DCHECK(soa != nullptr);\n"); for (const auto& member : soa.members) for (int i = 0; i < member.names.size(); ++i) { std::string name = member.names[i]; if (!member.arraySizes[i].empty()) { printf(" for (int c = 0; c < %s; ++c)\n", - member.arraySizes[i].c_str()); + member.arraySizes[i].c_str()); printf(" soa->%s[c][i] = a.%s[c];\n", name.c_str(), - name.c_str()); - } - else - printf(" soa->%s[i] = a.%s;\n", name.c_str(), - name.c_str()); + name.c_str()); + } else + printf(" soa->%s[i] = a.%s;\n", name.c_str(), name.c_str()); + } + printf(" }\n\n"); + printf(" KRR_CALLABLE void operator=(const GetSetIndirector &other) {\n"); + printf(" DCHECK(soa != nullptr && other.soa != nullptr);\n"); + for (const auto& member : soa.members) + for (int i = 0; i < member.names.size(); ++i) { + std::string name = member.names[i]; + if (!member.arraySizes[i].empty()) { + printf(" for (int c = 0; c < %s; ++c)\n", + member.arraySizes[i].c_str()); + printf(" soa->%s[c][i] = other.soa->%s[c][other.i];\n", name.c_str(), + name.c_str()); + } else + printf(" soa->%s[i] = other.soa->%s[other.i];\n", name.c_str(), name.c_str()); } printf(" }\n\n"); printf(" KRR_CALLABLE GetSetIndirector() = default;\n"); // default ctor