Skip to content

Commit 31a757f

Browse files
ritter-x2aritter-x2a
authored
Revert "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros (#157463)"
This reverts commit 02d3e6a.
1 parent e4124c0 commit 31a757f

File tree

10 files changed

+178
-3
lines changed

10 files changed

+178
-3
lines changed

‎clang/docs/AMDGPUSupport.rst‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -49,6 +49,10 @@ Predefined Macros
4949
- Defined as 1 if the CU mode is enabled and 0 if the WGP mode is enabled.
5050
* - ``__AMDGCN_UNSAFE_FP_ATOMICS__``
5151
- Defined if unsafe floating-point atomics are allowed.
52+
* - ``__AMDGCN_WAVEFRONT_SIZE__``
53+
- Defines the wavefront size. Allowed values are 32 and 64 (deprecated).
54+
* - ``__AMDGCN_WAVEFRONT_SIZE``
55+
- Alias to ``__AMDGCN_WAVEFRONT_SIZE__`` (deprecated).
5256
* - ``__HAS_FMAF__``
5357
- Defined if FMAF instruction is available (deprecated).
5458
* - ``__HAS_LDEXPF__``

‎clang/docs/HIPSupport.rst‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -178,7 +178,8 @@ Predefined Macros
178178
- Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
179179

180180
Note that some architecture specific AMDGPU macros will have default values when
181-
used from the HIP host compilation.
181+
used from the HIP host compilation. Other :doc:`AMDGPU macros <AMDGPUSupport>`
182+
like ``__AMDGCN_WAVEFRONT_SIZE__`` (deprecated) will default to 64 for example.
182183

183184
Compilation Modes
184185
=================

‎clang/lib/Basic/Targets/AMDGPU.cpp‎

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -356,6 +356,12 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
356356
if (hasFastFMA())
357357
Builder.defineMacro("FP_FAST_FMA");
358358

359+
Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE__", Twine(WavefrontSize),
360+
"compile-time-constant access to the wavefront size will "
361+
"be removed in a future release");
362+
Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE", Twine(WavefrontSize),
363+
"compile-time-constant access to the wavefront size will "
364+
"be removed in a future release");
359365
Builder.defineMacro("__AMDGCN_CUMODE__", Twine(CUMode));
360366
}
361367

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -20,7 +20,7 @@
2020
#define __maybe_undef __attribute__((maybe_undef))
2121
#define WARP_SIZE 64
2222

23-
static constexpr int warpSize = WARP_SIZE;
23+
static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE__;
2424

2525
__device__ static inline unsigned int __lane_id() {
2626
return __builtin_amdgcn_mbcnt_hi(

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// REQUIRES: amdgpu-registered-target
2-
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck -enable-var-scope %s
2+
// 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
33
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck -enable-var-scope %s
44
// 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
55
// 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,3 +48,7 @@ void test_read_exec_lo(global uint* out) {
4848
void test_read_exec_hi(global uint* out) {
4949
*out = __builtin_amdgcn_read_exec_hi();
5050
}
51+
52+
#if __AMDGCN_WAVEFRONT_SIZE != 32
53+
#error Wrong wavesize detected
54+
#endif

‎clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -50,3 +50,7 @@ void test_read_exec_lo(global ulong* out) {
5050
void test_read_exec_hi(global ulong* out) {
5151
*out = __builtin_amdgcn_read_exec_hi();
5252
}
53+
54+
#if defined(__AMDGCN_WAVEFRONT_SIZE__) && __AMDGCN_WAVEFRONT_SIZE__ != 64
55+
#error Wrong wavesize detected
56+
#endif

‎clang/test/Driver/amdgpu-macros.cl‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -153,10 +153,26 @@
153153
// ARCH-GCN-DAG: #define __[[CPU]]__ 1
154154
// ARCH-GCN-DAG: #define __[[FAMILY]]__ 1
155155
// ARCH-GCN-DAG: #define __amdgcn_processor__ "[[CPU]]"
156+
// ARCH-GCN-DAG: #define __AMDGCN_WAVEFRONT_SIZE [[WAVEFRONT_SIZE]]
156157
// ARCH-GCN-DAG: #define __GCC_DESTRUCTIVE_SIZE 128
157158
// ARCH-GCN-DAG: #define __GCC_CONSTRUCTIVE_SIZE 128
158159
// UNSAFEFPATOMIC-DAG: #define __AMDGCN_UNSAFE_FP_ATOMICS__ 1
159160

161+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
162+
// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
163+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
164+
// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
165+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
166+
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
167+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
168+
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE32 %s
169+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mno-wavefrontsize64 \
170+
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
171+
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mno-wavefrontsize64 \
172+
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
173+
// WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE 64
174+
// WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE 32
175+
160176
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 \
161177
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
162178
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mcumode \

‎clang/test/Driver/hip-macros.hip‎

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,27 @@
11
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang -E -dM --offload-arch=gfx906 -mwavefrontsize64 \
3+
// RUN: --cuda-device-only -nogpuinc -nogpulib \
4+
// RUN: %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
5+
// RUN: %clang -E -dM --offload-arch=gfx1010 -mwavefrontsize64 \
6+
// RUN: --cuda-device-only -nogpuinc -nogpulib \
7+
// RUN: %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
8+
// RUN: %clang -E -dM --offload-arch=gfx906 -mwavefrontsize64 \
9+
// RUN: --cuda-device-only -nogpuinc -nogpulib \
10+
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
11+
// RUN: %clang -E -dM --offload-arch=gfx1010 -mwavefrontsize64 \
12+
// RUN: --cuda-device-only -nogpuinc -nogpulib \
13+
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE32 %s
14+
// RUN: %clang -E -dM --offload-arch=gfx906 -mno-wavefrontsize64 \
15+
// RUN: --cuda-device-only -nogpuinc -nogpulib \
16+
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
17+
// RUN: %clang -E -dM --offload-arch=gfx1010 -mno-wavefrontsize64 \
18+
// RUN: --cuda-device-only -nogpuinc -nogpulib \
19+
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefixes=WAVE64 %s
20+
// WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE__ 64
21+
// WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE__ 32
22+
// WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE 64
23+
// WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE 32
24+
225
// RUN: %clang -E -dM --offload-arch=gfx906 --cuda-device-only -nogpuinc -nogpulib \
326
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
427
// RUN: %clang -E -dM --offload-arch=gfx906 --cuda-device-only -nogpuinc -nogpulib -mcumode \

‎clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip‎

Lines changed: 115 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,115 @@
1+
// REQUIRES: amdgpu-registered-target
2+
// RUN: %clang -xhip --offload-arch=gfx1030 --offload-host-only -pedantic -nogpuinc -nogpulib -nobuiltininc -fsyntax-only -Xclang -verify %s
3+
// RUN: %clang -xhip --offload-arch=gfx1030 --offload-device-only -pedantic -nogpuinc -nogpulib -nobuiltininc -fsyntax-only -Xclang -verify %s
4+
5+
// Test that deprecation warnings for the wavefront size macro are emitted properly.
6+
7+
#define WRAPPED __AMDGCN_WAVEFRONT_SIZE__
8+
9+
#define DOUBLE_WRAPPED (WRAPPED)
10+
11+
template <bool C, class T = void> struct my_enable_if {};
12+
13+
template <class T> struct my_enable_if<true, T> {
14+
typedef T type;
15+
};
16+
17+
__attribute__((host, device)) void use(int, const char*);
18+
19+
template<int N> __attribute__((host, device)) int templatify(int x) {
20+
return x + N;
21+
}
22+
23+
__attribute__((device)) const int GlobalConst = __AMDGCN_WAVEFRONT_SIZE__; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
24+
constexpr int GlobalConstExpr = __AMDGCN_WAVEFRONT_SIZE__; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
25+
26+
#if defined(__HIP_DEVICE_COMPILE__) && (__AMDGCN_WAVEFRONT_SIZE__ == 64) // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
27+
int foo(void);
28+
#endif
29+
30+
__attribute__((device)) int device_var = __AMDGCN_WAVEFRONT_SIZE__; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
31+
32+
__attribute__((device))
33+
void device_fun() {
34+
use(__AMDGCN_WAVEFRONT_SIZE, "device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
35+
use(__AMDGCN_WAVEFRONT_SIZE__, "device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
36+
use(WRAPPED, "device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
37+
use(DOUBLE_WRAPPED, "device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
38+
use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
39+
use(GlobalConst, "device function");
40+
use(GlobalConstExpr, "device function");
41+
}
42+
43+
__attribute__((global))
44+
void global_fun() {
45+
// no warnings expected
46+
use(__AMDGCN_WAVEFRONT_SIZE, "global function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
47+
use(__AMDGCN_WAVEFRONT_SIZE__, "global function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
48+
use(WRAPPED, "global function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
49+
use(DOUBLE_WRAPPED, "global function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
50+
use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "global function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
51+
}
52+
53+
int host_var = __AMDGCN_WAVEFRONT_SIZE__; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
54+
int host_var_alt = __AMDGCN_WAVEFRONT_SIZE; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
55+
int host_var_wrapped = WRAPPED; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
56+
int host_var_double_wrapped = DOUBLE_WRAPPED; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
57+
58+
__attribute__((host))
59+
void host_fun() {
60+
use(__AMDGCN_WAVEFRONT_SIZE, "host function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE' has been marked as deprecated}}
61+
use(__AMDGCN_WAVEFRONT_SIZE__, "host function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
62+
use(WRAPPED, "host function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
63+
use(DOUBLE_WRAPPED, "host function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
64+
use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "host function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
65+
use(GlobalConst, "host function");
66+
use(GlobalConstExpr, "host function");
67+
}
68+
69+
__attribute((host, device))
70+
void host_device_fun() {
71+
use(__AMDGCN_WAVEFRONT_SIZE__, "host device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
72+
use(WRAPPED, "host device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
73+
use(DOUBLE_WRAPPED, "host device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
74+
use(templatify<__AMDGCN_WAVEFRONT_SIZE__>(42), "host device function"); // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
75+
}
76+
77+
template <unsigned int OuterWarpSize = __AMDGCN_WAVEFRONT_SIZE__> // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
78+
class FunSelector {
79+
public:
80+
template<unsigned int FunWarpSize = OuterWarpSize>
81+
__attribute__((device))
82+
auto fun(void)
83+
-> typename my_enable_if<(FunWarpSize <= __AMDGCN_WAVEFRONT_SIZE__), void>::type // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
84+
{
85+
use(1, "yay!");
86+
}
87+
88+
template<unsigned int FunWarpSize = OuterWarpSize>
89+
__attribute__((device))
90+
auto fun(void)
91+
-> typename my_enable_if<(FunWarpSize > __AMDGCN_WAVEFRONT_SIZE__), void>::type // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
92+
{
93+
use(0, "nay!");
94+
}
95+
};
96+
97+
__attribute__((device))
98+
void device_fun_selector_user() {
99+
FunSelector<> f;
100+
f.fun<>();
101+
f.fun<1>();
102+
f.fun<1000>();
103+
104+
my_enable_if<(1 <= __AMDGCN_WAVEFRONT_SIZE__), int>::type x = 42; // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
105+
}
106+
107+
__attribute__((device)) my_enable_if<(1 <= __AMDGCN_WAVEFRONT_SIZE__), int>::type DeviceFunTemplateRet(void) { // expected-warning {{macro '__AMDGCN_WAVEFRONT_SIZE__' has been marked as deprecated}}
108+
return 42;
109+
}
110+
111+
__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}}
112+
return x;
113+
}
114+
115+
// expected-note@* 0+ {{macro marked 'deprecated' here}}

‎clang/test/Preprocessor/predefined-arch-macros.c‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4410,6 +4410,7 @@
44104410
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FMAF__
44114411
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FP64__
44124412
// CHECK_AMDGCN_NONE-NOT: #define __HAS_LDEXPF__
4413+
// CHECK_AMDGCN_NONE-NOT: #define __AMDGCN_WAVEFRONT_SIZE__
44134414

44144415
// Begin r600 tests ----------------
44154416

@@ -4430,6 +4431,7 @@
44304431
// RUN: %clang -x hip -E -dM %s -o - 2>&1 --offload-host-only -nogpulib \
44314432
// RUN: -nogpuinc --offload-arch=gfx803 -target x86_64-unknown-linux \
44324433
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_HIP_HOST
4434+
// CHECK_HIP_HOST: #define __AMDGCN_WAVEFRONT_SIZE__ 64
44334435
// CHECK_HIP_HOST: #define __AMDGPU__ 1
44344436
// CHECK_HIP_HOST: #define __AMD__ 1
44354437

0 commit comments

Comments
 (0)