From 204eb555222d86837d865007c1526018e391bf67 Mon Sep 17 00:00:00 2001
From: Justin Lebar <jlebar@google.com>
Date: Sat, 8 Oct 2016 22:16:08 +0000
Subject: [PATCH] [CUDA] Rename cuda_builtin_vars.h to
__clang_cuda_builtin_vars.h.
Summary: This matches the idiom we use for our other CUDA wrapper headers.
Reviewers: tra
Subscribers: beanz, mgorny, cfe-commits
Differential Revision: https://reviews.llvm.org/D24978
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@283679 91177308-0d34-0410-b5e6-96231b3b80d8
---
lib/Frontend/CompilerInvocation.cpp | 7 ++++---
lib/Headers/CMakeLists.txt | 2 +-
...uiltin_vars.h => __clang_cuda_builtin_vars.h} | 0
lib/Headers/__clang_cuda_runtime_wrapper.h | 16 ++++++++--------
test/CodeGenCUDA/cuda-builtin-vars.cu | 2 +-
test/SemaCUDA/cuda-builtin-vars.cu | 12 ++++++------
6 files changed, 20 insertions(+), 19 deletions(-)
rename lib/Headers/{cuda_builtin_vars.h => __clang_cuda_builtin_vars.h} (100%)
diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp
index 0dd05b5d9a9..cf39f357012 100644
--- a/lib/Frontend/CompilerInvocation.cpp
+++ b/lib/Frontend/CompilerInvocation.cpp
@@ -2012,9 +2012,10 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
// enabled for Microsoft Extensions or Borland Extensions, here.
//
// FIXME: __declspec is also currently enabled for CUDA, but isn't really a
- // CUDA extension, however it is required for supporting cuda_builtin_vars.h,
- // which uses __declspec(property). Once that has been rewritten in terms of
- // something more generic, remove the Opts.CUDA term here.
+ // CUDA extension. However, it is required for supporting
+ // __clang_cuda_builtin_vars.h, which uses __declspec(property). Once that has
+ // been rewritten in terms of something more generic, remove the Opts.CUDA
+ // term here.
Opts.DeclSpecKeyword =
Args.hasFlag(OPT_fdeclspec, OPT_fno_declspec,
(Opts.MicrosoftExt || Opts.Borland || Opts.CUDA));
diff --git a/lib/Headers/CMakeLists.txt b/lib/Headers/CMakeLists.txt
index 958038e0cda..be18ea8dbf5 100644
--- a/lib/Headers/CMakeLists.txt
+++ b/lib/Headers/CMakeLists.txt
@@ -22,12 +22,12 @@ set(files
avxintrin.h
bmi2intrin.h
bmiintrin.h
+ __clang_cuda_builtin_vars.h
__clang_cuda_cmath.h
__clang_cuda_intrinsics.h
__clang_cuda_math_forward_declares.h
__clang_cuda_runtime_wrapper.h
cpuid.h
- cuda_builtin_vars.h
clflushoptintrin.h
emmintrin.h
f16cintrin.h
diff --git a/lib/Headers/cuda_builtin_vars.h b/lib/Headers/__clang_cuda_builtin_vars.h
similarity index 100%
rename from lib/Headers/cuda_builtin_vars.h
rename to lib/Headers/__clang_cuda_builtin_vars.h
diff --git a/lib/Headers/__clang_cuda_runtime_wrapper.h b/lib/Headers/__clang_cuda_runtime_wrapper.h
index 0cf8b17def5..6c6dff86adc 100644
--- a/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ b/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -72,9 +72,9 @@
#define __CUDA_ARCH__ 350
#endif
-#include "cuda_builtin_vars.h"
+#include "__clang_cuda_builtin_vars.h"
-// No need for device_launch_parameters.h as cuda_builtin_vars.h above
+// No need for device_launch_parameters.h as __clang_cuda_builtin_vars.h above
// has taken care of builtin variables declared in the file.
#define __DEVICE_LAUNCH_PARAMETERS_H__
@@ -283,8 +283,8 @@ __device__ static inline void *malloc(size_t __size) {
}
} // namespace std
-// Out-of-line implementations from cuda_builtin_vars.h. These need to come
-// after we've pulled in the definition of uint3 and dim3.
+// Out-of-line implementations from __clang_cuda_builtin_vars.h. These need to
+// come after we've pulled in the definition of uint3 and dim3.
__device__ inline __cuda_builtin_threadIdx_t::operator uint3() const {
uint3 ret;
@@ -315,10 +315,10 @@ __device__ inline __cuda_builtin_gridDim_t::operator dim3() const {
// curand_mtgp32_kernel helpfully redeclares blockDim and threadIdx in host
// mode, giving them their "proper" types of dim3 and uint3. This is
-// incompatible with the types we give in cuda_builtin_vars.h. As as hack,
-// force-include the header (nvcc doesn't include it by default) but redefine
-// dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are only
-// used here for the redeclarations of blockDim and threadIdx.)
+// incompatible with the types we give in __clang_cuda_builtin_vars.h. As as
+// hack, force-include the header (nvcc doesn't include it by default) but
+// redefine dim3 and uint3 to our builtin types. (Thankfully dim3 and uint3 are
+// only used here for the redeclarations of blockDim and threadIdx.)
#pragma push_macro("dim3")
#pragma push_macro("uint3")
#define dim3 __cuda_builtin_blockDim_t
diff --git a/test/CodeGenCUDA/cuda-builtin-vars.cu b/test/CodeGenCUDA/cuda-builtin-vars.cu
index c2159f5af14..c1edff936a0 100644
--- a/test/CodeGenCUDA/cuda-builtin-vars.cu
+++ b/test/CodeGenCUDA/cuda-builtin-vars.cu
@@ -1,6 +1,6 @@
// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
-#include "cuda_builtin_vars.h"
+#include "__clang_cuda_builtin_vars.h"
// CHECK: define void @_Z6kernelPi(i32* %out)
__attribute__((global))
diff --git a/test/SemaCUDA/cuda-builtin-vars.cu b/test/SemaCUDA/cuda-builtin-vars.cu
index 108e75cae76..27a9c5abd70 100644
--- a/test/SemaCUDA/cuda-builtin-vars.cu
+++ b/test/SemaCUDA/cuda-builtin-vars.cu
@@ -1,6 +1,6 @@
// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify %s
-#include "cuda_builtin_vars.h"
+#include "__clang_cuda_builtin_vars.h"
__attribute__((global))
void kernel(int *out) {
int i = 0;
@@ -34,20 +34,20 @@ void kernel(int *out) {
out[i++] = warpSize;
warpSize = 0; // expected-error {{cannot assign to variable 'warpSize' with const-qualified type 'const int'}}
- // expected-note@cuda_builtin_vars.h:* {{variable 'warpSize' declared const here}}
+ // expected-note@__clang_cuda_builtin_vars.h:* {{variable 'warpSize' declared const here}}
// Make sure we can't construct or assign to the special variables.
__cuda_builtin_threadIdx_t x; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
- // expected-note@cuda_builtin_vars.h:* {{declared private here}}
+ // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}
__cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
- // expected-note@cuda_builtin_vars.h:* {{declared private here}}
+ // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}
threadIdx = threadIdx; // expected-error {{'operator=' is a private member of '__cuda_builtin_threadIdx_t'}}
- // expected-note@cuda_builtin_vars.h:* {{declared private here}}
+ // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}
void *ptr = &threadIdx; // expected-error {{'operator&' is a private member of '__cuda_builtin_threadIdx_t'}}
- // expected-note@cuda_builtin_vars.h:* {{declared private here}}
+ // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}
// Following line should've caused an error as one is not allowed to
// take address of a built-in variable in CUDA. Alas there's no way
--
GitLab