[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

Author: yxsamliu

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");
+    }
   }
 }
 

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;

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