feat: min and max compute capabilities for CUDA 12.8 & 13.0

Author: Tom94

CMakeLists.txt

@@ -1,5 +1,5 @@
 # Copyright (c) 2020-2025, NVIDIA CORPORATION.  All rights reserved.
-# 
+#
 # Redistribution and use in source and binary forms, with or without modification, are permitted
 # provided that the following conditions are met:
 #     * Redistributions of source code must retain the above copyright notice, this list of
@@ -10,7 +10,7 @@
 #     * Neither the name of the NVIDIA CORPORATION nor the names of its contributors may be used
 #       to endorse or promote products derived from this software without specific prior written
 #       permission.
-# 
+#
 # THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR
 # IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND
 # FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE
@@ -165,7 +165,9 @@ else()
 	set(LATEST_SUPPORTED_CUDA_ARCHITECTURE 120)
 endif()
 
-if (CUDA_VERSION VERSION_GREATER_EQUAL 12.0)
+if (CUDA_VERSION VERSION_GREATER_EQUAL 13.0)
+	set(EARLIEST_SUPPORTED_CUDA_ARCHITECTURE 75)
+elseif (CUDA_VERSION VERSION_GREATER_EQUAL 12.0)
 	set(EARLIEST_SUPPORTED_CUDA_ARCHITECTURE 50)
 else()
 	set(EARLIEST_SUPPORTED_CUDA_ARCHITECTURE 20)

bindings/torch/setup.py

@@ -14,7 +14,9 @@
 ROOT_DIR = os.path.dirname(os.path.dirname(SCRIPT_DIR))
 
 def min_supported_compute_capability(cuda_version):
-	if cuda_version >= parse_version("12.0"):
+	if cuda_version >= parse_version("13.0"):
+		return 75
+	elif cuda_version >= parse_version("12.0"):
 		return 50
 	else:
 		return 20

src/common_host.cu

@@ -41,10 +41,7 @@
 
 namespace tcnn {
 
-static_assert(
-	__CUDACC_VER_MAJOR__ > 10 || (__CUDACC_VER_MAJOR__ == 10 && __CUDACC_VER_MINOR__ >= 2),
-	"tiny-cuda-nn requires at least CUDA 10.2"
-);
+static_assert(__CUDACC_VER_MAJOR__ > 10 || (__CUDACC_VER_MAJOR__ == 10 && __CUDACC_VER_MINOR__ >= 2), "tiny-cuda-nn requires at least CUDA 10.2");
 
 std::function<void(LogSeverity, const std::string&)> g_log_callback = [](LogSeverity severity, const std::string& msg) {
 	switch (severity) {
@@ -214,9 +211,7 @@ int cuda_device() {
 	return device;
 }
 
-void set_cuda_device(int device) {
-	CUDA_CHECK_THROW(cudaSetDevice(device));
-}
+void set_cuda_device(int device) { CUDA_CHECK_THROW(cudaSetDevice(device)); }
 
 int cuda_device_count() {
 	int device_count;
@@ -244,9 +239,7 @@ const cudaDeviceProp& cuda_get_device_properties(int device) {
 	return cuda_device_properties().at(device);
 }
 
-std::string cuda_device_name(int device) {
-	return cuda_get_device_properties(device).name;
-}
+std::string cuda_device_name(int device) { return cuda_get_device_properties(device).name; }
 
 uint32_t cuda_compute_capability(int device) {
 	const auto& props = cuda_get_device_properties(device);
@@ -261,22 +254,20 @@ uint32_t cuda_max_supported_compute_capability() {
 		return 80;
 	} else if (cuda_version < 11080) {
 		return 86;
-	} else {
+	} else if (cuda_version < 12080) {
 		return 90;
+	} else {
+		return 120;
 	}
 }
 
 uint32_t cuda_supported_compute_capability(int device) {
 	return std::min(cuda_compute_capability(device), cuda_max_supported_compute_capability());
 }
 
-size_t cuda_max_shmem(int device) {
-	return cuda_get_device_properties(device).sharedMemPerBlockOptin;
-}
+size_t cuda_max_shmem(int device) { return cuda_get_device_properties(device).sharedMemPerBlockOptin; }
 
-uint32_t cuda_max_registers(int device) {
-	return (uint32_t)cuda_get_device_properties(device).regsPerBlock;
-}
+uint32_t cuda_max_registers(int device) { return (uint32_t)cuda_get_device_properties(device).regsPerBlock; }
 
 size_t cuda_memory_granularity(int device) {
 	size_t granularity;
@@ -358,4 +349,4 @@ template <> std::string type_to_string<double>() { return "double"; }
 template <> std::string type_to_string<float>() { return "float"; }
 template <> std::string type_to_string<__half>() { return "__half"; }
 
-}
+} // namespace tcnn