[AMDGPU] Rename predefined macro __AMDGCN_WAVEFRONT_SIZE

[yaxunl]

Summary

rename it to __AMDGCN_WAVEFRONT_SIZE__ for consistency.

__AMDGCN_WAVEFRONT_SIZE will be deprecated in the future.

[arsenm]

clang/lib/Basic/Targets/AMDGPU.cpp

318

If you're renaming it anyway, might as well go for __AMDGPU?

[yaxunl]

clang/lib/Basic/Targets/AMDGPU.cpp

318

but this will cause inconsistency with __AMDGCN_CUMODE__ and __AMDGCN_UNSAFE_FP_ATOMICS__, and all the other __amdgcn_* macros.

[jdoerfert]

I think this is a step in the right direction. @arsenm ?

[arsenm]

lgtm except for dropping test coverage

clang/test/Driver/hip-macros.hip

20

This shouldn't drop the check for the old one

[yaxunl]

revised by comments

clang/lib/Basic/Targets/AMDGPU.cpp

if (hasFastFMAF())
     Builder.defineMacro("FP_FAST_FMAF");
   if (hasLDEXPF())
     Builder.defineMacro("__HAS_LDEXPF__");
   if (hasFP64())
     Builder.defineMacro("__HAS_FP64__");
   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));
 }
 
 void AMDGPUTargetInfo::setAuxTarget(const TargetInfo *Aux) {
   assert(HalfFormat == Aux->HalfFormat);
   assert(FloatFormat == Aux->FloatFormat);
   assert(DoubleFormat == Aux->DoubleFormat);

clang/test/CodeGenHIP/maybe_undef-attr-verify.hip

 
 // CHECK: define linkonce_odr noundef i32 @_Z11__shfl_synciii(i32 noundef [[TMP1:%.*]], i32 noundef [[TMP2:%.*]], i32 noundef [[TMP3:%.*]])
 
 #define __global__ __attribute__((global))
 #define __device__ __attribute__((device))
 #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(
         -1, __builtin_amdgcn_mbcnt_lo(-1, 0));
 }
 
 __device__
 inline

clang/test/Driver/hip-macros.hip

 // RUN:   --cuda-device-only -nogpuinc -nogpulib \
 // RUN:   -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE32 %s
 // RUN: %clang -E -dM --offload-arch=gfx906 -mno-wavefrontsize64 \
 // RUN:   --cuda-device-only -nogpuinc -nogpulib \
 // RUN:   -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
 // 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
 
 // RUN: %clang -E -dM --offload-arch=gfx906 --cuda-device-only -nogpuinc -nogpulib \
 // RUN:   %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
 // RUN: %clang -E -dM --offload-arch=gfx906 --cuda-device-only -nogpuinc -nogpulib -mcumode \
 // RUN:   %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
 // RUN: %clang -E -dM --offload-arch=gfx906 --cuda-device-only -nogpuinc -nogpulib -mno-cumode \
 // RUN:   %s 2>&1 | FileCheck --check-prefixes=CUMODE-ON,WARN-CUMODE %s