Skip to content

[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros #157463

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
merged 1 commit into from
Sep 15, 2025
Merged

[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros #157463

merged 1 commit into from
Sep 15, 2025

Conversation

@ritter-x2a
Copy link
Member

@ritter-x2a ritter-x2a commented Sep 8, 2025

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.

@ritter-x2a
[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros
2b0505e 
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.
@ritter-x2a Graphite App
Copy link
Member Author

ritter-x2a commented Sep 8, 2025

This stack of pull requests is managed by Graphite. Learn more about stacking.

@ritter-x2a ritter-x2a added clang Clang issues not falling into any other category backend:AMDGPU labels Sep 8, 2025 — with Graphite App
@llvmbot
Copy link
Member

llvmbot commented Sep 8, 2025 •
edited
Loading

@llvm/pr-subscribers-clang-driver
@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-clang

Author: Fabian Ritter (ritter-x2a)

Changes

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.

Full diff: https://github.com/llvm/llvm-project/pull/157463.diff

10 Files Affected:

  • (modified) clang/docs/AMDGPUSupport.rst (-4)
  • (modified) clang/docs/HIPSupport.rst (+1-2)
  • (modified) clang/lib/Basic/Targets/AMDGPU.cpp (-6)
  • (modified) clang/test/CodeGenHIP/maybe_undef-attr-verify.hip (+1-1)
  • (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl (+1-5)
  • (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl (-4)
  • (modified) clang/test/Driver/amdgpu-macros.cl (-16)
  • (modified) clang/test/Driver/hip-macros.hip (-23)
  • (removed) clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip (-115)
  • (modified) clang/test/Preprocessor/predefined-arch-macros.c (-2)
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

@ritter-x2a ritter-x2a marked this pull request as ready for review September 8, 2025 13:40
@llvmbot llvmbot added clang:driver 'clang' and 'clang++' user-facing binaries. Not 'clang-cl' clang:frontend Language frontend issues, e.g. anything involving "Sema" labels Sep 8, 2025
jhuber6
jhuber6 approved these changes Sep 8, 2025
View reviewed changes
Copy link
Contributor

@jhuber6 jhuber6 left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LG on my end, the OpenMP and libc runtimes are already using the builtin for things that require the wavefront size.

@ritter-x2a Graphite App
Copy link
Member Author

ritter-x2a commented Sep 15, 2025 •
edited
Loading

Merge activity

  • Sep 15, 7:31 AM UTC: A user started a stack merge that includes this pull request via Graphite.
  • Sep 15, 7:33 AM UTC: @ritter-x2a merged this pull request with Graphite.

@ritter-x2a ritter-x2a merged commit 02d3e6a into main Sep 15, 2025
17 checks passed
@ritter-x2a ritter-x2a deleted the users/ritter-x2a/09-08-_hip_clang_remove___amdgcn_wavefront_size_macros branch September 15, 2025 07:33
ritter-x2a added a commit that referenced this pull request Sep 15, 2025
@ritter-x2a
Revert "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros (#157463)"
31a757f 
This reverts commit 02d3e6a.
@ritter-x2a ritter-x2a mentioned this pull request Sep 15, 2025
Merged
ritter-x2a added a commit that referenced this pull request Sep 15, 2025
@ritter-x2a
Revert "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (#158566)
78bf682 
Reverts #157463
The PR breaks buildbots with old ROCm versions, so revert it and reapply
when buildbots are updated.
llvm-sync bot pushed a commit to arm/arm-toolchain that referenced this pull request Sep 16, 2025
@ritter-x2a @github-actions
Automerge: Revert "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros…
bbb3a83 
…" (#158566)

Reverts llvm/llvm-project#157463
The PR breaks buildbots with old ROCm versions, so revert it and reapply
when buildbots are updated.
ritter-x2a added a commit to ritter-x2a/llvm-project that referenced this pull request Oct 20, 2025
@ritter-x2a
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros"
2910885 
This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.
@ritter-x2a ritter-x2a mentioned this pull request Oct 20, 2025
Merged
ritter-x2a added a commit to ritter-x2a/llvm-project that referenced this pull request Oct 30, 2025
@ritter-x2a
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros"
6f9819d 
This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.
ritter-x2a added a commit that referenced this pull request Oct 30, 2025
@ritter-x2a
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (#164217)
ea03447 
This reverts commit 78bf682.

Original PR: #157463
Revert PR: #158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.
ritter-x2a added a commit to ROCm/llvm-project that referenced this pull request Oct 30, 2025
@ritter-x2a
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (llvm#16…
276271f 
…4217)

This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.

(cherry picked from commit ea03447)
@ritter-x2a ritter-x2a mentioned this pull request Oct 30, 2025
Merged
aokblast pushed a commit to aokblast/llvm-project that referenced this pull request Oct 30, 2025
@ritter-x2a @aokblast
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (llvm#16…
2fc2173 
…4217)

This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.
kzhuravl pushed a commit to ROCm/llvm-project that referenced this pull request Oct 31, 2025
@ritter-x2a
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (llvm#16…
94f12a2 
…4217) (#434)

This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.

(cherry picked from commit ea03447)
luciechoi pushed a commit to luciechoi/llvm-project that referenced this pull request Nov 1, 2025
@ritter-x2a @luciechoi
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (llvm#16…
f8d913a 
…4217)

This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.
DEBADRIBASAK pushed a commit to DEBADRIBASAK/llvm-project that referenced this pull request Nov 3, 2025
@ritter-x2a
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (llvm#16…
570e242 
…4217)

This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.
david-salinas pushed a commit to ROCm/llvm-project that referenced this pull request Nov 4, 2025
@ritter-x2a
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (llvm#16…
47f994c 
…4217)

This reverts commit 78bf682.

Original PR: llvm#157463
Revert PR: llvm#158566

The relevant buildbots have been updated to a ROCm version that does not
use the macros anymore to avoid the failures.

Implements SWDEV-522062.
@david-salinas david-salinas mentioned this pull request Nov 4, 2025
Merged
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@shiltian shiltian shiltian approved these changes

@jhuber6 jhuber6 jhuber6 approved these changes

@AlexVlx AlexVlx Awaiting requested review from AlexVlx

@arsenm arsenm Awaiting requested review from arsenm

@yxsamliu yxsamliu Awaiting requested review from yxsamliu

Assignees

No one assigned

Labels

backend:AMDGPU clang:driver 'clang' and 'clang++' user-facing binaries. Not 'clang-cl' clang:frontend Language frontend issues, e.g. anything involving "Sema" clang Clang issues not falling into any other category

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

5 participants

@ritter-x2a @llvmbot @shiltian @jhuber6