shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity

Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (#164217)

Browse Source 
This reverts commit 78bf682cb9033cf6a5bbc733e062c7b7d825fdaf.

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.
This commit is contained in:
Fabian Ritter 2025-10-30 13:42:32 +01:00 committed by GitHub
parent a55a7207c7
commit ea034477fd
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
10 changed files with 3 additions and 178 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
clang/docs/AMDGPUSupport.rst
View File

@ -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__``

3
clang/docs/HIPSupport.rst
View File

@ -180,8 +180,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
=================

6
clang/lib/Basic/Targets/AMDGPU.cpp
View File

@ -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));
}

2
clang/test/CodeGenHIP/maybe_undef-attr-verify.hip
View File

@ -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(

6
clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl
View File

@ -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

4
clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl
View File

@ -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

16
clang/test/Driver/amdgpu-macros.cl
View File

@ -154,26 +154,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 \

23
clang/test/Driver/hip-macros.hip
View File

@ -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 \

115
clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip
View File

@ -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}}

2
clang/test/Preprocessor/predefined-arch-macros.c
View File

@ -4418,7 +4418,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 ----------------
@ -4439,7 +4438,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