[clang] [Clang][HIP] Deprecate the AMDGCN_WAVEFRONT_SIZE macros (PR #112849)

Joseph Huber via cfe-commits cfe-commits at lists.llvm.org
Mon Oct 21 04:30:01 PDT 2024


================
@@ -337,9 +337,12 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
   if (hasFastFMA())
     Builder.defineMacro("FP_FAST_FMA");
 
-  Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE__", Twine(WavefrontSize));
-  // ToDo: deprecate this macro for naming consistency.
-  Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE", Twine(WavefrontSize));
+  Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE__", Twine(WavefrontSize),
+                      "compile-time-constant access to the wavefront size will "
+                      "be removed in a future release");
+  Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE", Twine(WavefrontSize),
+                      "compile-time-constant access to the wavefront size will "
+                      "be removed in a future release");
----------------
jhuber6 wrote:

This never should've been defined on the host side for sure. We do this in the OpenMP device runtime.
```c
return llvm::omp::getAMDGPUGridValues<__AMDGCN_WAVEFRONT_SIZE>();
```
That could probably be replaced with
```c
return __builtin_amdgcn_wavefrontsize();
```

https://github.com/llvm/llvm-project/pull/112849


More information about the cfe-commits mailing list