graphdeco-inria · amd-wangfan · Mar 31, 2025
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -11,26 +11,56 @@
 
 cmake_minimum_required(VERSION 3.20)
 
-project(DiffRast LANGUAGES CUDA CXX)
+project(DiffRast LANGUAGES CXX)
+find_package(HIP)
+if(HIP_FOUND)
+	set(HIP_ENABLED ON)
+	project(DiffRast LANGUAGES HIP)
+else()
+	project(DiffRast LANGUAGES CUDA)
+endif()
+
 
 set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_EXTENSIONS OFF)
-set(CMAKE_CUDA_STANDARD 17)
-
-set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
 
-add_library(CudaRasterizer
-	cuda_rasterizer/backward.h
-	cuda_rasterizer/backward.cu
-	cuda_rasterizer/forward.h
-	cuda_rasterizer/forward.cu
-	cuda_rasterizer/auxiliary.h
-	cuda_rasterizer/rasterizer_impl.cu
-	cuda_rasterizer/rasterizer_impl.h
-	cuda_rasterizer/rasterizer.h
-)
+if(HIP_ENABLED)
+	set(CMAKE_HIP_STANDARD 17)
+	message(STATUS "Building with HIP support")
+else()
+	set(CMAKE_CUDA_STANDARD 17)
+	message(STATUS "Building with CUDA support")
+endif()
 
-set_target_properties(CudaRasterizer PROPERTIES CUDA_ARCHITECTURES "70;75;86")
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
 
-target_include_directories(CudaRasterizer PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/cuda_rasterizer)
-target_include_directories(CudaRasterizer PRIVATE third_party/glm ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
+if(HIP_ENABLED)
+	add_library(HipRasterizer
+		hip_rasterizer/backward.h
+		hip_rasterizer/backward.hip
+		hip_rasterizer/forward.h
+		hip_rasterizer/forward.hip
+		hip_rasterizer/auxiliary.h
+		hip_rasterizer/rasterizer_impl.hip
+		hip_rasterizer/rasterizer_impl.h
+		hip_rasterizer/rasterizer.h
+	)
+	set_target_properties(HipRasterizer PROPERTIES HIP_ARCHITECTURES "gfx908;gfx90a;gfx942")
+	target_include_directories(HipRasterizer PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/hip_rasterizer)
+    target_include_directories(HipRasterizer PRIVATE ${CMAKE_HIP_TOOLKIT_INCLUDE_DIRECTORIES})
+	target_include_directories(HipRasterizer PRIVATE third_party/glm)
+else()
+	add_library(CudaRasterizer
+		cuda_rasterizer/backward.h
+		cuda_rasterizer/backward.cu
+		cuda_rasterizer/forward.h
+		cuda_rasterizer/forward.cu
+		cuda_rasterizer/auxiliary.h
+		cuda_rasterizer/rasterizer_impl.cu
+		cuda_rasterizer/rasterizer_impl.h
+		cuda_rasterizer/rasterizer.h
+	)
+	set_target_properties(CudaRasterizer PROPERTIES CUDA_ARCHITECTURES "70;75;86")
+	target_include_directories(CudaRasterizer PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/cuda_rasterizer)
+    target_include_directories(CudaRasterizer PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
+endif()
diff --git a/README.md b/README.md
@@ -2,6 +2,23 @@
 
 Used as the rasterization engine for the paper "3D Gaussian Splatting for Real-Time Rendering of Radiance Fields". If you can make use of it in your own research, please be so kind to cite us.
 
+## Install:
+```
+git clone https://github.com/graphdeco-inria/diff-gaussian-rasterization.git
+git submodule update --init --recursive
+cd diff-gaussian-rasterization
+pip insatll -e .
+
+# or
+git clone --recursive https://github.com/graphdeco-inria/diff-gaussian-rasterization.git
+cd diff-gaussian-rasterization
+pip insatll -e .
+```
+## Test:
+```
+python tests/test_forward.py
+```
+
 <section class="section" id="BibTeX">
   <div class="container is-max-desktop content">
     <h2 class="title">BibTeX</h2>
@@ -16,4 +33,4 @@ Used as the rasterization engine for the paper "3D Gaussian Splatting for Real-T
       url          = {https://repo-sam.inria.fr/fungraph/3d-gaussian-splatting/}
 }</code></pre>
   </div>
-</section>
+</section>
diff --git a/cuda_rasterizer/auxiliary.h b/cuda_rasterizer/auxiliary.h
@@ -18,6 +18,13 @@
 #define BLOCK_SIZE (BLOCK_X * BLOCK_Y)
 #define NUM_WARPS (BLOCK_SIZE/32)
 
+#if defined(USE_ROCM) || defined(__HIP_PLATFORM_AMD__)
+  #include <assert.h>
+  #define DEVICE_TRAP() assert(false)
+#else
+  __device__ __forceinline__ void DEVICE_TRAP() { __trap(); }
+#endif
+
 // Spherical harmonics coefficients
 __device__ const float SH_C0 = 0.28209479177387814f;
 __device__ const float SH_C1 = 0.4886025119029199f;
@@ -156,7 +163,7 @@ __forceinline__ __device__ bool in_frustum(int idx,
 		if (prefiltered)
 		{
 			printf("Point is filtered although prefiltered is set. This shouldn't happen!");
-			__trap();
+			DEVICE_TRAP();
 		}
 		return false;
 	}

diff --git a/cuda_rasterizer/backward.cu b/cuda_rasterizer/backward.cu
@@ -12,7 +12,6 @@
 #include "backward.h"
 #include "auxiliary.h"
 #include <cooperative_groups.h>
-#include <cooperative_groups/reduce.h>
 namespace cg = cooperative_groups;
 
 // Backward pass for conversion of spherical harmonics to RGB for
@@ -584,7 +583,7 @@ void BACKWARD::preprocess(
 	// Somewhat long, thus it is its own kernel rather than being part of 
 	// "preprocess". When done, loss gradient w.r.t. 3D means has been
 	// modified and gradient w.r.t. 3D covariance matrix has been computed.	
-	computeCov2DCUDA << <(P + 255) / 256, 256 >> > (
+	computeCov2DCUDA <<<(P + 255) / 256, 256 >>> (
 		P,
 		means3D,
 		radii,
@@ -601,7 +600,7 @@ void BACKWARD::preprocess(
 	// Propagate gradients for remaining steps: finish 3D mean gradients,
 	// propagate color gradients to SH (if desireD), propagate 3D covariance
 	// matrix gradients to scale and rotation.
-	preprocessCUDA<NUM_CHANNELS> << < (P + 255) / 256, 256 >> > (
+	preprocessCUDA<NUM_CHANNELS> <<< (P + 255) / 256, 256 >>> (
 		P, D, M,
 		(float3*)means3D,
 		radii,
@@ -638,7 +637,7 @@ void BACKWARD::render(
 	float* dL_dopacity,
 	float* dL_dcolors)
 {
-	renderCUDA<NUM_CHANNELS> << <grid, block >> >(
+	renderCUDA<NUM_CHANNELS> <<<grid, block >>>(
 		ranges,
 		point_list,
 		W, H,

diff --git a/cuda_rasterizer/backward.h b/cuda_rasterizer/backward.h
@@ -14,7 +14,6 @@
 
 #include <cuda.h>
 #include "cuda_runtime.h"
-#include "device_launch_parameters.h"
 #define GLM_FORCE_CUDA
 #include <glm/glm.hpp>