uvos commited on
Commit
c3467c7
·
1 Parent(s): 6799437

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

Browse files
Files changed (1) hide show
  1. ggml/src/ggml-cuda/common.cuh +3 -3
ggml/src/ggml-cuda/common.cuh CHANGED
@@ -262,11 +262,11 @@ static bool cp_async_available(const int cc) {
262
  }
263
 
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;
267
  #else
268
  return 32;
269
- #endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__)
270
  }
271
 
272
  [[noreturn]]
 
262
  }
263
 
264
  static constexpr __device__ int ggml_cuda_get_physical_warp_size() {
265
+ #if defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) || defined(__GFX8__))
266
+ return 64;
267
  #else
268
  return 32;
269
+ #endif // defined(GGML_USE_HIP) && defined(__HIP_PLATFORM_AMD__) && (defined(__GFX9__) || defined(__GFX8__))
270
  }
271
 
272
  [[noreturn]]