Skip to content

Commit ea03447

Browse files
ritter-x2aritter-x2a
authored
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (#164217)
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.
1 parent a55a720 commit ea03447

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
@@ -180,8 +180,7 @@ Predefined Macros
180180
- Alias to ``__HIP_API_PER_THREAD_DEFAULT_STREAM__``. Deprecated.
181181

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

186185
Compilation Modes
187186
=================

‎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
@@ -154,26 +154,10 @@
154154
// ARCH-GCN-DAG: #define __[[CPU]]__ 1
155155
// ARCH-GCN-DAG: #define __[[FAMILY]]__ 1
156156
// ARCH-GCN-DAG: #define __amdgcn_processor__ "[[CPU]]"
157-
// ARCH-GCN-DAG: #define __AMDGCN_WAVEFRONT_SIZE [[WAVEFRONT_SIZE]]
158157
// ARCH-GCN-DAG: #define __GCC_DESTRUCTIVE_SIZE 128
159158
// ARCH-GCN-DAG: #define __GCC_CONSTRUCTIVE_SIZE 128
160159
// UNSAFEFPATOMIC-DAG: #define __AMDGCN_UNSAFE_FP_ATOMICS__ 1
161160

162-
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
163-
// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
164-
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
165-
// RUN: %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
166-
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mwavefrontsize64 \
167-
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
168-
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mwavefrontsize64 \
169-
// RUN: -mno-wavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE32 %s
170-
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 -mno-wavefrontsize64 \
171-
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
172-
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx1010 -mno-wavefrontsize64 \
173-
// RUN: -mwavefrontsize64 %s 2>&1 | FileCheck --check-prefix=WAVE64 %s
174-
// WAVE64-DAG: #define __AMDGCN_WAVEFRONT_SIZE 64
175-
// WAVE32-DAG: #define __AMDGCN_WAVEFRONT_SIZE 32
176-
177161
// RUN: %clang -E -dM -target amdgcn -mcpu=gfx906 \
178162
// RUN: %s 2>&1 | FileCheck --check-prefix=CUMODE-ON %s
179163
// 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
@@ -4418,7 +4418,6 @@
44184418
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FMAF__
44194419
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FP64__
44204420
// CHECK_AMDGCN_NONE-NOT: #define __HAS_LDEXPF__
4421-
// CHECK_AMDGCN_NONE-NOT: #define __AMDGCN_WAVEFRONT_SIZE__
44224421

44234422
// Begin r600 tests ----------------
44244423

@@ -4439,7 +4438,6 @@
44394438
// RUN: %clang -x hip -E -dM %s -o - 2>&1 --offload-host-only -nogpulib \
44404439
// RUN: -nogpuinc --offload-arch=gfx803 -target x86_64-unknown-linux \
44414440
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_HIP_HOST
4442-
// CHECK_HIP_HOST: #define __AMDGCN_WAVEFRONT_SIZE__ 64
44434441
// CHECK_HIP_HOST: #define __AMDGPU__ 1
44444442
// CHECK_HIP_HOST: #define __AMD__ 1
44454443

0 commit comments

Comments
 (0)