fixes

cuteday · Apr 16, 2024 · c2d02bf · c2d02bf
1 parent dec0beb
commit c2d02bf
Show file tree

Hide file tree

Showing 4 changed files with 30 additions and 17 deletions.
diff --git 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<void *>(ptr.get()));
+		Log(Debug, "thrust_cached_resource::do_deallocate(): ptr == %p", reinterpret_cast<void *>(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
@@ -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
@@ -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
@@ -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