https://github.com/ritter-x2a created
https://github.com/llvm/llvm-project/pull/157463
Remove definitions, test uses, and documentation of the macros, which were
deprecated in November 2024 with PR #112849 / #115507.
Where required, the wavefront size should instead be queried via means provided
by the HIP runtime: the (non-constexpr) `warpSize` variable in device code, or
`hipGetDeviceProperties` in host code.
This change passed AMD-internal testing.
Implements SWDEV-522062.
>From 2b0505ef56b5afaace232020c74dbf67e31424a9 Mon Sep 17 00:00:00 2001
From: Fabian Ritter <fabian.rit...@amd.com>
Date: Mon, 8 Sep 2025 04:56:31 -0400
Subject: [PATCH] [HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros
Remove definitions, test uses, and documentation of the macros, which were
deprecated in November 2024 with PR #112849 / #115507.
Where required, the wavefront size should instead be queried via means provided
by the HIP runtime: the (non-constexpr) `warpSize` variable in device code, or
`hipGetDeviceProperties` in host code.
This change passed AMD-internal testing.
Implements SWDEV-522062.
---
clang/docs/AMDGPUSupport.rst | 4 -
clang/docs/HIPSupport.rst | 3 +-
clang/lib/Basic/Targets/AMDGPU.cpp | 6 -
.../CodeGenHIP/maybe_undef-attr-verify.hip | 2 +-
.../CodeGenOpenCL/builtins-amdgcn-wave32.cl | 6 +-
.../CodeGenOpenCL/builtins-amdgcn-wave64.cl | 4 -
clang/test/Driver/amdgpu-macros.cl | 16 ---
clang/test/Driver/hip-macros.hip | 23 ----
...wavefront-size-deprecation-diagnostics.hip | 115 ------------------
.../Preprocessor/predefined-arch-macros.c | 2 -
10 files changed, 3 insertions(+), 178 deletions(-)
delete mode 100644
clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip
diff --git a/clang/docs/AMDGPUSupport.rst b/clang/docs/AMDGPUSupport.rst
index 3eada5f900613..18e3de8abe92a 100644
--- a/clang/docs/AMDGPUSupport.rst
+++ b/clang/docs/AMDGPUSupport.rst
@@ -49,10 +49,6 @@ Predefined Macros
- Defined as 1 if the CU mode is enabled and 0 if the WGP mode is enabled.
* - ``__AMDGCN_UNSAFE_FP_ATOMICS__``
- Defined if unsafe floating-point atomics are allowed.
- * - ``__AMDGCN_WAVEFRONT_SIZE__``
- - Defines the wavefront size. Allowed values are 32 and 64 (deprecated).
- * - ``__AMDGCN_WAVEFRONT_SIZE``
- - Alias to ``__AMDGCN_WAVEFRONT_SIZE__`` (deprecated).
* - ``__HAS_FMAF__``
- Defined if FMAF instruction is available (deprecated).
* - ``__HAS_LDEXPF__``
diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst
index b4a671e3cfa3c..0d04b842af025 100644
--- a/clang/docs/HIPSupport.rst
+++ b/clang/docs/HIPSupport.rst
@@ -178,8 +178,7 @@ Predefined Macros
- Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
Note that some architecture specific AMDGPU macros will have default values
when
-used from the HIP host compilation. Other :doc:`AMDGPU macros <AMDGPUSupport>`
-like ``__AMDGCN_WAVEFRONT_SIZE__`` (deprecated) will default to 64 for example.
+used from the HIP host compilation.
Compilation Modes
=================
diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp
b/clang/lib/Basic/Targets/AMDGPU.cpp
index 87de9e6865e71..443dfbc93a182 100644
--- a/clang/lib/Basic/Targets/AMDGPU.cpp
+++ b/clang/lib/Basic/Targets/AMDGPU.cpp
@@ -356,12 +356,6 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions
&Opts,
if (hasFastFMA())
Builder.defineMacro("FP_FAST_FMA");
- 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");
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 571fba148f5cc..6dc57c4fcc5fc 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 = WARP_SIZE;
__device__ static inline unsigned int __lane_id() {
return __builtin_amdgcn_mbcnt_hi(
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
index d390418523694..31fd0e7bceaf5 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
@@ -1,5 +1,5 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown
-D__AMDGCN_WAVEFRONT_SIZE=32 -target-feature +wavefrontsize32 -emit-llvm -o -
%s | FileCheck -enable-var-scope %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown
-target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck
-enable-var-scope %s
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu
gfx1010 -emit-llvm -o - %s | FileCheck -enable-var-scope %s
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu
gfx1010 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck
-enable-var-scope %s
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu
gfx1100 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck
-enable-var-scope %s
@@ -48,7 +48,3 @@ void test_read_exec_lo(global uint* out) {
void test_read_exec_hi(global uint* out) {
*out = __builtin_amdgcn_read_exec_hi();
}
-
-#if __AMDGCN_WAVEFRONT_SIZE != 32
-#error Wrong wavesize detected
-#endif
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
index d851ec7e6734f..758b5aa532d73 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
@@ -50,7 +50,3 @@ void test_read_exec_lo(global ulong* out) {
void test_read_exec_hi(global ulong* out) {
*out = __builtin_amdgcn_read_exec_hi();
}
-
-#if defined(__AMDGCN_WAVEFRONT_SIZE__) && __AMDGCN_WAVEFRONT_SIZE__ != 64
-#error Wrong wavesize detected
-#endif
diff --git a/clang/test/Driver/amdgpu-macros.cl
b/clang/test/Driver/amdgpu-macros.cl
index a60593f2ab9ed..dd6fcc773a32b 100644
--- a/clang/test/Driver/amdgpu-macros.cl
+++ b/clang/test/Driver/amdgpu-macros.cl
@@ -153,26 +153,10 @@
// ARCH-GCN-DAG: #define __[[CPU]]__ 1
// ARCH-GCN-DAG: #define __[[FAMILY]]__ 1
// ARCH-GCN-DAG: #define __amdgcn_processor__ "[[CPU]]"
-// ARCH-GCN-DAG: #define __AMDGCN_WAVEFRONT_SIZE [[WAVEFRONT_SIZE]]
// ARCH-GCN-DAG: #define __GCC_DESTRUCTIVE_SIZE 128
// ARCH-GCN-DAG: #define __GCC_CONSTRUCTIVE_SIZE 128
// UNSAFEFPATOMIC-DAG: #define __AMDGCN_UNSAFE_FP_ATOMICS__ 1
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
-// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
-// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
-// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
-// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE32 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mno-wavefrontsize64 \
-// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mno-wavefrontsize64 \
-// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
-// WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE 64
-// WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE 32
-
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 \
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mcumode \
diff --git a/clang/test/Driver/hip-macros.hip b/clang/test/Driver/hip-macros.hip
index 516e01a6c4743..4c460d50bf39a 100644
--- a/clang/test/Driver/hip-macros.hip
+++ b/clang/test/Driver/hip-macros.hip
@@ -1,27 +1,4 @@
// REQUIRES: amdgpu-registered-target
-// RUN: %clang -E -dM --offload-arch=gfx906 -mwavefrontsize64 \
-// RUN: --cuda-device-only -nogpuinc -nogpulib \
-// RUN: %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
-// RUN: %clang -E -dM --offload-arch=gfx1010 -mwavefrontsize64 \
-// RUN: --cuda-device-only -nogpuinc -nogpulib \
-// RUN: %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
-// RUN: %clang -E -dM --offload-arch=gfx906 -mwavefrontsize64 \
-// RUN: --cuda-device-only -nogpuinc -nogpulib \
-// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
-// RUN: %clang -E -dM --offload-arch=gfx1010 -mwavefrontsize64 \
-// 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 \
diff --git a/clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip
b/clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip
deleted file mode 100644
index 8a60f5a150048..0000000000000
--- a/clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip
+++ /dev/null
@@ -1,115 +0,0 @@
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang -xhip --offload-arch=gfx1030 --offload-host-only -pedantic
-nogpuinc -nogpulib -nobuiltininc -fsyntax-only -Xclang -verify %s
-// RUN: %clang -xhip --offload-arch=gfx1030 --offload-device-only -pedantic
-nogpuinc -nogpulib -nobuiltininc -fsyntax-only -Xclang -verify %s
-
-// Test that deprecation warnings for the wavefront size macro are emitted
properly.
-
-#define WRAPPED __AMDGCN_WAVEFRONT_SIZE__
-
-#define DOUBLE_WRAPPED (WRAPPED)
-
-template <bool C, class T = void> struct my_enable_if {};
-
-template <class T> struct my_enable_if<true, T> {
- typedef T type;
-};
-
-__attribute__((host, device)) void use(int, const char*);
-
-template<int N> __attribute__((host, device)) int templatify(int x) {
- return x + N;
-}
-
-__attribute__((device)) const int GlobalConst = __AMDGCN_WAVEFRONT_SIZE__; //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
-constexpr int GlobalConstExpr = __AMDGCN_WAVEFRONT_SIZE__; // expected-warning
{{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
-
-#if defined(__HIP_DEVICE_COMPILE__) && (__AMDGCN_WAVEFRONT_SIZE__ == 64) //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
-int foo(void);
-#endif
-
-__attribute__((device)) int device_var = __AMDGCN_WAVEFRONT_SIZE__; //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
-
-__attribute__((device))
-void device_fun() {
- use(__AMDGCN_WAVEFRONT_SIZE, "device function"); // expected-warning
{{macro '__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
- use(__AMDGCN_WAVEFRONT_SIZE__, "device function"); // expected-warning
{{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(WRAPPED, "device function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(DOUBLE_WRAPPED, "device function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "device function"); //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
- use(GlobalConst, "device function");
- use(GlobalConstExpr, "device function");
-}
-
-__attribute__((global))
-void global_fun() {
- // no warnings expected
- use(__AMDGCN_WAVEFRONT_SIZE, "global function"); // expected-warning
{{macro '__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
- use(__AMDGCN_WAVEFRONT_SIZE__, "global function"); // expected-warning
{{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(WRAPPED, "global function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(DOUBLE_WRAPPED, "global function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "global function"); //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
-}
-
-int host_var = __AMDGCN_WAVEFRONT_SIZE__; // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
-int host_var_alt = __AMDGCN_WAVEFRONT_SIZE; // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
-int host_var_wrapped = WRAPPED; // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
-int host_var_double_wrapped = DOUBLE_WRAPPED; // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
-
-__attribute__((host))
-void host_fun() {
- use(__AMDGCN_WAVEFRONT_SIZE, "host function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
- use(__AMDGCN_WAVEFRONT_SIZE__, "host function"); // expected-warning
{{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(WRAPPED, "host function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(DOUBLE_WRAPPED, "host function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "host function"); //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
- use(GlobalConst, "host function");
- use(GlobalConstExpr, "host function");
-}
-
-__attribute((host, device))
-void host_device_fun() {
- use(__AMDGCN_WAVEFRONT_SIZE__, "host device function"); //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
- use(WRAPPED, "host device function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(DOUBLE_WRAPPED, "host device function"); // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "host device function"); //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
-}
-
-template <unsigned int OuterWarpSize = __AMDGCN_WAVEFRONT_SIZE__> //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
-class FunSelector {
-public:
- template<unsigned int FunWarpSize = OuterWarpSize>
- __attribute__((device))
- auto fun(void)
- -> typename my_enable_if<(FunWarpSize <= __AMDGCN_WAVEFRONT_SIZE__),
void>::type // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been
marked as deprecated}}
- {
- use(1, "yay!");
- }
-
- template<unsigned int FunWarpSize = OuterWarpSize>
- __attribute__((device))
- auto fun(void)
- -> typename my_enable_if<(FunWarpSize > __AMDGCN_WAVEFRONT_SIZE__),
void>::type // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been
marked as deprecated}}
- {
- use(0, "nay!");
- }
-};
-
-__attribute__((device))
-void device_fun_selector_user() {
- FunSelector<> f;
- f.fun<>();
- f.fun<1>();
- f.fun<1000>();
-
- my_enable_if<(1 <= __AMDGCN_WAVEFRONT_SIZE__), int>::type x = 42; //
expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as
deprecated}}
-}
-
-__attribute__((device)) my_enable_if<(1 <= __AMDGCN_WAVEFRONT_SIZE__),
int>::type DeviceFunTemplateRet(void) { // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- return 42;
-}
-
-__attribute__((device)) int DeviceFunTemplateArg(my_enable_if<(1 <=
__AMDGCN_WAVEFRONT_SIZE__), int>::type x) { // expected-warning {{macro
'__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
- return x;
-}
-
-// expected-note@* 0+ {{macro marked 'deprecated' here}}
diff --git a/clang/test/Preprocessor/predefined-arch-macros.c
b/clang/test/Preprocessor/predefined-arch-macros.c
index ecddf130a5c51..ebdfc8b79e063 100644
--- a/clang/test/Preprocessor/predefined-arch-macros.c
+++ b/clang/test/Preprocessor/predefined-arch-macros.c
@@ -4410,7 +4410,6 @@
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FMAF__
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FP64__
// CHECK_AMDGCN_NONE-NOT: #define __HAS_LDEXPF__
-// CHECK_AMDGCN_NONE-NOT: #define __AMDGCN_WAVEFRONT_SIZE__
// Begin r600 tests ----------------
@@ -4431,7 +4430,6 @@
// RUN: %clang -x hip -E -dM %s -o - 2>&1 --offload-host-only -nogpulib \
// RUN: -nogpuinc --offload-arch=gfx803 -target x86_64-unknown-linux \
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_HIP_HOST
-// CHECK_HIP_HOST: #define __AMDGCN_WAVEFRONT_SIZE__ 64
// CHECK_HIP_HOST: #define __AMDGPU__ 1
// CHECK_HIP_HOST: #define __AMD__ 1
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
