From d2fa393a6c7d524b6e81a99f2e570727ffa1ce6f Mon Sep 17 00:00:00 2001
From: F1red0g3 <201348850@qq.com>
Date: Wed, 20 May 2026 22:49:06 +0800
Subject: [PATCH 1/2] CUDA 13.x API compatibility

Signed-off-by: Firedoge <201348850@qq.com>
---
 physx/source/compiler/cmakegpu/CMakeLists.txt  |  2 +-
 .../cudamanager/src/CudaContextManager.cpp     |  7 ++++++-
 .../cudamanager/src/CudaKernelWrangler.cpp     | 18 ++++++++++++++++++
 3 files changed, 25 insertions(+), 2 deletions(-)

diff --git a/physx/source/compiler/cmakegpu/CMakeLists.txt b/physx/source/compiler/cmakegpu/CMakeLists.txt
index 625102ddd2..1bdf0afbc3 100644
--- a/physx/source/compiler/cmakegpu/CMakeLists.txt
+++ b/physx/source/compiler/cmakegpu/CMakeLists.txt
@@ -88,7 +88,7 @@ IF(PX_GENERATE_GPU_REDUCED_ARCHITECTURES)
 	GENERATE_ARCH_CODE_LIST(SASS "80,86,89,90,100,120" PTX "120")
 ELSE()
 	# Volta is the minimum compute arch required because NGC supports V100
-	GENERATE_ARCH_CODE_LIST(SASS "70,80,86,89,90,100,120" PTX "120")
+	GENERATE_ARCH_CODE_LIST(SASS "80,86,89,90,100,120" PTX "120")
 ENDIF()
 
 # Force response files off because clangd does not parse them
diff --git a/physx/source/cudamanager/src/CudaContextManager.cpp b/physx/source/cudamanager/src/CudaContextManager.cpp
index a6a511c674..56c3b3c205 100644
--- a/physx/source/cudamanager/src/CudaContextManager.cpp
+++ b/physx/source/cudamanager/src/CudaContextManager.cpp
@@ -574,7 +574,12 @@ CudaCtxMgr::CudaCtxMgr(const PxCudaContextManagerDesc& desc, PxErrorCallback& er
 				return;
 			}
 			
-			status = cuCtxCreate(&mCtx, (unsigned int)flags, mDevHandle);
+			#if CUDA_VERSION >= 13000
+				status = cuCtxCreate(&mCtx, nullptr, (unsigned int)flags, mDevHandle);
+			#else
+				status = cuCtxCreate(&mCtx, (unsigned int)flags, mDevHandle);
+			#endif
+			
 			if (CUDA_SUCCESS != status)
 			{
 				const size_t bufferSize = 128;
diff --git a/physx/source/cudamanager/src/CudaKernelWrangler.cpp b/physx/source/cudamanager/src/CudaKernelWrangler.cpp
index 05c49deb43..7812c910c9 100644
--- a/physx/source/cudamanager/src/CudaKernelWrangler.cpp
+++ b/physx/source/cudamanager/src/CudaKernelWrangler.cpp
@@ -189,6 +189,24 @@ void CUDARTAPI __cudaRegisterFunction(void** fatCubinHandle, const char*,
 
 /* These functions are implemented just to resolve link dependencies */
 
+extern "C" 
+cudaError_t CUDARTAPI __cudaLaunchKernel(
+    const void*,
+    dim3,
+    dim3,
+    void**,
+    size_t,
+    cudaStream_t)
+{
+    return cudaSuccess;
+}
+
+extern "C" 
+const void* CUDARTAPI __cudaGetKernel(const void* func)
+{
+    return func;
+}
+
 extern "C"
 cudaError_t CUDARTAPI cudaLaunch(const char* entry)
 {

From 126212f0ab3687ea121ae4c23f7471106b1e3586 Mon Sep 17 00:00:00 2001
From: F1red0g3 <201348850@qq.com>
Date: Wed, 20 May 2026 22:49:34 +0800
Subject: [PATCH 2/2] Adjust GPU architecture settings for CUDA versions

Signed-off-by: Firedoge <201348850@qq.com>
---
 physx/source/compiler/cmakegpu/CMakeLists.txt | 13 +++++++++----
 1 file changed, 9 insertions(+), 4 deletions(-)

diff --git a/physx/source/compiler/cmakegpu/CMakeLists.txt b/physx/source/compiler/cmakegpu/CMakeLists.txt
index 1bdf0afbc3..6d7c72e370 100644
--- a/physx/source/compiler/cmakegpu/CMakeLists.txt
+++ b/physx/source/compiler/cmakegpu/CMakeLists.txt
@@ -85,10 +85,15 @@ SET(SOURCE_DISTRO_FILE_LIST "")
 # since that's the default set in: compiler\internal\CMakeLists.txt (using CMAKE_CUDA_ARCHITECTURES)
 # see policy CMP0104
 IF(PX_GENERATE_GPU_REDUCED_ARCHITECTURES)
-	GENERATE_ARCH_CODE_LIST(SASS "80,86,89,90,100,120" PTX "120")
-ELSE()
-	# Volta is the minimum compute arch required because NGC supports V100
-	GENERATE_ARCH_CODE_LIST(SASS "80,86,89,90,100,120" PTX "120")
+	IF(CMAKE_CUDA_COMPILER_VERSION VERSION_LESS 13.0)
+		# CUDA 12.x still supports offline compilation for Volta / sm_70.
+		SET(PHYSX_GPU_SASS_ARCHS "70,80,86,89,90,100,120")
+	ELSE()
+		# CUDA 13.x removed offline compilation support for pre-Turing architectures.
+		SET(PHYSX_GPU_SASS_ARCHS "80,86,89,90,100,120")
+	ENDIF()
+
+GENERATE_ARCH_CODE_LIST(SASS "${PHYSX_GPU_SASS_ARCHS}" PTX "120")
 ENDIF()
 
 # Force response files off because clangd does not parse them