Skip to content

Commit 02d3e6a

Browse files
ritter-x2aritter-x2a
authored
[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros (#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.
1 parent c60972a commit 02d3e6a

File tree

10 files changed

+3
-178
lines changed

10 files changed

+3
-178
lines changed

‎clang/docs/AMDGPUSupport.rst‎

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -49,10 +49,6 @@ 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).
5652
* - ``__HAS_FMAF__``
5753
- Defined if FMAF instruction is available (deprecated).
5854
* - ``__HAS_LDEXPF__``

‎clang/docs/HIPSupport.rst‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -178,8 +178,7 @@ 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. Other :doc:`AMDGPU macros <AMDGPUSupport>`
182-
like ``__AMDGCN_WAVEFRONT_SIZE__`` (deprecated) will default to 64 for example.
181+
used from the HIP host compilation.
183182

184183
Compilation Modes
185184
=================

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

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -356,12 +356,6 @@ 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");
365359
Builder.defineMacro("__AMDGCN_CUMODE__", Twine(CUMode));
366360
}
367361

‎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 = __AMDGCN_WAVEFRONT_SIZE__;
23+
static constexpr int warpSize = WARP_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: 1 addition & 5 deletions
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 -D__AMDGCN_WAVEFRONT_SIZE=32 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck -enable-var-scope %s
2+
// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -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,7 +48,3 @@ 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: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,3 @@ 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: 0 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -153,26 +153,10 @@
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]]
157156
// ARCH-GCN-DAG: #define __GCC_DESTRUCTIVE_SIZE 128
158157
// ARCH-GCN-DAG: #define __GCC_CONSTRUCTIVE_SIZE 128
159158
// UNSAFEFPATOMIC-DAG: #define __AMDGCN_UNSAFE_FP_ATOMICS__ 1
160159

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-
176160
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 \
177161
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
178162
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mcumode \

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

Lines changed: 0 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -1,27 +1,4 @@
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-
252
// RUN: %clang -E -dM --offload-arch=gfx906 --cuda-device-only -nogpuinc -nogpulib \
263
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
274
// RUN: %clang -E -dM --offload-arch=gfx906 --cuda-device-only -nogpuinc -nogpulib -mcumode \

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

Lines changed: 0 additions & 115 deletions
This file was deleted.

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

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4410,7 +4410,6 @@
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__
44144413

44154414
// Begin r600 tests ----------------
44164415

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

0 commit comments

Comments
 (0)