[OpenCL] Add distinct file extension for C++ for OpenCL

[Anastasia]

Summary

Files compiled with C++ for OpenCL can have a distinct file extension - clcpp, then clang driver will pick the compilation mode automatically without the use of -cl-std=clc++

See also RFC: https://lists.llvm.org/pipermail/cfe-dev/2021-February/067703.html

[mantognini]

I feel this would be a valuable addition, indeed. Only a minor question from me.

clang/lib/Driver/Types.cpp

252

I'm not sure we want that -- I'm actually fine if we don't -- but I see below c++ and cxx are supported in addition to cpp. Should we therefore also have clc++ and clcxx as file valid extensions for consistency? I wonder what the general opinion is.

[svenvh]

clang/lib/Driver/Types.cpp

252

And there is also .cc above. Since you're asking for opinions, my personal opinion is that one extension should be enough, and providing a CL counterpart for all existing C++ file extensions does not bring more value.

[Anastasia]

clang/lib/Driver/Types.cpp

252

FYI, there are some more discussions on RFC thread https://lists.llvm.org/pipermail/cfe-dev/2021-February/067703.html

[awarzynski]

Hi @Anastasia , thank you for working on this!

IIUC, this patch introduces:

• new input types: TY_CLCXX and TY_CLCXXHeader
• new language: OpenCLCXX

Based on the attached test, this is only to remove the need for -cl-std=clc++. Do you expect any other benefits of introducing the above ^^^? And are you planning to to update the OpenCL C++ tests not to use -cl-std=clc++? Just trying to make sure I understand the reason behind this :)

clang/lib/Driver/Types.cpp

163

This is not used.

[Anastasia]

In D96771#2571582, @awarzynski wrote:

Hi @Anastasia , thank you for working on this!

IIUC, this patch introduces:

• new input types: TY_CLCXX and TY_CLCXXHeader
• new language: OpenCLCXX

Based on the attached test, this is only to remove the need for -cl-std=clc++. Do you expect any other benefits of introducing the above ^^^?

This is only the initial patch and for the moment the primary goal is to remove the need for the flag at least from the clang perspective.

And are you planning to to update the OpenCL C++ tests not to use -cl-std=clc++?

I was not sure about this. I just dislike the fact that moving files in the repo complicates the commit history lookup so I was not sure if the renaming was good or bad thing. Do you have any suggestions?

I was thinking to introduce as a guideline that the new tests should definitely use the new extension though.

I find the way driver is setup for the languages quite different so I was wondering if we have any guidelines?

clang/lib/Driver/Types.cpp

163

This is used in the patch that is planned to be committed hopefully soon:
https://reviews.llvm.org/D96515

[awarzynski]

In D96771#2571855, @Anastasia wrote:

This is only the initial patch and for the moment the primary goal is to remove the need for the flag at least from the clang perspective.

Shorter compiler invocation is always nice! Does OpenCL (will) require a flag to specify the C++ standard used in the kernel? Or is that going to be controlled with -cl-std={cl2.1|cl2.2|cl2.3? This would be very weird -cl-std=clc++ -cl-std=cl2.1, so I see why you would want to remove -cl-std=clc++.

I just dislike the fact that moving files in the repo complicates the commit history lookup so I was not sure if the renaming was good or bad thing. Do you have any suggestions?

Have you tried it? Git is very clever in this respect and will track the file changes very accurately anyway. IMHO this wouldn't be disruptive in terms of browsing the history at all. If you are making this change to remove the need for -cl-std=clc++ and to simplify things, you could as well organize tests per kernel language (i.e. OpenCL C vs OpenCL C++). Also, -cl-std=clc++ in tests will look very confusing (i.e., why keep OpenCL C++ tests in OpenCL C tests and force the standard with -cl-std=clc++)? Lastly, such change would be a very through verification of your changes :)

I was thinking to introduce as a guideline that the new tests should definitely use the new extension though.

+1

I find the way driver is setup for the languages quite different so I was wondering if we have any guidelines?

I'm not aware of any documentation for the driver, sorry.

By scanning your patch I get the impression that you don't need TY_CLCXX just yet. The language for the kernel is set in clang/lib/Frontend/CompilerInvocation.cpp:

.Case("clcpp", Language::OpenCLCXX)

In Types.cpp this should be sufficient for now:

.Case("clcpp", TY_CLCXX)

Basically, you could limit your changes to the frontend driver and the changes in the compiler driver keep to a tiny minimum. I think! :)

[Anastasia]

• Renamed extensions to cppcl as it was more popular in RFC
• Minimized number of changes in driver setup
• Changed file extension in all C++ for OpenCL tests

[Anastasia]

In D96771#2574638, @awarzynski wrote:

In D96771#2571855, @Anastasia wrote:

This is only the initial patch and for the moment the primary goal is to remove the need for the flag at least from the clang perspective.

Shorter compiler invocation is always nice! Does OpenCL (will) require a flag to specify the C++ standard used in the kernel? Or is that going to be controlled with -cl-std={cl2.1|cl2.2|cl2.3? This would be very weird -cl-std=clc++ -cl-std=cl2.1, so I see why you would want to remove -cl-std=clc++.

Thanks a lot for all the feedback! Just to clarify I am not removing the flag I am just removing the need for it to be passed. FYI I have modified the driver test slightly to indicate that the flag is still available. The reason for it is backward compatibility. OpenCL (not only C++ for OpenCL) has always used this flag to override the default settings per file extensions. Some IDEs and other toolchains started relying on this feature. For example, they would compile .c extension file with -cl-std=cl. Not sure why it was set up this way but it is not something that would be easy to change without introducing backward compatibility issue.

I just dislike the fact that moving files in the repo complicates the commit history lookup so I was not sure if the renaming was good or bad thing. Do you have any suggestions?

Have you tried it? Git is very clever in this respect and will track the file changes very accurately anyway. IMHO this wouldn't be disruptive in terms of browsing the history at all. If you are making this change to remove the need for -cl-std=clc++ and to simplify things, you could as well organize tests per kernel language (i.e. OpenCL C vs OpenCL C++). Also, -cl-std=clc++ in tests will look very confusing (i.e., why keep OpenCL C++ tests in OpenCL C tests and force the standard with -cl-std=clc++)? Lastly, such change would be a very through verification of your changes :)

Yes, I agree. I have renamed the tests. The only drawback is that git log would need an extra option to see the full history but it should be acceptable and I am up for consistency. Also if we are ever to rename those tests now is better than later.

By scanning your patch I get the impression that you don't need TY_CLCXX just yet. The language for the kernel is set in clang/lib/Frontend/CompilerInvocation.cpp:

.Case("clcpp", Language::OpenCLCXX)

In Types.cpp this should be sufficient for now:

.Case("clcpp", TY_CLCXX)

Basically, you could limit your changes to the frontend driver and the changes in the compiler driver keep to a tiny minimum. I think! :)

Yes, I have removed most of the redundant code. It was mainly related to the headers btw. It looks a lot cleaner now. :)

[svenvh]

LGTM, but maybe give this another 24h before landing in case there are any remaining concerns.

[awarzynski]

Thank you for addressing my comments and for working on this! I've left one small suggestion, but that's a [nit].

clang/test/Driver/cxx_for_opencl.cppcl

1 ↗

(On Diff #327915)

This is a very neat test. Would it make sense to add something more basic too? For example (in a separate file):

// RUN: %clang %s -fstynax-only -### | FileCheck %s

// CHECK: -x -cppcl

This way you would have a bit more explicit test to verifiy that the compiler driver picks the right language based on the extension.

[Anastasia]

Updated Driver test to check -x option.

[Anastasia]

I had a second thought about the extension name and I realized that the reason why I initially wanted to use clcpp is that it aligns better with clc++ which is used in -cl-std. Even though from the RFC the preference was towards cppcl it felt like there was no objection to clcpp either. So I just want to check one last time whether it would make sense to align with clc++ and use clcpp. Perhaps, it make clang interface a bit more inconsistent?

clang/test/Driver/cxx_for_opencl.cppcl

1 ↗

(On Diff #327915)

Good idea! Thanks!

[svenvh]

In D96771#2626507, @Anastasia wrote:

I had a second thought about the extension name and I realized that the reason why I initially wanted to use clcpp is that it aligns better with clc++ which is used in -cl-std. Even though from the RFC the preference was towards cppcl it felt like there was no objection to clcpp either. So I just want to check one last time whether it would make sense to align with clc++ and use clcpp. Perhaps, it make clang interface a bit more inconsistent?

That is a good point, in my opinion aligning with the -cl-std flag is a good reason to go for .clcpp. I would recommend mentioning this on the RfC thread too, to ensure that people who suggested a slight preference for the other extension are okay with .clcpp too.

[awarzynski]

In D96771#2628691, @svenvh wrote:

In D96771#2626507, @Anastasia wrote:

I had a second thought about the extension name and I realized that the reason why I initially wanted to use clcpp is that it aligns better with clc++ which is used in -cl-std. Even though from the RFC the preference was towards cppcl it felt like there was no objection to clcpp either. So I just want to check one last time whether it would make sense to align with clc++ and use clcpp. Perhaps, it make clang interface a bit more inconsistent?

That is a good point, in my opinion aligning with the -cl-std flag is a good reason to go for .clcpp. I would recommend mentioning this on the RfC thread too, to ensure that people who suggested a slight preference for the other extension are okay with .clcpp too.

+1 for consistency.

[Anastasia]

Back to .clcpp!

[Anastasia]

Thanks, I have changed it to '.clcpp' and also replied in RFC: https://lists.llvm.org/pipermail/cfe-dev/2021-March/067936.html

Let's wait a few days for some more feedback.

[svenvh]

LGTM; just one minor comment, which you can address as part of your commit.

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

1–2

Nitpick: this -cl-std can be removed too.

[gribozavr2]

Could you also update the switch statement in LLDB here: llvm-project/lldb/source/Plugins/TypeSystem/Clang/TypeSystemClang.cpp:472 ?

[gribozavr2]

Pushed the most minimal fix in https://github.com/llvm/llvm-project/commit/41454c30f6a38c3e107d857e63da0561610fd141, please double check.

[Anastasia]

In D96771#2647848, @gribozavr2 wrote:

Pushed the most minimal fix in https://github.com/llvm/llvm-project/commit/41454c30f6a38c3e107d857e63da0561610fd141, please double check.

This looks great! Thanks for fixing this!

clang/include/clang/Basic/LangStandard.h

enum class Language : uint8_t {
   LLVM_IR,
 
   ///@{ Languages that the frontend can parse and compile.
   C,
   CXX,
   ObjC,
   ObjCXX,
   OpenCL,
+  OpenCLCXX,
   CUDA,
   RenderScript,
   HIP,
   ///@}
 };
 
 enum LangFeatures {
   LineComment = (1 << 0),

clang/include/clang/Driver/Types.def

 // options flag string will be replaced with this variadic list.
 // Most of the options in Flags have been removed in favor of subsuming their
 // meaning from the phases list.
 
 // C family source language (with and without preprocessing).
 TYPE("cpp-output",               PP_C,         INVALID,         "i",      phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("c",                        C,            PP_C,            "c",      phases::Preprocess, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("cl",                       CL,           PP_C,            "cl",     phases::Preprocess, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
+TYPE("clcpp",                    CLCXX,        PP_CXX,          "clcpp",  phases::Preprocess, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("cuda-cpp-output",          PP_CUDA,      INVALID,         "cui",    phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("cuda",                     CUDA,         PP_CUDA,         "cu",     phases::Preprocess, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("cuda",                     CUDA_DEVICE,  PP_CUDA,         "cu",     phases::Preprocess, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("hip-cpp-output",           PP_HIP,       INVALID,         "cui",    phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("hip",                      HIP,          PP_HIP,          "cu",     phases::Preprocess, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("hip",                      HIP_DEVICE,   PP_HIP,          "cu",     phases::Preprocess, phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("objective-c-cpp-output",   PP_ObjC,      INVALID,         "mi",     phases::Compile, phases::Backend, phases::Assemble, phases::Link)
 TYPE("objc-cpp-output",          PP_ObjC_Alias, INVALID,        "mi",     phases::Compile, phases::Backend, phases::Assemble, phases::Link)

clang/lib/Driver/Types.cpp

 
 bool types::isAcceptedByClang(ID Id) {
   switch (Id) {
   default:
     return false;
 
   case TY_Asm:
   case TY_C: case TY_PP_C:
-  case TY_CL:
+  case TY_CL: case TY_CLCXX:
   case TY_CUDA: case TY_PP_CUDA:
   case TY_CUDA_DEVICE:
   case TY_HIP:
   case TY_PP_HIP:
   case TY_HIP_DEVICE:
   case TY_ObjC: case TY_PP_ObjC: case TY_PP_ObjC_Alias:
   case TY_CXX: case TY_PP_CXX:
   case TY_ObjCXX: case TY_PP_ObjCXX: case TY_PP_ObjCXX_Alias:
bool types::isObjC(ID Id) {
   case TY_ObjC: case TY_PP_ObjC: case TY_PP_ObjC_Alias:
   case TY_ObjCXX: case TY_PP_ObjCXX:
   case TY_ObjCHeader: case TY_PP_ObjCHeader:
   case TY_ObjCXXHeader: case TY_PP_ObjCXXHeader: case TY_PP_ObjCXX_Alias:
     return true;
   }
 }
 
-bool types::isOpenCL(ID Id) { return Id == TY_CL; }
+bool types::isOpenCL(ID Id) { return Id == TY_CL || Id == TY_CLCXX; }
 
 bool types::isCXX(ID Id) {
   switch (Id) {
   default:
     return false;
 
   case TY_CXX: case TY_PP_CXX:
   case TY_ObjCXX: case TY_PP_ObjCXX: case TY_PP_ObjCXX_Alias:
return llvm::StringSwitch<types::ID>(Ext)
            .Case("M", TY_ObjCXX)
            .Case("o", TY_Object)
            .Case("S", TY_Asm)
            .Case("s", TY_PP_Asm)
            .Case("bc", TY_LLVM_BC)
            .Case("cc", TY_CXX)
            .Case("CC", TY_CXX)
            .Case("cl", TY_CL)
+           .Case("clcpp", TY_CLCXX)
            .Case("cp", TY_CXX)
            .Case("cu", TY_CUDA)
            .Case("hh", TY_CXXHeader)
            .Case("ii", TY_PP_CXX)
            .Case("ll", TY_LLVM_IR)
            .Case("mi", TY_PP_ObjC)
            .Case("mm", TY_ObjCXX)
            .Case("rs", TY_RenderScript)
ID types::lookupHeaderTypeForSourceType(ID Id) {
   case types::TY_CXX:
   case types::TY_CXXModule:
     return types::TY_CXXHeader;
   case types::TY_ObjC:
     return types::TY_ObjCHeader;
   case types::TY_ObjCXX:
     return types::TY_ObjCXXHeader;
   case types::TY_CL:
+  case types::TY_CLCXX:
     return types::TY_CLHeader;
   }
 }

clang/lib/Frontend/CompilerInvocation.cpp

 
 //===----------------------------------------------------------------------===//
 // Deserialization (from args)
 //===----------------------------------------------------------------------===//
 
 static unsigned getOptimizationLevel(ArgList &Args, InputKind IK,
                                      DiagnosticsEngine &Diags) {
   unsigned DefaultOpt = llvm::CodeGenOpt::None;
-  if (IK.getLanguage() == Language::OpenCL && !Args.hasArg(OPT_cl_opt_disable))
+  if ((IK.getLanguage() == Language::OpenCL ||
+       IK.getLanguage() == Language::OpenCLCXX) &&
+      !Args.hasArg(OPT_cl_opt_disable))
     DefaultOpt = llvm::CodeGenOpt::Default;
 
   if (Arg *A = Args.getLastArg(options::OPT_O_Group)) {
     if (A->getOption().matches(options::OPT_O0))
       return llvm::CodeGenOpt::None;
 
     if (A->getOption().matches(options::OPT_Ofast))
       return llvm::CodeGenOpt::Aggressive;
if (Opts.AuxTargetFeatures.hasValue())
     StringRef Lang;
     switch (Opts.DashX.getLanguage()) {
     case Language::C:
       Lang = "c";
       break;
     case Language::OpenCL:
       Lang = "cl";
       break;
+    case Language::OpenCLCXX:
+      Lang = "clcpp";
+      break;
     case Language::CUDA:
       Lang = "cuda";
       break;
     case Language::HIP:
       Lang = "hip";
       break;
     case Language::CXX:
       Lang = "c++";
if (const Arg *A = Args.getLastArg(OPT_x)) {
     IsHeaderFile = !Preprocessed && !ModuleMap &&
                    XValue != "precompiled-header" &&
                    XValue.consume_back("-header");
 
     // Principal languages.
     DashX = llvm::StringSwitch<InputKind>(XValue)
                 .Case("c", Language::C)
                 .Case("cl", Language::OpenCL)
+                .Case("clcpp", Language::OpenCLCXX)
                 .Case("cuda", Language::CUDA)
                 .Case("hip", Language::HIP)
                 .Case("c++", Language::CXX)
                 .Case("objective-c", Language::ObjC)
                 .Case("objective-c++", Language::ObjCXX)
                 .Case("renderscript", Language::RenderScript)
                 .Default(Language::Unknown);
 
if (LangStd == LangStandard::lang_unspecified) {
     // Based on the base language, pick one.
     switch (IK.getLanguage()) {
     case Language::Unknown:
     case Language::LLVM_IR:
       llvm_unreachable("Invalid input kind!");
     case Language::OpenCL:
       LangStd = LangStandard::lang_opencl10;
       break;
+    case Language::OpenCLCXX:
+      LangStd = LangStandard::lang_openclcpp;
+      break;
     case Language::CUDA:
       LangStd = LangStandard::lang_cuda;
       break;
     case Language::Asm:
     case Language::C:
 #if defined(CLANG_DEFAULT_STD_C)
       LangStd = CLANG_DEFAULT_STD_C;
 #else
case Language::LLVM_IR:
     llvm_unreachable("should not parse language flags for this input");
 
   case Language::C:
   case Language::ObjC:
   case Language::RenderScript:
     return S.getLanguage() == Language::C;
 
   case Language::OpenCL:
-    return S.getLanguage() == Language::OpenCL;
+    return S.getLanguage() == Language::OpenCL ||
+           S.getLanguage() == Language::OpenCLCXX;
+
+  case Language::OpenCLCXX:
+    return S.getLanguage() == Language::OpenCLCXX;
 
   case Language::CXX:
   case Language::ObjCXX:
     return S.getLanguage() == Language::CXX;
 
   case Language::CUDA:
     // FIXME: What -std= values should be permitted for CUDA compilations?
     return S.getLanguage() == Language::CUDA ||
static const StringRef GetInputKindName(InputKind IK) {
   case Language::ObjC:
     return "Objective-C";
   case Language::CXX:
     return "C++";
   case Language::ObjCXX:
     return "Objective-C++";
   case Language::OpenCL:
     return "OpenCL";
+  case Language::OpenCLCXX:
+    return "C++ for OpenCL";
   case Language::CUDA:
     return "CUDA";
   case Language::RenderScript:
     return "RenderScript";
   case Language::HIP:
     return "HIP";
 
   case Language::Asm:

clang/lib/Frontend/FrontendActions.cpp

 
 void PrintPreambleAction::ExecuteAction() {
   switch (getCurrentFileKind().getLanguage()) {
   case Language::C:
   case Language::CXX:
   case Language::ObjC:
   case Language::ObjCXX:
   case Language::OpenCL:
+  case Language::OpenCLCXX:
   case Language::CUDA:
   case Language::HIP:
     break;
 
   case Language::Unknown:
   case Language::Asm:
   case Language::LLVM_IR:
   case Language::RenderScript:

clang/lib/Frontend/FrontendOptions.cpp

return llvm::StringSwitch<InputKind>(Extension)
       .Case("mi", InputKind(Language::ObjC).getPreprocessed())
       .Cases("mm", "M", Language::ObjCXX)
       .Case("mii", InputKind(Language::ObjCXX).getPreprocessed())
       .Cases("C", "cc", "cp", Language::CXX)
       .Cases("cpp", "CPP", "c++", "cxx", "hpp", "hxx", Language::CXX)
       .Case("cppm", Language::CXX)
       .Case("iim", InputKind(Language::CXX).getPreprocessed())
       .Case("cl", Language::OpenCL)
+      .Case("clcpp", Language::OpenCLCXX)
       .Cases("cu", "cuh", Language::CUDA)
       .Case("hip", Language::HIP)
       .Cases("ll", "bc", Language::LLVM_IR)
       .Default(Language::Unknown);
 }

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

-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -O0 -emit-llvm -o - | FileCheck %s -check-prefixes=COMMON,PTR
-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -O0 -emit-llvm -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s -check-prefixes=COMMON,PTR
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -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 -triple spir-unknown-unknown -cl-std=clc++ -O0 -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -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 -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -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 -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir -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 -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir -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 -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s --check-prefixes="COMMON,IMPL"
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL | FileCheck %s --check-prefixes="COMMON,EXPL"
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -pedantic -verify -O0 -o - -DDECL -DUSE_DEFLT | FileCheck %s --check-prefixes="COMMON,IMPL"
+// RUN: %clang_cc1 %s -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 -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -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++ -triple spir -emit-llvm -o - -O0 | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -emit-llvm -o - -O0 | 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 -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CLC++ -emit-llvm -O0 -o - | FileCheck %s --check-prefix=CHECK-DEFINITIONS
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -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 -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -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/CodeGenOpenCLCXX/atexit.clcpp

-//RUN: %clang_cc1 %s -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
 
 struct S {
   ~S(){};
 };
 S s;
 
 //CHECK-LABEL: define internal spir_func void @__cxx_global_var_init()
 //CHECK: call spir_func i32 @__cxa_atexit(void (i8*)* bitcast (void (%struct.S addrspace(4)*)* @_ZNU3AS41SD1Ev to void (i8*)*), i8* null, i8 addrspace(1)* @__dso_handle)
 
 //CHECK: declare spir_func i32 @__cxa_atexit(void (i8*)*, i8*, i8 addrspace(1)*)

clang/test/CodeGenOpenCLCXX/constexpr.clcpp

-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -O0 -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -O0 -emit-llvm -o - | FileCheck %s
 
 typedef int int2 __attribute__((ext_vector_type(2)));
 typedef int int4 __attribute__((ext_vector_type(4)));
 
 struct Storage final {
   constexpr const float& operator[](const int index) const noexcept {
     return InternalStorage[index];
   }

clang/test/CodeGenOpenCLCXX/global_init.clcpp

-// RUN: %clang_cc1 %s -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
 
 struct S {
   S() {}
 };
 
 S s;
 
 //CHECK: define internal spir_kernel void @_GLOBAL__sub_I_{{.*}}!kernel_arg_addr_space [[ARGMD:![0-9]+]] !kernel_arg_access_qual [[ARGMD]] !kernel_arg_type [[ARGMD]] !kernel_arg_base_type [[ARGMD]] !kernel_arg_type_qual [[ARGMD]]
 // Check that parameters are empty.
 //CHECK: [[ARGMD]] = !{}

clang/test/CodeGenOpenCLCXX/local_addrspace_init.clcpp

-// RUN: %clang_cc1 %s -triple spir -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+// RUN: %clang_cc1 %s -triple spir -emit-llvm -O0 -o - | FileCheck %s
 
 // Test that we don't initialize local address space objects.
 //CHECK: @_ZZ4testE1i = internal addrspace(3) global i32 undef
 //CHECK: @_ZZ4testE2ii = internal addrspace(3) global %class.C undef
 class C {
   int i;
 };
 

clang/test/CodeGenOpenCLCXX/method-overload-address-space.clcpp

-//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -emit-llvm -O0 -o - | FileCheck %s
+//RUN: %clang_cc1 %s -triple spir-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
 
 struct C {
   void foo() __local;
   void foo() __global;
   void foo();
   void bar();
 };
 

clang/test/CodeGenOpenCLCXX/template-address-spaces.clcpp

-// RUN: %clang_cc1 -cl-std=clc++ %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
 
 template <typename T>
 struct S{
   T a;
   T foo();
 };
 
 template<typename T>

clang/test/Driver/cxx_for_opencl.clcpp

+// RUN: %clang %s -Xclang -verify -fsyntax-only
+// RUN: %clang %s -cl-std=clc++ -Xclang -verify -fsyntax-only
+// RUN: %clang %s -cl-std=cl2.0 -Xclang -verify -fsyntax-only
+// RUN: %clang %s -### 2>&1 | FileCheck %s
+
+// CHECK: "-x" "clcpp"
+
+#ifdef __OPENCL_CPP_VERSION__
+//expected-no-diagnostics
+#endif
+
+kernel void k(){
+  auto a = get_local_id(1);
+#ifndef __OPENCL_CPP_VERSION__
+//expected-error@-2{{OpenCL C version 2.0 does not support the 'auto' storage class specifier}}
+//expected-warning@-3{{type specifier missing, defaults to 'int'}}
+#endif
+}

clang/test/Driver/lit.local.cfg

 config.suffixes = ['.c', '.cpp', '.h', '.m', '.mm', '.S', '.s', '.f90', '.F90', '.f95',
-                   '.cu', '.rs', '.cl', '.hip']
+                   '.cu', '.rs', '.cl', '.clcpp', '.hip']
 config.substitutions = list(config.substitutions)
 config.substitutions.insert(0,
     ('%clang_cc1',
      """*** Do not use 'clang -cc1' in Driver tests. ***""") )
 
 # Remove harmful environmental variables for clang Driver tests.
 # Some might be useful for other tests so they are only removed here.
 driver_overwrite_env_vars = ['MACOSX_DEPLOYMENT_TARGET',
                              'IPHONEOS_DEPLOYMENT_TARGET',
                              'SDKROOT', 'CCC_OVERRIDE_OPTIONS',
                              'CC_PRINT_OPTIONS', 'CC_PRINT_HEADERS',
                              'CC_LOG_DIAGNOSTICS', 'CC_PRINT_PROC_STAT']
 
 for name in driver_overwrite_env_vars:
   if name in config.environment:
     del config.environment[name]

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

-//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
+//RUN: %clang_cc1 %s -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++ -pedantic -verify
+// RUN: %clang_cc1 %s -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++ -pedantic -ast-dump -verify | FileCheck %s
+//RUN: %clang_cc1 %s -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++ -pedantic -ast-dump -verify | FileCheck %s
-//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify -triple i386-windows | FileCheck %s
+//RUN: %clang_cc1 %s -pedantic -ast-dump -verify | FileCheck %s
+//RUN: %clang_cc1 %s -pedantic -ast-dump -verify -triple i386-windows | 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++;};

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

-//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify
+//RUN: %clang_cc1 %s -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 -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -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 -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -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++ -pedantic -verify -ast-dump  | FileCheck %s
+//RUN: %clang_cc1 %s -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 -verify -pedantic -fsyntax-only -cl-std=clc++
+// RUN: %clang_cc1 %s -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++ -pedantic -ast-dump -verify | FileCheck %s
+//RUN: %clang_cc1 %s -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_cast.clcpp

-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -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 -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify -ast-dump | FileCheck %s
+// RUN: %clang_cc1 %s -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);
 }

clang/test/SemaOpenCLCXX/invalid-kernel.clcpp

-// RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -pedantic -verify -fsyntax-only
 
 struct C {
   kernel void m(); //expected-error{{kernel functions cannot be class members}}
 };
 
 template <typename T>
 kernel void templ(T par) { //expected-error{{kernel functions cannot be used in a template declaration, instantiation or specialization}}
 }
 
 template <int>
 kernel void bar(int par) { //expected-error{{kernel functions cannot be used in a template declaration, instantiation or specialization}}
 }
 
 kernel void foo(int); //expected-note{{previous declaration is here}}
 
 kernel void foo(float); //expected-error{{conflicting types for 'foo'}}

clang/test/SemaOpenCLCXX/members.clcpp

-//RUN: %clang_cc1 %s -triple spir -cl-std=clc++ -verify -fsyntax-only
-//RUN: %clang_cc1 %s -triple spir -cl-std=clc++ -verify -fsyntax-only -DFUNCPTREXT
+//RUN: %clang_cc1 %s -triple spir -verify -fsyntax-only
+//RUN: %clang_cc1 %s -triple spir -verify -fsyntax-only -DFUNCPTREXT
 
 #ifdef FUNCPTREXT
 #pragma OPENCL EXTENSION __cl_clang_function_pointers : enable
 //expected-no-diagnostics
 #endif
 
 // Check that pointer to member functions are diagnosed
 // unless specific clang extension is enabled.

clang/test/SemaOpenCLCXX/method-overload-address-space.clcpp

-//RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify
+//RUN: %clang_cc1 %s -triple spir-unknown-unknown -pedantic -verify
 
 struct C {
   void m1() __local __local; //expected-warning{{multiple identical address spaces specified for type}}
   //expected-note@-1{{candidate function}}
   void m1() __global;
   //expected-note@-1{{candidate function}}
   void m2() __global __local; //expected-error{{multiple address spaces specified for type}}
 };

clang/test/SemaOpenCLCXX/newdelete.clcpp

-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
 
 class A {
   public:
   A() : x(21) {}
   int x;
 };
 
 typedef __SIZE_TYPE__ size_t;

clang/test/SemaOpenCLCXX/references.clcpp

-//RUN: %clang_cc1 %s -cl-std=clc++ -verify -fsyntax-only -triple spir
-//RUN: %clang_cc1 %s -cl-std=clc++ -verify -fsyntax-only -DFPTREXT -triple spir
+//RUN: %clang_cc1 %s -verify -fsyntax-only -triple spir
+//RUN: %clang_cc1 %s -verify -fsyntax-only -DFPTREXT -triple spir
 
 #ifdef FPTREXT
 #pragma OPENCL EXTENSION __cl_clang_function_pointers : enable
 #endif // FPTREXT
 
 // References to functions are not allowed.
 struct myclass {
 //FIXME: Here we provide incorrect diagnostic.

clang/test/SemaOpenCLCXX/restricted.clcpp

-// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=clc++ -pedantic -verify -fsyntax-only
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -pedantic -verify -fsyntax-only
 
 // This test checks that various C/C++/OpenCL C constructs are not available in
 // C++ for OpenCL.
 
 // Test that typeid is not available.
 namespace std {
   // Provide a dummy std::type_info so that we can use typeid.
   class type_info {

clang/test/lit.cfg.py

 # testFormat: The test format to use to interpret tests.
 #
 # For now we require '&&' between commands, until they get globally killed and
 # the test runner updated.
 config.test_format = lit.formats.ShTest(not llvm_config.use_lit_shell)
 
 # suffixes: A list of file extensions to treat as test files.
 config.suffixes = ['.c', '.cpp', '.i', '.cppm', '.m', '.mm', '.cu',
-                   '.ll', '.cl', '.s', '.S', '.modulemap', '.test', '.rs', '.ifs']
+                   '.ll', '.cl', '.clcpp', '.s', '.S', '.modulemap', '.test', '.rs', '.ifs']
 
 # excludes: A list of directories to exclude from the testsuite. The 'Inputs'
 # subdirectories contain auxiliary inputs for various tests in their parent
 # directories.
 config.excludes = ['Inputs', 'CMakeLists.txt', 'README.txt', 'LICENSE.txt', 'debuginfo-tests']
 
 # test_source_root: The root path where tests are located.
 config.test_source_root = os.path.dirname(__file__)