HIP: Replace usage of depricated preprocessor macro __AMDGCN_WAVEFRONT_SIZE__ (#14183)

IMbackK · web-flow · commit 2c2caa444341 · 2025-06-15T15:45:27.000+02:00
diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh
@@ -262,11 +262,11 @@ static bool cp_async_available(const int cc) {
 }
 
 static constexpr __device__ int ggml_cuda_get_physical_warp_size() {
-#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
-    return __AMDGCN_WAVEFRONT_SIZE;
+#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) || defined(__GFX8__))
+    return 64;
 #else
     return 32;
-#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
+#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) || defined(__GFX8__))
 }
 
 [[noreturn]]

Original file line number	Diff line number	Diff line change
`@@ -262,11 +262,11 @@ static bool cp_async_available(const int cc) {`
`262`	`262`	`}`
`263`	`263`
`264`	`264`	`static constexpr __device__ int ggml_cuda_get_physical_warp_size() {`
`265`		`-#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)`
`266`		`- return __AMDGCN_WAVEFRONT_SIZE;`
	`265`	`+#if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) \|\| defined(__GFX8__))`
	`266`	`+ return 64;`
`267`	`267`	`#else`
`268`	`268`	`return 32;`
`269`		`-#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)`
	`269`	`+#endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) \|\| defined(__GFX8__))`
`270`	`270`	`}`
`271`	`271`
`272`	`272`	`[[noreturn]]`