Author: Fabian Ritter Date: 2025-09-15T10:04:06+02:00 New Revision: 31a757f2a38c493da868bf97f557c2e30bf24cca
URL: https://github.com/llvm/llvm-project/commit/31a757f2a38c493da868bf97f557c2e30bf24cca DIFF: https://github.com/llvm/llvm-project/commit/31a757f2a38c493da868bf97f557c2e30bf24cca.diff LOG: Revert "[HIP][Clang] Remove __AMDGCN_WAVEFRONT_SIZE macros (#157463)" This reverts commit 02d3e6ac75e776041fb1782efc4dfccfe6b46218. Added: clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip Modified: clang/docs/AMDGPUSupport.rst clang/docs/HIPSupport.rst clang/lib/Basic/Targets/AMDGPU.cpp clang/test/CodeGenHIP/maybe_undef-attr-verify.hip clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl clang/test/Driver/amdgpu-macros.cl clang/test/Driver/hip-macros.hip clang/test/Preprocessor/predefined-arch-macros.c Removed: ################################################################################ diff --git a/clang/docs/AMDGPUSupport.rst b/clang/docs/AMDGPUSupport.rst index 18e3de8abe92a..3eada5f900613 100644 --- a/clang/docs/AMDGPUSupport.rst +++ b/clang/docs/AMDGPUSupport.rst @@ -49,6 +49,10 @@ 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__`` diff --git a/clang/docs/HIPSupport.rst b/clang/docs/HIPSupport.rst index 0d04b842af025..b4a671e3cfa3c 100644 --- a/clang/docs/HIPSupport.rst +++ b/clang/docs/HIPSupport.rst @@ -178,7 +178,8 @@ 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. +used from the HIP host compilation. Other :doc:`AMDGPU macros <AMDGPUSupport>` +like ``__AMDGCN_WAVEFRONT_SIZE__`` (deprecated) will default to 64 for example. Compilation Modes ================= diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 443dfbc93a182..87de9e6865e71 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -356,6 +356,12 @@ 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)); } diff --git a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip index 6dc57c4fcc5fc..571fba148f5cc 100644 --- a/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip +++ b/clang/test/CodeGenHIP/maybe_undef-attr-verify.hip @@ -20,7 +20,7 @@ #define __maybe_undef __attribute__((maybe_undef)) #define WARP_SIZE 64 -static constexpr int warpSize = WARP_SIZE; +static constexpr int warpSize = __AMDGCN_WAVEFRONT_SIZE__; __device__ static inline unsigned int __lane_id() { return __builtin_amdgcn_mbcnt_hi( diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl index 31fd0e7bceaf5..d390418523694 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl @@ -1,5 +1,5 @@ // REQUIRES: amdgpu-registered-target -// 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 -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-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,3 +48,7 @@ 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 diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl index 758b5aa532d73..d851ec7e6734f 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-wave64.cl @@ -50,3 +50,7 @@ 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 diff --git a/clang/test/Driver/amdgpu-macros.cl b/clang/test/Driver/amdgpu-macros.cl index dd6fcc773a32b..a60593f2ab9ed 100644 --- a/clang/test/Driver/amdgpu-macros.cl +++ b/clang/test/Driver/amdgpu-macros.cl @@ -153,10 +153,26 @@ // 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 \ diff --git a/clang/test/Driver/hip-macros.hip b/clang/test/Driver/hip-macros.hip index 4c460d50bf39a..516e01a6c4743 100644 --- a/clang/test/Driver/hip-macros.hip +++ b/clang/test/Driver/hip-macros.hip @@ -1,4 +1,27 @@ // 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 \ diff --git a/clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip b/clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip new file mode 100644 index 0000000000000..8a60f5a150048 --- /dev/null +++ b/clang/test/Driver/hip-wavefront-size-deprecation-diagnostics.hip @@ -0,0 +1,115 @@ +// 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}} diff --git a/clang/test/Preprocessor/predefined-arch-macros.c b/clang/test/Preprocessor/predefined-arch-macros.c index ebdfc8b79e063..ecddf130a5c51 100644 --- a/clang/test/Preprocessor/predefined-arch-macros.c +++ b/clang/test/Preprocessor/predefined-arch-macros.c @@ -4410,6 +4410,7 @@ // 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 ---------------- @@ -4430,6 +4431,7 @@ // 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 _______________________________________________ llvm-branch-commits mailing list llvm-branch-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
