[PATCH] D155480: [HIP] Rename predefined macros

[Yaxun Liu via Phabricator via cfe-commits]

This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rGf0a955d3fa54: [HIP] Rename predefined macros (authored by yaxunl).
Herald added a project: clang.

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D155480/new/

https://reviews.llvm.org/D155480

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


Index: clang/test/Driver/hip-macros.hip
===================================================================
--- clang/test/Driver/hip-macros.hip
+++ 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
Index: clang/test/CodeGenCUDA/Inputs/cuda.h
===================================================================
--- clang/test/CodeGenCUDA/Inputs/cuda.h
+++ clang/test/CodeGenCUDA/Inputs/cuda.h
@@ -35,7 +35,7 @@
 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 @@
                                       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;
Index: clang/lib/Frontend/InitPreprocessor.cpp
===================================================================
--- clang/lib/Frontend/InitPreprocessor.cpp
+++ clang/lib/Frontend/InitPreprocessor.cpp
@@ -588,12 +588,18 @@
     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");
+    }
   }
 }
 


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D155480.541191.patch
Type: text/x-patch
Size: 3580 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230717/d7f7b1b9/attachment.bin>