[ROCm] Add an AMD GPU build with HIP

.gitignore

@@ -12,3 +12,6 @@ x86/
 
 # Visual Studio 2015/2017 cache/options directory
 .vs/
+
+# CMake build directory (ROCm port)
+build/

CMakeLists.txt

@@ -0,0 +1,107 @@
+cmake_minimum_required(VERSION 3.24)
+
+project(Velvet LANGUAGES C CXX)
+
+option(USE_HIP "Build with HIP for AMD GPUs" OFF)
+
+# GPU language setup
+if(USE_HIP)
+  enable_language(HIP)
+  if(NOT DEFINED CMAKE_HIP_ARCHITECTURES OR CMAKE_HIP_ARCHITECTURES STREQUAL "")
+    set(CMAKE_HIP_ARCHITECTURES "gfx90a")
+  endif()
+  message(STATUS "Building with HIP, architectures: ${CMAKE_HIP_ARCHITECTURES}")
+else()
+  enable_language(CUDA)
+  message(STATUS "Building with CUDA")
+endif()
+
+set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+
+# Find required packages (all provided via vcpkg)
+find_package(OpenGL REQUIRED)
+find_package(glfw3 CONFIG REQUIRED)
+find_package(glad CONFIG REQUIRED)
+find_package(glm CONFIG REQUIRED)
+find_package(fmt CONFIG REQUIRED)
+find_package(assimp CONFIG REQUIRED)
+find_package(imgui CONFIG REQUIRED)
+
+# HIP-specific dependencies
+if(USE_HIP)
+  find_package(hipcub REQUIRED)
+  find_package(rocthrust REQUIRED)
+endif()
+
+# Host/GPU sources. Upstream keeps these as .cpp (built by MSVC) plus the two
+# .cu files (built by nvcc) in Velvet.vcxproj; that layout is preserved here.
+# The .cpp files include device code through Common.cuh, so the HIP build below
+# compiles them as HIP via LANGUAGE without renaming, leaving the Visual Studio
+# CUDA build (which references the .cpp names) intact.
+set(CPP_SOURCES
+  Velvet/Actor.cpp
+  Velvet/Component.cpp
+  Velvet/GUI.cpp
+  Velvet/GameInstance.cpp
+  Velvet/Helper.cpp
+  Velvet/Input.cpp
+  Velvet/MeshRenderer.cpp
+  Velvet/Timer.cpp
+  Velvet/VtEngine.cpp
+  Velvet/stb_image.cpp
+  Velvet/main.cpp
+)
+set(CU_SOURCES
+  Velvet/VtClothSolverGPU.cu
+  Velvet/SpatialHashGPU.cu
+)
+set(SOURCES ${CPP_SOURCES} ${CU_SOURCES})
+
+# Create executable
+add_executable(Velvet ${SOURCES})
+
+if(USE_HIP)
+  # Compile every source as HIP. A mixed CXX/HIP target would apply the HIP
+  # arch flags to the .cpp sources as plain C++ and fail, so mark them all HIP.
+  set_source_files_properties(${SOURCES} PROPERTIES LANGUAGE HIP)
+  target_compile_definitions(Velvet PRIVATE USE_HIP)
+  target_link_libraries(Velvet PRIVATE hip::hipcub roc::rocthrust)
+  set_target_properties(Velvet PROPERTIES
+    HIP_ARCHITECTURES "${CMAKE_HIP_ARCHITECTURES}"
+  )
+endif()
+# For the non-HIP (CUDA) build, language auto-detection drives the split:
+# .cpp -> CXX, .cu -> CUDA, matching upstream's Visual Studio project.
+
+# Include directories
+target_include_directories(Velvet PRIVATE
+  ${CMAKE_CURRENT_SOURCE_DIR}/Velvet
+  ${CMAKE_CURRENT_SOURCE_DIR}/Velvet/External
+  ${CMAKE_CURRENT_SOURCE_DIR}/Velvet/External/cuda
+)
+
+# Link libraries
+target_link_libraries(Velvet PRIVATE
+  OpenGL::GL
+  glfw
+  glad::glad
+  glm::glm
+  fmt::fmt
+  assimp::assimp
+  imgui::imgui
+)
+
+# Set output directory
+set_target_properties(Velvet PROPERTIES
+  RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/bin"
+)
+
+# Copy assets to build directory
+if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/Velvet/Assets")
+  add_custom_command(TARGET Velvet POST_BUILD
+    COMMAND ${CMAKE_COMMAND} -E copy_directory
+    "${CMAKE_CURRENT_SOURCE_DIR}/Velvet/Assets"
+    "$<TARGET_FILE_DIR:Velvet>/Assets"
+  )
+endif()

README.md

@@ -53,6 +53,24 @@ If you want to build from source by yourself, dependencies can be installed usin
 ./vcpkg.exe install imgui[core, opengl3-binding, glfw-binding]:x64-windows
 ```
 
+### Building for AMD GPUs (ROCm/HIP)
+
+Velvet also builds for AMD GPUs with ROCm/HIP through CMake. The same dependencies are available from vcpkg using the `x64-linux` triplet:
+
+```bash
+./vcpkg install glfw3 glad fmt glm assimp "imgui[core,opengl3-binding,glfw-binding]" --triplet x64-linux
+```
+
+Then configure with `USE_HIP=ON`, selecting your GPU architecture with `CMAKE_HIP_ARCHITECTURES` (for example `gfx1100` or `gfx1201`):
+
+```bash
+cmake -B build -DUSE_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1100 \
+      -DCMAKE_TOOLCHAIN_FILE=<path-to-vcpkg>/scripts/buildsystems/vcpkg.cmake
+cmake --build build -j
+```
+
+Velvet renders with OpenGL and shares buffers with the GPU, so it needs an AMD GPU with a graphics pipeline (RDNA, such as the Radeon RX 7000/9000 series). Compute-only datacenter GPUs (CDNA, such as the Instinct MI series) cannot create the OpenGL context Velvet requires.
+
 ## Implementation Details
 
 In computer graphics, building your own wheel can often be unevitable. But what fears most is that sometimes you don't even have recipe for the wheel you want to build. There are lots of great paper describing their methods, but many of the implementation details are left out or scattered across the internet.

Velvet/Common.cuh

@@ -1,16 +1,27 @@
 #pragma once
 
 #include <tuple>
+#include <cstring>
+#include <cstdlib>
+
+// cuda_to_hip.h MUST come before glm so GLM_FORCE_CUDA et al. are defined
+#include "cuda_to_hip.h"
 
 #include <fmt/format.h>
 #include <glm/glm.hpp>
 
-#include <cuda_runtime.h> 
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+// HIP path: hip_runtime.h is included via cuda_to_hip.h
+#else
+// CUDA path
+#include <cuda_runtime.h>
 #include <cuda_runtime_api.h>
 #include <device_launch_parameters.h>
 #include <helper_cuda.h>
 #include <helper_math.h>
+#endif
 
+// cuda_to_hip.h already defines THRUST_DEVICE_COMPILER=5 for HIP backend
 #include <thrust/device_ptr.h>
 #include <thrust/transform.h>
 #include <thrust/sort.h>
@@ -20,7 +31,7 @@
 #define GET_CUDA_ID_NO_RETURN(id, maxID) 	uint id = blockIdx.x * blockDim.x + threadIdx.x
 #define EPSILON					1e-6f
 
-#ifdef __CUDACC__ 
+#if defined(__CUDACC__) || defined(__HIPCC__)
 #define CUDA_CALL(func, totalThreads)  \
 	if (totalThreads == 0) return; \
 	uint func ## _numBlocks, func ## _numThreads; \

Velvet/Common.hpp

@@ -7,11 +7,12 @@
 
 #define IMGUI_LEFT_LABEL(func, label, ...) (ImGui::TextUnformatted(label), ImGui::SameLine(), func("##" label, __VA_ARGS__))
 
-// Only initialize value on host. 
-// Since CUDA doesn't allow dynamics initialization, 
-// we use this macro to ignore initialization when compiling with NVCC.
-#ifdef __CUDA_ARCH__
-	#define HOST_INIT(val) 
+// Only initialize value on host.
+// Since CUDA/HIP does not allow dynamic initialization for __device__ variables,
+// suppress default member initializers when compiling with NVCC or hipcc.
+// The struct is used both as a host variable and as __device__ __constant__.
+#if defined(__CUDACC__) || defined(__HIPCC__)
+	#define HOST_INIT(val)
 #else
 	#define HOST_INIT(val) = val
 #endif
@@ -87,12 +88,15 @@ class VtCallback
 		m_funcs.push_back(func);
 	}
 
-	template <class... TArgs>
-	void Invoke(TArgs... args)
+	// Named Args, not TArgs: VtCallback already has a TArgs parameter pack, and
+	// reusing that name here shadows it, which nvcc/gcc reject ("shadows
+	// template parameter"). Keep this name distinct.
+	template <class... Args>
+	void Invoke(Args... args)
 	{
 		for (const auto& func : m_funcs)
 		{
-			func(std::forward<TArgs>(args)...);
+			func(std::forward<Args>(args)...);
 		}
 	}
 

Velvet/Component.hpp

@@ -1,6 +1,7 @@
 #pragma once
 
 #include <iostream>
+#include <memory>
 #include <string>
 
 #include "Transform.hpp"

Velvet/GameInstance.hpp

@@ -11,6 +11,7 @@
 
 #include "Component.hpp"
 #include "Common.hpp"
+#include "Actor.hpp"
 
 namespace Velvet
 {
@@ -19,7 +20,6 @@ namespace Velvet
 	class Light;
 	class RenderPipeline;
 	class GUI;
-	class Actor;
 	class Timer;
 
 	class GameInstance

Velvet/Helper.cpp

@@ -1,5 +1,5 @@
 #include "Helper.hpp"
-#include <glm\ext\matrix_transform.hpp>
+#include <glm/ext/matrix_transform.hpp>
 
 namespace Velvet
 {

Velvet/Helper.hpp

@@ -1,20 +1,20 @@
 #pragma once
 
+#include <algorithm>
 #include <fmt/format.h>
-//#include <GLFW/glfw3.h>
 #include <glm/glm.hpp>
 
 template <>
 struct fmt::formatter<glm::vec3> : fmt::formatter<std::string> {
-	auto format(glm::vec3 p, format_context& ctx) {
+	auto format(glm::vec3 p, format_context& ctx) const {
 		return formatter<std::string>::format(
 			fmt::format("[{:.2f}, {:.2f}, {:.2f}]", p.x, p.y, p.z), ctx);
 	}
 };
 
 template <>
 struct fmt::formatter<glm::vec2> : fmt::formatter<std::string> {
-	auto format(glm::vec2 p, format_context& ctx) {
+	auto format(glm::vec2 p, format_context& ctx) const {
 		return formatter<std::string>::format(
 			fmt::format("[{:.2f}, {:.2f}]", p.x, p.y), ctx);
 	}
@@ -35,7 +35,7 @@ namespace Velvet
 		template <class T>
 		T Lerp(T value1, T value2, float a)
 		{
-			a = min(max(a, 0.0f), 1.0f);
+			a = std::min(std::max(a, 0.0f), 1.0f);
 			return a * value2 + (1 - a) * value1;
 		}
 	}

Velvet/SpatialHashGPU.cu

@@ -1,6 +1,11 @@
 #include "SpatialHashGPU.cuh"
 
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+#include <hipcub/hipcub.hpp>
+namespace cub = hipcub;
+#else
 #include <cub/device/device_radix_sort.cuh>
+#endif
 
 #include "Timer.hpp"
 #include "VtBuffer.hpp"

Velvet/SpatialHashGPU.hpp

@@ -3,10 +3,12 @@
 #include <unordered_set>
 
 #include <glm/glm.hpp>
+#include <fmt/ranges.h>
 
+#include "cuda_to_hip.h"
 #include "VtBuffer.hpp"
 #include "Global.hpp"
-#include "SpatialhashGPU.cuh"
+#include "SpatialHashGPU.cuh"
 
 using namespace std;
 

Velvet/Timer.hpp

@@ -7,7 +7,12 @@
 #include <glad/glad.h>
 #include <GLFW/glfw3.h>
 #include <fmt/printf.h>
+
+#include "cuda_to_hip.h"
+
+#if !defined(USE_HIP) && !defined(__HIP_PLATFORM_AMD__)
 #include <cuda_runtime.h>
+#endif
 
 //#include "Global.hpp"
 

Velvet/VtBuffer.hpp

@@ -182,7 +182,11 @@ namespace Velvet
 		size_t m_numBytes = 0;
 		T* m_buffer = nullptr;
 		T* m_bufferCPU = nullptr;
+#if defined(USE_HIP) || defined(__HIP_PLATFORM_AMD__)
+		hipGraphicsResource* m_cudaVboResource = nullptr;
+#else
 		struct cudaGraphicsResource* m_cudaVboResource = nullptr;
+#endif
 	};
 
 	template <class T>

Velvet/VtClothSolverGPU.hpp

@@ -4,12 +4,18 @@
 
 #include <glad/glad.h>
 #include <glm/glm.hpp>
+
+#include "cuda_to_hip.h"
+
+#if !defined(USE_HIP) && !defined(__HIP_PLATFORM_AMD__)
 #include <cuda_runtime.h>
 #include <cuda_gl_interop.h>
+#include "helper_cuda.h"
+#endif
+
 #include <thrust/device_ptr.h>
 #include <thrust/transform.h>
 
-#include "helper_cuda.h"
 #include "Mesh.hpp"
 #include "VtClothSolverGPU.cuh"
 #include "VtBuffer.hpp"

Velvet/VtEngine.cpp

@@ -89,9 +89,9 @@ int VtEngine::Run()
 #pragma warning( push )
 #pragma warning( disable : 4129)
 		fmt::print(
-			"��{0:\-^{2}}��\n"
-			"��{1: ^{2}}��\n"
-			"��{0:\-^{2}}��\n", "", "Hello, Velvet!", 30);
+			"��{0:\-^{2}}��\n"
+			"��{1: ^{2}}��\n"
+			"��{0:\-^{2}}��\n", "", "Hello, Velvet!", 30);
 #pragma warning( pop ) 
 
 		m_game = make_shared<GameInstance>(m_window, m_gui);

Velvet/VtEngine.hpp

@@ -1,6 +1,7 @@
 #pragma once
 
 #include <iostream>
+#include <memory>
 #include <string>
 #include <vector>