[OpenCL] Reuse C++ for OpenCL 1.0 address space tests for version 2021

[Topotuna]

Summary

Adds RUN lines in C++ for OpenCL address space tests to cover
version 2021. With all optional core features enabled by default,
C++ for OpenCL 2021 is expected to behave exactly the same way as
C++ for OpenCL 1.0. Therefore, no adjustments were needed in test
code itself.

[Topotuna]

Test coverage expanded

clang/test/CodeGenOpenCLCXX/address-space-castoperators.cpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 void test_reinterpret_cast(){
 __private float x;
 __private float& y = x; 
 // We don't need bitcast to cast pointer type and
 // address space at the same time.
 //CHECK: addrspacecast float* %x to i32 addrspace(4)*
 //CHECK: [[REG:%[0-9]+]] = load float*, float** %y
 //CHECK: addrspacecast float* [[REG]] to i32 addrspace(4)*
 //CHECK-NOT: bitcast 
 __generic int& rc1 = reinterpret_cast<__generic int&>(x);
 __generic int& rc2 = reinterpret_cast<__generic int&>(y);
 }

clang/test/CodeGenOpenCLCXX/address-space-deduction.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s -check-prefixes=COMMON,PTR
 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s -check-prefixes=COMMON,PTR
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF
 
 #ifdef REF
 #define PTR &
 #define ADR(x) x
 #else
 #define PTR *
 #define ADR(x) &x
 #endif

clang/test/CodeGenOpenCLCXX/address-space-deduction2.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 class P {
 public:
   P(const P &Rhs) = default;
 
   long A;
   long B;
 };

clang/test/CodeGenOpenCLCXX/addrspace-conversion.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 void bar(__generic volatile unsigned int* ptr)
 {
   //CHECK: addrspacecast i32 addrspace(4)* %{{.+}} to i32 addrspace(1)*
   auto gptr = (__global volatile unsigned int*)ptr;
 }

clang/test/CodeGenOpenCLCXX/addrspace-derived-base.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 struct B {
   int mb;
 };
 
 class D : public B {
 public:
   int getmb() { return mb; }

clang/test/CodeGenOpenCLCXX/addrspace-new-delete.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 typedef __SIZE_TYPE__ size_t;
 
 class A {
 public:
   void* operator new(size_t);
   void operator delete(void *ptr);
 };
 
 void test_new_delete(A **a) {
 // CHECK: %{{.*}} = call spir_func i8 addrspace(4)* @_ZNU3AS41AnwEj(i32 {{.*}})
     *a = new A;
 // CHECK: call spir_func void @_ZNU3AS41AdlEPU3AS4v(i8 addrspace(4)* {{.*}})
     delete *a;
 }

clang/test/CodeGenOpenCLCXX/addrspace-of-this.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL"
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL"
 // expected-no-diagnostics
 
 // Test that the 'this' pointer is in the __generic address space.
 
 #ifdef USE_DEFLT
 #define DEFAULT =default
 #else
 #define DEFAULT

clang/test/CodeGenOpenCLCXX/addrspace-operators.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 enum E {
   a,
   b,
 };
 
 class C {
 public:

clang/test/CodeGenOpenCLCXX/addrspace-references.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 typedef short short2 __attribute__((ext_vector_type(2)));
 
 int bar(const unsigned int &i);
 
 class C {
 public:
   void bar(const short2 &);

clang/test/CodeGenOpenCLCXX/addrspace-with-class.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s --check-prefix=CHECK-DEFINITIONS
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s --check-prefix=CHECK-DEFINITIONS
 
 // This test ensures the proper address spaces and address space cast are used
 // for constructors, member functions and destructors.
 // See also atexit.cl and global_init.cl for other specific tests.
 
 // CHECK: %struct.MyType = type { i32 }
 struct MyType {
   MyType(int i) : i(i) {}

clang/test/CodeGenOpenCLCXX/addrspace_cast.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 //CHECK-LABEL: define{{.*}} spir_func void @_Z3barPU3AS1i
 void bar(global int *gl) {
   //CHECK: addrspacecast i32 addrspace(1)* %{{[0-9]+}} to i32 addrspace(4)*
   int *gen = addrspace_cast<int *>(gl);
 }

clang/test/SemaOpenCLCXX/address-space-castoperators.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
 
 void nester_ptr() {
   local int * * locgen;
   constant int * * congen;
   int * * gengen;
 
   gengen = const_cast<int**>(locgen); //expected-error{{const_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}}
   gengen = static_cast<int**>(locgen); //expected-error{{static_cast from '__local int *__generic *' to '__generic int *__generic *' is not allowed}}
 // CHECK-NOT: AddressSpaceConversion
   gengen = reinterpret_cast<int**>(locgen); //expected-warning{{reinterpret_cast from '__local int *__generic *' to '__generic int *__generic *' changes address space of nested pointers}}
 
   gengen = const_cast<int**>(congen); //expected-error{{const_cast from '__constant int *__generic *' to '__generic int *__generic *' is not allowed}}
   gengen = static_cast<int**>(congen); //expected-error{{static_cast from '__constant int *__generic *' to '__generic int *__generic *' is not allowed}}
 // CHECK-NOT: AddressSpaceConversion
   gengen = reinterpret_cast<int**>(congen); //expected-warning{{reinterpret_cast from '__constant int *__generic *' to '__generic int *__generic *' changes address space of nested pointers}}
 }

clang/test/SemaOpenCLCXX/address-space-cond.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -verify
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -verify
 
 namespace PointerRvalues {
 
 void f(__global int *__constant *a, const __global int *__constant *b) {
   using T = decltype(true ? +a : +b);
   using T = const __global int *const __constant *;
 }
 

clang/test/SemaOpenCLCXX/address-space-deduction.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
 
 //expected-no-diagnostics
 
 //CHECK: |-VarDecl {{.*}} foo 'const __global int'
 constexpr int foo = 0;
 
 //CHECK: |-VarDecl {{.*}} foo1 'T' cinit
 //CHECK: `-VarTemplateSpecializationDecl {{.*}} used foo1 '__global long':'__global long' cinit

clang/test/SemaOpenCLCXX/address-space-lambda.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple i386-windows -pedantic -ast-dump -verify | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
 
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (__private int){{.*}} const __generic'
 auto glambda = [](auto a) { return a; };
 
 __kernel void test() {
   int i;
 //CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __generic'
   auto  llambda = [&]() {i++;};

 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __private'
   auto priv1 = []() __private {};
   priv1();
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __generic'
   auto priv2 = []() __generic {};
   priv2();
   auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}}
 #if defined(_WIN32) && !defined(_WIN64)
-  //expected-note@35{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
+  //expected-note@-2{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
 #else
-  //expected-note@35{{conversion candidate of type 'void (*)()'}}
+  //expected-note@-4{{conversion candidate of type 'void (*)()'}}
 #endif
   priv3(); //expected-error{{no matching function for call to object of type}}
 
   __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}}
 #if defined(_WIN32) && !defined(_WIN64)
-  //expected-note@43{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
+  //expected-note@-2{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
 #else
-  //expected-note@43{{conversion candidate of type 'void (*)()'}}
+  //expected-note@-4{{conversion candidate of type 'void (*)()'}}
 #endif
   const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
   __constant auto const2 = []() __generic{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__generic'}}
 #if defined(_WIN32) && !defined(_WIN64)
-  //expected-note@50{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
+  //expected-note@-2{{conversion candidate of type 'void (*)() __attribute__((thiscall))'}}
 #else
-  //expected-note@50{{conversion candidate of type 'void (*)()'}}
+  //expected-note@-4{{conversion candidate of type 'void (*)()'}}
 #endif
   const2(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __constant'
   __constant auto const3 = []() __constant{};
   const3();
 
   [&] () __global {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__global'}}
   [&] () __private {} (); //expected-error{{no matching function for call to object of type '(lambda at}} expected-note{{candidate function not viable: 'this' object is in default address space, but method expects object in address space '__private'}}

clang/test/SemaOpenCLCXX/address-space-of-this-class-scope.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -verify
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -verify
 
 struct C {
   auto fGlob() __global -> decltype(this);
   auto fGen() -> decltype(this);
   auto fErr() __global __local -> decltype(this); //expected-error{{multiple address spaces specified for type}}
 };
 
 void bar(__local C*);

clang/test/SemaOpenCLCXX/address-space-of-this.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
 // expected-no-diagnostics
 
 // Extract from PR38614
 struct C {};
 
 C f1() {
  return C{};
 }
 
 C f2(){
     C c;
     return c;
 }

clang/test/SemaOpenCLCXX/address-space-references.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
 
 __global const int& f(__global float &ref) {
   return ref; // expected-error{{reference of type 'const __global int &' cannot bind to a temporary object because of address space mismatch}}
 }
 
 int bar(const __global unsigned int &i); // expected-note{{passing argument to parameter 'i' here}}
 //FIXME: With the overload below the call should be resolved
 // successfully. However, current overload resolution logic

clang/test/SemaOpenCLCXX/address-space-templates.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -verify -ast-dump | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -verify -ast-dump | FileCheck %s
 
 template <typename T>
 struct S {
   T a;        // expected-error{{field may not be qualified with an address space}}
   T f1();     // we ignore address space on a return types.
   void f2(T); // expected-error{{parameter may not be qualified with an address space}}
 };
 

clang/test/SemaOpenCLCXX/address_space_overloading.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -verify -pedantic -fsyntax-only
 // expected-no-diagnostics
 
 struct RetGlob {
   int dummy;
 };
 
 struct RetGen {
   char dummy;

clang/test/SemaOpenCLCXX/addrspace-auto.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
 
 __constant int i = 1;
 //CHECK: |-VarDecl {{.*}} ai '__global int':'__global int'
 auto ai = i;
 
 kernel void test() {
   int i;
   //CHECK: VarDecl {{.*}} ai '__private int':'__private int'

clang/test/SemaOpenCLCXX/addrspace-constructors.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -ast-dump -verify | FileCheck %s
 
 __constant int g1; // expected-error {{variable in constant address space must be initialized}}
 __constant int g2 = 0;
 
 struct X {
   int x;
 //CHECK:  CXXConstructorDecl
 //CHECK-NOT:  used

clang/test/SemaOpenCLCXX/addrspace_cast.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
 
 void foo(global int *gl, const global int *gl_const, global int &gl_ref) {
   //FIXME: Diagnostics can be improved to be more specific in some cases.
   float *gen_fl = addrspace_cast<float *>(gl); //expected-error{{addrspace_cast from '__global int *__private' to '__generic float *' is not allowed}}
 
   int i = addrspace_cast<int>(gl); //expected-error{{addrspace_cast from '__global int *__private' to 'int' is not allowed}}
 
   int *gen = addrspace_cast<int *>(*gl); //expected-error{{addrspace_cast from '__global int' to '__generic int *' is not allowed}}

clang/test/SemaOpenCLCXX/addrspace_cast_ast_dump.clcpp

 // RUN: %clang_cc1 %s -cl-std=clc++1.0 -triple spir-unknown-unknown -pedantic -verify -ast-dump | FileCheck %s
+// RUN: %clang_cc1 %s -cl-std=clc++2021 -triple spir-unknown-unknown -pedantic -verify -ast-dump | FileCheck %s
 
 // expected-no-diagnostics
 
 //CHECK:`-FunctionDecl {{.*}} bar 'void (__global int *__private)'
 //CHECK:  |-ParmVarDecl {{.*}} used gl '__global int *__private'
 //CHECK:      `-VarDecl {{.*}} gen '__generic int *__private' cinit
 //CHECK:        `-CXXAddrspaceCastExpr {{.*}} '__generic int *' addrspace_cast<__generic int *> <AddressSpaceConversion>
 //CHECK:          `-DeclRefExpr {{.*}} '__global int *__private' lvalue ParmVar {{.*}} 'gl' '__global int *__private'
 
 void bar(global int *gl) {
   int *gen = addrspace_cast<int *>(gl);
 }