Skip to content

Commit 276271f

Browse files
ritter-x2aritter-x2a
committed
Reapply "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros" (llvm#164217)
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)
1 parent 4affeb6 commit 276271f

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
@@ -354,12 +354,6 @@ void AMDGPUTargetInfo::getTargetDefines(const LangOptions &Opts,
354354
if (hasFastFMA())
355355
Builder.defineMacro("FP_FAST_FMA");
356356

357-
Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE__", Twine(WavefrontSize),
358-
"compile-time-constant access to the wavefront size will "
359-
"be removed in a future release");
360-
Builder.defineMacro("__AMDGCN_WAVEFRONT_SIZE", Twine(WavefrontSize),
361-
"compile-time-constant access to the wavefront size will "
362-
"be removed in a future release");
363357
Builder.defineMacro("__AMDGCN_CUMODE__", Twine(CUMode));
364358
}
365359

‎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
@@ -41,7 +41,3 @@ void test_read_exec_lo(global uint* out) {
4141
void test_read_exec_hi(global uint* out) {
4242
*out = __builtin_amdgcn_read_exec_hi();
4343
}
44-
45-
#if __AMDGCN_WAVEFRONT_SIZE != 32
46-
#error Wrong wavesize detected
47-
#endif

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

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,3 @@ void test_read_exec_lo(global ulong* out) {
4343
void test_read_exec_hi(global ulong* out) {
4444
*out = __builtin_amdgcn_read_exec_hi();
4545
}
46-
47-
#if defined(__AMDGCN_WAVEFRONT_SIZE__) && __AMDGCN_WAVEFRONT_SIZE__ != 64
48-
#error Wrong wavesize detected
49-
#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
@@ -4460,7 +4460,6 @@
44604460
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FMAF__
44614461
// CHECK_AMDGCN_NONE-NOT: #define __HAS_FP64__
44624462
// CHECK_AMDGCN_NONE-NOT: #define __HAS_LDEXPF__
4463-
// CHECK_AMDGCN_NONE-NOT: #define __AMDGCN_WAVEFRONT_SIZE__
44644463

44654464
// Begin r600 tests ----------------
44664465

@@ -4481,7 +4480,6 @@
44814480
// RUN: %clang -x hip -E -dM %s -o - 2>&1 --offload-host-only -nogpulib \
44824481
// RUN: -nogpuinc --offload-arch=gfx803 -target x86_64-unknown-linux \
44834482
// RUN: | FileCheck -match-full-lines %s -check-prefixes=CHECK_HIP_HOST
4484-
// CHECK_HIP_HOST: #define __AMDGCN_WAVEFRONT_SIZE__ 64
44854483
// CHECK_HIP_HOST: #define __AMDGPU__ 1
44864484
// CHECK_HIP_HOST: #define __AMD__ 1
44874485

0 commit comments

Comments
 (0)