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");
+ }
}
}
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");
+ }
}
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
