[clang] f0a955d - [HIP] Rename predefined macros

[Yaxun Liu via cfe-commits]

Author: Yaxun (Sam) Liu
Date: 2023-07-17T15:45:49-04:00
New Revision: f0a955d3fa54d2b8600e646bbfb6bf3217a18d0c

URL: https://github.com/llvm/llvm-project/commit/f0a955d3fa54d2b8600e646bbfb6bf3217a18d0c
DIFF: https://github.com/llvm/llvm-project/commit/f0a955d3fa54d2b8600e646bbfb6bf3217a18d0c.diff

LOG: [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

Added: 
    

Modified: 
    clang/lib/Frontend/InitPreprocessor.cpp
    clang/test/CodeGenCUDA/Inputs/cuda.h
    clang/test/Driver/hip-macros.hip

Removed: 
    


################################################################################
diff  --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp
index 16dd0c01bcb443..2ab9e652b10c3c 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 06399659c0b53e..8df425b77ac823 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 6289ef91dd04fa..9954ddf075e66f 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