[clang] ad96f25 - [AMDGPU] Rename predefined macro __AMDGCN_WAVEFRONT_SIZE

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2023-07-02T11:10:06-04:00
New Revision: ad96f25b9365d65459740142a24db4a339705099

URL: https://github.com/llvm/llvm-project/commit/ad96f25b9365d65459740142a24db4a339705099
DIFF: https://github.com/llvm/llvm-project/commit/ad96f25b9365d65459740142a24db4a339705099.diff

LOG: [AMDGPU] Rename predefined macro __AMDGCN_WAVEFRONT_SIZE

rename it to __AMDGCN_WAVEFRONT_SIZE__ for consistency.

__AMDGCN_WAVEFRONT_SIZE will be deprecated in the future.

Reviewed by: Matt Arsenault, Johannes Doerfert

Differential Revision: https://reviews.llvm.org/D154207

Added: 
    

Modified: 
    clang/lib/Basic/Targets/AMDGPU.cpp
    clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
    clang/test/Driver/hip-macros.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp
index 7025c7c4840017..f15216ac15a100 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -315,6 +315,8 @@ 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_CUMODE__", Twine(CUMode));
 }

diff  --git a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
index b7ab32ac408f1f..571fba148f5ccc 100644
--- a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
+++ b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
@@ -20,7 +20,7 @@
 #define __maybe_undef __attribute__((maybe_undef))
 #define WARP_SIZE 64
 
-static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE;
+static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE__;
 
 __device__ static inline unsigned int __lane_id() {
     return  __builtin_amdgcn_mbcnt_hi(

diff  --git a/clang/test/Driver/hip-macros.hip b/clang/test/Driver/hip-macros.hip
index 6056e785f2cee8..6289ef91dd04fa 100644
--- a/clang/test/Driver/hip-macros.hip
+++ b/clang/test/Driver/hip-macros.hip
@@ -17,6 +17,8 @@
 // RUN: %clang -E -dM --offload-arch=gfx1010 -mno-wavefrontsize64 \
 // RUN:   --cuda-device-only -nogpuinc -nogpulib \
 // RUN:   -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
+// WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE__ 64
+// WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE__ 32
 // WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE 64
 // WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE 32