From f0a955d3fa54d2b8600e646bbfb6bf3217a18d0c Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 17 Jul 2023 11:57:38 -0400 Subject: [PATCH] [HIP] Rename predefined macros Rename HIP_API_PER_THREAD_DEFAULT_STREAM and __HIP_NO_IMAGE_SUPPORT so that they follow the convention with prefix and postfix __. Reviewed by: Artem Belevich Differential Revision: https://reviews.llvm.org/D155480 --- clang/lib/Frontend/InitPreprocessor.cpp | 10 ++++++++-- clang/test/CodeGenCUDA/Inputs/cuda.h | 4 ++-- clang/test/Driver/hip-macros.hip | 17 ++++++++++++++++- 3 files changed, 26 insertions(+), 5 deletions(-) diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 16dd0c0..2ab9e65 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -588,12 +588,18 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI, Builder.defineMacro("__HIP_MEMORY_SCOPE_SYSTEM", "5"); if (LangOpts.CUDAIsDevice) { Builder.defineMacro("__HIP_DEVICE_COMPILE__"); - if (!TI.hasHIPImageSupport()) + if (!TI.hasHIPImageSupport()) { + Builder.defineMacro("__HIP_NO_IMAGE_SUPPORT__", "1"); + // Deprecated. Builder.defineMacro("__HIP_NO_IMAGE_SUPPORT", "1"); + } } if (LangOpts.GPUDefaultStream == - LangOptions::GPUDefaultStreamKind::PerThread) + LangOptions::GPUDefaultStreamKind::PerThread) { + Builder.defineMacro("__HIP_API_PER_THREAD_DEFAULT_STREAM__"); + // Deprecated. Builder.defineMacro("HIP_API_PER_THREAD_DEFAULT_STREAM"); + } } } diff --git a/clang/test/CodeGenCUDA/Inputs/cuda.h b/clang/test/CodeGenCUDA/Inputs/cuda.h index 0639965..8df425b 100644 --- a/clang/test/CodeGenCUDA/Inputs/cuda.h +++ b/clang/test/CodeGenCUDA/Inputs/cuda.h @@ -35,7 +35,7 @@ int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0, extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0, hipStream_t stream = 0); -#ifndef HIP_API_PER_THREAD_DEFAULT_STREAM +#ifndef __HIP_API_PER_THREAD_DEFAULT_STREAM__ extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, @@ -45,7 +45,7 @@ extern "C" hipError_t hipLaunchKernel_spt(const void *func, dim3 gridDim, dim3 blockDim, void **args, size_t sharedMem, hipStream_t stream); -#endif //HIP_API_PER_THREAD_DEFAULT_STREAM +#endif // __HIP_API_PER_THREAD_DEFAULT_STREAM__ #else typedef struct cudaStream *cudaStream_t; typedef enum cudaError {} cudaError_t; diff --git a/clang/test/Driver/hip-macros.hip b/clang/test/Driver/hip-macros.hip index 6289ef9..9954ddf 100644 --- a/clang/test/Driver/hip-macros.hip +++ b/clang/test/Driver/hip-macros.hip @@ -61,5 +61,20 @@ // RUN: -Xclang -target-feature -Xclang "+image-insts" %s 2>&1 | FileCheck --check-prefixes=NOIMAGE,WARN %s // NOWARN-NOT: warning // WARN: warning: feature flag '{{[+|-]}}image-insts' is ignored since the feature is read only [-Winvalid-command-line-argument] +// IMAGE-NOT: #define __HIP_NO_IMAGE_SUPPORT__ // IMAGE-NOT: #define __HIP_NO_IMAGE_SUPPORT -// NOIMAGE: #define __HIP_NO_IMAGE_SUPPORT 1 +// NOIMAGE-DAG: #define __HIP_NO_IMAGE_SUPPORT__ 1 +// NOIMAGE-DAG: #define __HIP_NO_IMAGE_SUPPORT 1 + +// RUN: %clang -E -dM --offload-arch=gfx1100 -nogpuinc -nogpulib \ +// RUN: -fgpu-default-stream=per-thread %s 2>&1 | FileCheck --check-prefixes=PTS %s +// RUN: %clang -E -dM --offload-arch=gfx940 --cuda-device-only -nogpuinc -nogpulib \ +// RUN: -fgpu-default-stream=legacy %s 2>&1 | FileCheck --check-prefixes=NOPTS %s +// RUN: %clang -E -dM --offload-arch=gfx940 --cuda-device-only -nogpuinc -nogpulib \ +// RUN: %s 2>&1 | FileCheck --check-prefixes=NOPTS %s +// PTS-DAG: #define __HIP_API_PER_THREAD_DEFAULT_STREAM__ 1 +// PTS-DAG: #define __HIP_API_PER_THREAD_DEFAULT_STREAM__ 1 +// PTS-DAG: #define HIP_API_PER_THREAD_DEFAULT_STREAM 1 +// PTS-DAG: #define HIP_API_PER_THREAD_DEFAULT_STREAM 1 +// NOPTS-NOT: #define __HIP_API_PER_THREAD_DEFAULT_STREAM__ +// NOPTS-NOT: #define HIP_API_PER_THREAD_DEFAULT_STREAM -- 2.7.4