From a7973d8f676f9fdec8956df858a3684b0a53fdd0 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Thomas=20M=C3=BCller?= <thomas94@gmx.net>
Date: Sun, 30 Oct 2022 16:45:44 +0100
Subject: [PATCH] Better compatibility with CUDA clang

---
 dependencies/tiny-cuda-nn                     |  2 +-
 include/neural-graphics-primitives/common.h   |  6 ++---
 include/neural-graphics-primitives/nerf.h     |  2 +-
 .../neural-graphics-primitives/nerf_loader.h  |  2 +-
 include/neural-graphics-primitives/sdf.h      |  3 ++-
 include/neural-graphics-primitives/testbed.h  |  4 ++--
 src/nerf_loader.cu                            | 24 +++++++++----------
 src/tinyexr_wrapper.cu                        | 16 ++++++-------
 8 files changed, 30 insertions(+), 29 deletions(-)

diff --git a/dependencies/tiny-cuda-nn b/dependencies/tiny-cuda-nn
index 9acde94..8a64135 160000
--- a/dependencies/tiny-cuda-nn
+++ b/dependencies/tiny-cuda-nn
@@ -1 +1 @@
-Subproject commit 9acde9481f76c8f00fb8f0e75dae9c9f4b892de8
+Subproject commit 8a64135a109c305ca4fcf81f26d9bb15b16356a9
diff --git a/include/neural-graphics-primitives/common.h b/include/neural-graphics-primitives/common.h
index 52e3145..9b5c0cf 100644
--- a/include/neural-graphics-primitives/common.h
+++ b/include/neural-graphics-primitives/common.h
@@ -23,8 +23,8 @@
 // to emit this diagnostic.
 // nlohmann::json produces a comparison with zero in one of its templates,
 // which can also safely be ignored.
-#if defined(__NVCC__)
-#  if defined __NVCC_DIAG_PRAGMA_SUPPORT__
+#ifdef __NVCC__
+#  ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
 #    pragma nv_diag_suppress = esa_on_defaulted_function_ignored
 #    pragma nv_diag_suppress = unsigned_compare_with_zero
 #  else
@@ -185,7 +185,7 @@ struct Lens {
 	float params[7] = {};
 };
 
-#ifdef __NVCC__
+#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__))
 #define NGP_HOST_DEVICE __host__ __device__
 #else
 #define NGP_HOST_DEVICE
diff --git a/include/neural-graphics-primitives/nerf.h b/include/neural-graphics-primitives/nerf.h
index a61538c..27e3237 100644
--- a/include/neural-graphics-primitives/nerf.h
+++ b/include/neural-graphics-primitives/nerf.h
@@ -36,7 +36,7 @@ struct NerfPayload {
 };
 
 struct RaysNerfSoa {
-#ifdef __NVCC__
+#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__))
 	void copy_from_other_async(const RaysNerfSoa& other, cudaStream_t stream) {
 		CUDA_CHECK_THROW(cudaMemcpyAsync(rgba, other.rgba, size * sizeof(Eigen::Array4f), cudaMemcpyDeviceToDevice, stream));
 		CUDA_CHECK_THROW(cudaMemcpyAsync(depth, other.depth, size * sizeof(float), cudaMemcpyDeviceToDevice, stream));
diff --git a/include/neural-graphics-primitives/nerf_loader.h b/include/neural-graphics-primitives/nerf_loader.h
index aa9d50f..d863881 100644
--- a/include/neural-graphics-primitives/nerf_loader.h
+++ b/include/neural-graphics-primitives/nerf_loader.h
@@ -105,7 +105,7 @@ struct NerfDataset {
 		if (from_mitsuba) {
 			result *= -1;
 		} else {
-			result=Eigen::Vector3f(result.y(), result.z(), result.x());
+			result = Eigen::Vector3f(result.y(), result.z(), result.x());
 		}
 		return result;
 	}
diff --git a/include/neural-graphics-primitives/sdf.h b/include/neural-graphics-primitives/sdf.h
index 66e2a0f..ddef63c 100644
--- a/include/neural-graphics-primitives/sdf.h
+++ b/include/neural-graphics-primitives/sdf.h
@@ -28,7 +28,7 @@ struct SdfPayload {
 };
 
 struct RaysSdfSoa {
-#ifdef __NVCC__
+#if defined(__NVCC__) || (defined(__clang__) && defined(__CUDA__))
 	void enlarge(size_t n_elements) {
 		pos.enlarge(n_elements);
 		normal.enlarge(n_elements);
@@ -49,6 +49,7 @@ struct RaysSdfSoa {
 		CUDA_CHECK_THROW(cudaMemcpyAsync(payload.data(), other.payload.data(), n_elements * sizeof(SdfPayload), cudaMemcpyDeviceToDevice, stream));
 	}
 #endif
+
 	tcnn::GPUMemory<Eigen::Vector3f> pos;
 	tcnn::GPUMemory<Eigen::Vector3f> normal;
 	tcnn::GPUMemory<float> distance;
diff --git a/include/neural-graphics-primitives/testbed.h b/include/neural-graphics-primitives/testbed.h
index 7fa6e6c..1503674 100644
--- a/include/neural-graphics-primitives/testbed.h
+++ b/include/neural-graphics-primitives/testbed.h
@@ -33,8 +33,8 @@
 #include <filesystem/path.h>
 
 #ifdef NGP_PYTHON
-#include <pybind11/pybind11.h>
-#include <pybind11/numpy.h>
+#  include <pybind11/pybind11.h>
+#  include <pybind11/numpy.h>
 #endif
 
 #include <thread>
diff --git a/src/nerf_loader.cu b/src/nerf_loader.cu
index deb0e3c..ba21e6b 100644
--- a/src/nerf_loader.cu
+++ b/src/nerf_loader.cu
@@ -33,20 +33,20 @@
 
 #define STB_IMAGE_IMPLEMENTATION
 
-#if defined(__NVCC__)
-#if defined __NVCC_DIAG_PRAGMA_SUPPORT__
-#  pragma nv_diag_suppress 550
-#else
-#  pragma diag_suppress 550
-#endif
+#ifdef __NVCC__
+#  ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
+#    pragma nv_diag_suppress 550
+#  else
+#    pragma diag_suppress 550
+#  endif
 #endif
 #include <stb_image/stb_image.h>
-#if defined(__NVCC__)
-#if defined __NVCC_DIAG_PRAGMA_SUPPORT__
-#  pragma nv_diag_default 550
-#else
-#  pragma diag_default 550
-#endif
+#ifdef __NVCC__
+#  ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
+#    pragma nv_diag_default 550
+#  else
+#    pragma diag_default 550
+#  endif
 #endif
 
 using namespace tcnn;
diff --git a/src/tinyexr_wrapper.cu b/src/tinyexr_wrapper.cu
index 587ba1e..096716d 100644
--- a/src/tinyexr_wrapper.cu
+++ b/src/tinyexr_wrapper.cu
@@ -20,14 +20,14 @@
 
 #include <tiny-cuda-nn/gpu_memory.h>
 
-#if defined(__NVCC__)
-#if defined __NVCC_DIAG_PRAGMA_SUPPORT__
-#  pragma nv_diag_suppress 174
-#  pragma nv_diag_suppress 550
-#else
-#  pragma diag_suppress 174
-#  pragma diag_suppress 550
-#endif
+#ifdef __NVCC__
+#  ifdef __NVCC_DIAG_PRAGMA_SUPPORT__
+#    pragma nv_diag_suppress 174
+#    pragma nv_diag_suppress 550
+#  else
+#    pragma diag_suppress 174
+#    pragma diag_suppress 550
+#  endif
 #endif
 
 #define TINYEXR_IMPLEMENTATION
-- 
GitLab