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

[yaxunl]

Summary

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.

[tra]

clang/include/clang/Basic/LangStandards.def

196–197

It would make sense to bump C++ version for CUDA as well.

[yaxunl]

revised by Artem's comments, and fix tests

[yaxunl]

clang/include/clang/Basic/LangStandards.def

196–197

done

[tra]

clang/test/SemaCUDA/asm_delayed_diags.cu

31

Perhaps we should just remove the register.
TBH, I'm not quite sure what this test is supposed to do for NVPTX -- t is not a register class in NVPTX nor does it seem to be a generic constraint.

I guess the idea was that we should ignore asm errors in the host code, while we're compiling for device. In that case, we certainly do not need register

[yaxunl]

fix test

[yaxunl]

clang/test/SemaCUDA/asm_delayed_diags.cu

31

removed register.

I think this test is for deferred diags.

clang/include/clang/Basic/LangStandards.def

 LANGSTANDARD_ALIAS_DEPR(opencl10, "CL")
 LANGSTANDARD_ALIAS_DEPR(opencl11, "CL1.1")
 LANGSTANDARD_ALIAS_DEPR(opencl12, "CL1.2")
 LANGSTANDARD_ALIAS_DEPR(opencl20, "CL2.0")
 LANGSTANDARD_ALIAS_DEPR(opencl30, "CL3.0")
 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
 #undef LANGSTANDARD_ALIAS_DEPR

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

 // RUN: %clang_cc1 -fsyntax-only -verify %s
 
 template<typename T=int> struct S {};
 template<typename> void f();
 
 void foo(void) {
   foo<<<1;      // expected-error {{expected '>>>'}} expected-note {{to match this '<<<'}}
 
   foo<<<1,1>>>; // expected-error {{expected '('}}
 
   foo<<<>>>();  // expected-error {{expected expression}}
 
   // 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

+// 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

static __device__ __host__ unsigned t2(signed char input) {
   unsigned output;
   __asm__("xyz"
           : "=a"(output)
           : "0"(input));
   return output;
 }
 
 static __device__ __host__ double t3(double x) {
-  register long double result;
+  long double result;
   __asm __volatile("frndint"
                    : "=t"(result)
                    : "0"(x));
   return result;
 }
 
 static __device__ __host__ unsigned char t4(unsigned char a, unsigned char b) {
   unsigned int la = a;

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

-// 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))
 void kernel(int *out) {
   int i = 0;
   out[i++] = threadIdx.x;
   threadIdx.x = 0; // expected-error {{no setter defined for property 'x'}}
   out[i++] = threadIdx.y;
void kernel(int *out) {
   out[i++] = gridDim.z;
   gridDim.z = 0; // expected-error {{no setter defined for property 'z'}}
 
   out[i++] = warpSize;
   warpSize = 0; // expected-error {{cannot assign to variable 'warpSize' with const-qualified type 'const int'}}
   // 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
   // to prevent getting address of a 'const int', so the line
   // currently compiles without errors or warnings.
   const void *wsptr = &warpSize;
 }

clang/test/SemaCUDA/function-target.cu

 // RUN: %clang_cc1 -fsyntax-only -verify %s
 // RUN: %clang_cc1 -fsyntax-only -fcuda-is-device -verify=dev,expected %s
 
 #include "Inputs/cuda.h"
 
 __host__ void h1h(void);
 __device__ void h1d(void); // expected-note {{candidate function not viable: call to __device__ function from __host__ function}}
 __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}}
 };
 
 __host__ void h1(void) {
   h1h();
   h1d(); // expected-error {{no matching function}}
   h1hd();
   h1g<<<1, 1>>>();

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

 };
 
 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}}
 }
 
 //------------------------------------------------------------------------------
 // Test 2: collision between two fields
 
 struct C2_with_collision {
   A1_with_host_ctor aa;
   B1_with_device_ctor bb;
 };
 
 // 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}}
 
 }
 
 //------------------------------------------------------------------------------
 // Test 3: collision between a field and a base
 
 struct C3_with_collision : A1_with_host_ctor {
   B1_with_device_ctor bb;
 };
 
 // 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}}
 }