[CUDA][HIP] Change default lang std to c++14

Currently clang and nvcc use c++14 as default std for C++.
gcc 11 even uses c++17 as default std for C++. However,
clang uses c++98 as default std for CUDA/HIP.

As c++14 has been well adopted and became default for
clang, it seems reasonable to use c++14 as default std
for CUDA/HIP.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D103221

clang/include/clang/Basic/LangStandards.def

@@ -194,11 +194,11 @@ LANGSTANDARD_ALIAS_DEPR(openclcpp, "CLC++")
 
 // CUDA
 LANGSTANDARD(cuda, "cuda", CUDA, "NVIDIA CUDA(tm)",
-             LineComment | CPlusPlus | Digraphs)
+             LineComment | CPlusPlus | CPlusPlus11 | CPlusPlus14 | Digraphs)
 
 // HIP
 LANGSTANDARD(hip, "hip", HIP, "HIP",
-             LineComment | CPlusPlus | Digraphs)
+             LineComment | CPlusPlus | CPlusPlus11 | CPlusPlus14 | Digraphs)
 
 #undef LANGSTANDARD
 #undef LANGSTANDARD_ALIAS

clang/test/Parser/cuda-kernel-call.cu

@@ -12,8 +12,8 @@ void foo(void) {
 
   // The following two are parse errors because -std=c++11 is not enabled.
 
-  S<S<S<int>>> s; // expected-error 2{{use '> >'}}
-  S<S<S<>>> s1; // expected-error 2{{use '> >'}}
-  (void)(&f<S<S<int>>>==0); // expected-error 2{{use '> >'}}
-  (void)(&f<S<S<>>>==0); // expected-error 2{{use '> >'}}
+  S<S<S<int>>> s;
+  S<S<S<>>> s1;
+  (void)(&f<S<S<int>>>==0);
+  (void)(&f<S<S<>>>==0);
 }

clang/test/Preprocessor/lang-std.cu

@@ -0,0 +1,7 @@
+// RUN: %clang_cc1 -dM -E -x hip %s | FileCheck -check-prefix=CXX14 %s
+// RUN: %clang_cc1 -dM -E %s | FileCheck -check-prefix=CXX14 %s
+// RUN: %clang_cc1 -dM -E -std=c++98 -x hip %s | FileCheck -check-prefix=CXX98 %s
+// RUN: %clang_cc1 -dM -E -std=c++98 %s | FileCheck -check-prefix=CXX98 %s
+
+// CXX98: #define __cplusplus 199711L
+// CXX14: #define __cplusplus 201402L

clang/test/SemaCUDA/asm_delayed_diags.cu

@@ -28,7 +28,7 @@ static __device__ __host__ unsigned t2(signed char input) {
 }
 
 static __device__ __host__ double t3(double x) {
-  register long double result;
+  long double result;
   __asm __volatile("frndint"
                    : "=t"(result)
                    : "0"(x));

clang/test/SemaCUDA/cuda-builtin-vars.cu

@@ -1,4 +1,5 @@
-// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda" -fcuda-is-device -fsyntax-only -verify \
+// RUN:   -verify-ignore-unexpected=note %s
 
 #include "__clang_cuda_builtin_vars.h"
 __attribute__((global))
@@ -37,17 +38,13 @@ void kernel(int *out) {
   // 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@__clang_cuda_builtin_vars.h:* {{declared private here}}
+  __cuda_builtin_threadIdx_t x; // expected-error {{call to deleted constructor of '__cuda_builtin_threadIdx_t'}}
 
-  __cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{calling a private constructor of class '__cuda_builtin_threadIdx_t'}}
-  // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}
+  __cuda_builtin_threadIdx_t y = threadIdx; // expected-error {{call to deleted constructor of '__cuda_builtin_threadIdx_t'}}
 
-  threadIdx = threadIdx; // expected-error {{'operator=' is a private member of '__cuda_builtin_threadIdx_t'}}
-  // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}
+  threadIdx = threadIdx; // expected-error {{overload resolution selected deleted operator '='}}
 
-  void *ptr = &threadIdx; // expected-error {{'operator&' is a private member of '__cuda_builtin_threadIdx_t'}}
-  // expected-note@__clang_cuda_builtin_vars.h:* {{declared private here}}
+  void *ptr = &threadIdx; // expected-error {{overload resolution selected deleted operator '&'}}
 
   // 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

clang/test/SemaCUDA/function-target.cu

@@ -9,6 +9,7 @@ __host__ __device__ void h1hd(void);
 __global__ void h1g(void);
 
 struct h1ds { // expected-note {{requires 1 argument}}
+	      // expected-note@-1 {{candidate constructor (the implicit move constructor) not viable}}
   __device__ h1ds(); // expected-note {{candidate constructor not viable: call to __device__ function from __host__ function}}
 };
 

clang/test/SemaCUDA/implicit-member-target-collision.cu

@@ -19,6 +19,7 @@ struct C1_with_collision : A1_with_host_ctor, B1_with_device_ctor {
 // expected-note@-3 {{candidate constructor (the implicit default constructor}} not viable
 // expected-note@-4 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
 // expected-note@-5 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-6 {{candidate constructor (the implicit move constructor) not viable}}
 
 void hostfoo1() {
   C1_with_collision c; // expected-error {{no matching constructor}}
@@ -35,6 +36,7 @@ struct C2_with_collision {
 // expected-note@-5 {{candidate constructor (the implicit default constructor}} not viable
 // expected-note@-6 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
 // expected-note@-7 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-8 {{candidate constructor (the implicit move constructor) not viable}}
 
 void hostfoo2() {
   C2_with_collision c; // expected-error {{no matching constructor}}
@@ -51,6 +53,7 @@ struct C3_with_collision : A1_with_host_ctor {
 // expected-note@-4 {{candidate constructor (the implicit default constructor}} not viable
 // expected-note@-5 {{implicit default constructor inferred target collision: call to both __host__ and __device__ members}}
 // expected-note@-6 {{candidate constructor (the implicit copy constructor}} not viable
+// expected-note@-7 {{candidate constructor (the implicit move constructor) not viable}}
 
 void hostfoo3() {
   C3_with_collision c; // expected-error {{no matching constructor}}