// 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));
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
#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__
}
} // 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;
// 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
// 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))
// 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;
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