[SYCL][ABI Break] Remove "cl" namespace (#6518)

As part of the change, also start using "inline namespace _V1" to allow
possible future ABI-affecting changes.

Author: aelovikov-intel

clang/lib/CodeGen/CGSYCLRuntime.cpp

@@ -44,8 +44,8 @@ static bool isPFWI(const FunctionDecl &FD) {
   if (!MD)
     return false;
   static std::array<Util::DeclContextDesc, 3> Scopes = {
-      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
       Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "_V1"},
       Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, "group"}};
   if (!Util::matchQualifiedTypeName(MD->getParent(), Scopes))
     return false;

clang/lib/Sema/SemaDeclAttr.cpp

@@ -10426,8 +10426,8 @@ bool isDeviceAspectType(const QualType Ty) {
     return false;
 
   std::array<std::pair<Decl::Kind, StringRef>, 3> Scopes = {
-      MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
       MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
+      MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       MakeDeclContextDesc(Decl::Kind::Enum, "aspect")};
 
   const auto *Ctx = cast<DeclContext>(ET->getDecl());

clang/lib/Sema/SemaSYCL.cpp

@@ -925,13 +925,13 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor<MarkWIScopeFnVisitor> {
       return true;
     QualType Ty = Ctx.getRecordType(Call->getRecordDecl());
     if (!Util::isSyclType(Ty, "group", true /*Tmpl*/))
-      // not a member of cl::sycl::group - continue search
+      // not a member of sycl::group - continue search
       return true;
     auto Name = Callee->getName();
     if (((Name != "parallel_for_work_item") && (Name != "wait_for")) ||
         Callee->hasAttr<SYCLScopeAttr>())
       return true;
-    // it is a call to cl::sycl::group::parallel_for_work_item/wait_for -
+    // it is a call to sycl::group::parallel_for_work_item/wait_for -
     // mark the callee
     Callee->addAttr(
         SYCLScopeAttr::CreateImplicit(Ctx, SYCLScopeAttr::Level::WorkItem));
@@ -4215,7 +4215,7 @@ static const char *paramKind2Str(KernelParamKind K) {
 //   VB,
 //     std::array<T1, N>& VC, int param, T2 ... varargs) {
 //     ...
-//     deviceQueue.submit([&](cl::sycl::handler& cgh) {
+//     deviceQueue.submit([&](sycl::handler& cgh) {
 //       ...
 //       cgh.parallel_for<class SimpleVadd<T1, N, T2...>>(...)
 //       ...
@@ -4650,8 +4650,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
       FwdDeclEmitter.Visit(K.NameType);
   O << "\n";
 
-  O << "__SYCL_INLINE_NAMESPACE(cl) {\n";
   O << "namespace sycl {\n";
+  O << "__SYCL_INLINE_VER_NAMESPACE(_V1) {\n";
   O << "namespace detail {\n";
 
   // Generate declaration of variable of type __sycl_device_global_registration
@@ -4795,8 +4795,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   }
   O << "\n";
   O << "} // namespace detail\n";
+  O << "} // __SYCL_INLINE_VER_NAMESPACE(_V1)\n";
   O << "} // namespace sycl\n";
-  O << "} // __SYCL_INLINE_NAMESPACE(cl)\n";
   O << "\n";
 }
 
@@ -5076,8 +5076,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
       DeviceGlobOS << "\");\n";
     } else {
       EmittedFirstSpecConstant = true;
-      OS << "__SYCL_INLINE_NAMESPACE(cl) {\n";
       OS << "namespace sycl {\n";
+      OS << "__SYCL_INLINE_VER_NAMESPACE(_V1) {\n";
       OS << "namespace detail {\n";
       OS << "template<>\n";
       OS << "inline const char *get_spec_constant_symbolic_ID_impl<";
@@ -5095,8 +5095,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
       OS << "\";\n";
       OS << "}\n";
       OS << "} // namespace detail\n";
+      OS << "} // __SYCL_INLINE_VER_NAMESPACE(_V1)\n";
       OS << "} // namespace sycl\n";
-      OS << "} // __SYCL_INLINE_NAMESPACE(cl)\n";
     }
   }
 
@@ -5132,8 +5132,8 @@ bool Util::isSyclSpecialType(const QualType Ty) {
 
 bool Util::isSyclSpecConstantType(QualType Ty) {
   std::array<DeclContextDesc, 6> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental"),
@@ -5144,25 +5144,25 @@ bool Util::isSyclSpecConstantType(QualType Ty) {
 
 bool Util::isSyclSpecIdType(QualType Ty) {
   std::array<DeclContextDesc, 3> Scopes = {
-      Util::MakeDeclContextDesc(clang::Decl::Kind::Namespace, "cl"),
       Util::MakeDeclContextDesc(clang::Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
                                 "specialization_id")};
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
 bool Util::isSyclKernelHandlerType(QualType Ty) {
   std::array<DeclContextDesc, 3> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "kernel_handler")};
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
 bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) {
   std::array<DeclContextDesc, 7> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "property"),
@@ -5174,8 +5174,8 @@ bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) {
 
 bool Util::isSyclBufferLocationType(QualType Ty) {
   std::array<DeclContextDesc, 7> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "intel"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "property"),
@@ -5189,16 +5189,16 @@ bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) {
   Decl::Kind ClassDeclKind =
       Tmpl ? Decl::Kind::ClassTemplateSpecialization : Decl::Kind::CXXRecord;
   std::array<DeclContextDesc, 3> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       Util::MakeDeclContextDesc(ClassDeclKind, Name)};
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
 bool Util::isAccessorPropertyListType(QualType Ty) {
   std::array<DeclContextDesc, 5> Scopes = {
-      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
+      Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
       Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -10,8 +10,8 @@ __spirv_ControlBarrier(int, int, int) noexcept;
 #endif
 
 // Dummy runtime classes to model SYCL API.
-inline namespace cl {
 namespace sycl {
+inline namespace _V1 {
 struct sampler_impl {
 #ifdef __SYCL_DEVICE_ONLY__
   __ocl_sampler_t m_Sampler;
@@ -491,7 +491,7 @@ class __attribute__((sycl_special_class)) stream {
   void __finalize() {}
 
 private:
-  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read_write> Acc;
+  sycl::accessor<char, 1, sycl::access::mode::read_write> Acc;
   int FlushBufferSize;
 };
 
@@ -621,5 +621,5 @@ class image {
   }
 };
 
+} // inline namespace _V1
 } // namespace sycl
-} // namespace cl

clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp

@@ -3,8 +3,8 @@
 // Test which verifies that readonly attribute is generated for unexpected access mode value.
 
 // Dummy library with unexpected access::mode enum value.
-namespace cl {
 namespace sycl {
+inline namespace _V1 {
 
 namespace access {
 
@@ -71,16 +71,16 @@ class __attribute__((sycl_special_class)) accessor {
               range<dimensions> MemRange, id<dimensions> Offset) {}
 };
 
+} // inline namespace _V1
 } // namespace sycl
-} // namespace cl
 
 template <typename name, typename Func>
 __attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
   kernelFunc();
 }
 
 int main() {
-  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read> Acc;
+  sycl::accessor<int, 1, sycl::access::mode::read> Acc;
   // CHECK: spir_kernel{{.*}}fake_kernel
   // CHECK-SAME: readonly
   kernel_single_task<class fake_kernel>([=]() {

clang/test/CodeGenSYCL/accessor-readonly.cpp

@@ -5,9 +5,9 @@
 #include "Inputs/sycl.hpp"
 
 // CHECK-NOT: spir_kernel{{.*}}f0_kernel{{.*}}readonly
-void f0(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
-  myQueue.submit([&](cl::sycl::handler &cgh) {
-    auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
+void f0(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
+  myQueue.submit([&](sycl::handler &cgh) {
+    auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
     cgh.single_task<class f0_kernel>([write_acc] {});
   });
 }
@@ -16,10 +16,10 @@ void f0(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
 // CHECK-NOT: readonly
 // CHECK-SAME: %_arg_write_acc{{.*}}%_arg_write_acc1{{.*}}%_arg_write_acc2{{.*}}%_arg_write_acc3
 // CHECK-SAME:  readonly align 4 %_arg_read_acc
-void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
-  myQueue.submit([&](cl::sycl::handler &cgh) {
-    auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
-    auto read_acc = in_buf.get_access<cl::sycl::access::mode::read>(cgh);
+void f1(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
+  myQueue.submit([&](sycl::handler &cgh) {
+    auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
+    auto read_acc = in_buf.get_access<sycl::access::mode::read>(cgh);
     cgh.single_task<class f1_kernel>([write_acc, read_acc] {});
   });
 }
@@ -28,10 +28,10 @@ void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
 // CHECK-SAME: readonly align 4 %_arg_read_acc
 // CHECK-NOT: readonly
 // CHECK-SAME: %_arg_write_acc
-void f2(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
-  myQueue.submit([&](cl::sycl::handler &cgh) {
-    auto read_acc = in_buf.get_access<cl::sycl::access::mode::read>(cgh);
-    auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
+void f2(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
+  myQueue.submit([&](sycl::handler &cgh) {
+    auto read_acc = in_buf.get_access<sycl::access::mode::read>(cgh);
+    auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
     cgh.single_task<class f2_kernel>([read_acc, write_acc] {});
   });
 }

clang/test/CodeGenSYCL/accessor_inheritance.cpp

@@ -3,26 +3,26 @@
 
 struct Base {
   int A, B;
-  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
+  sycl::accessor<char, 1, sycl::access::mode::read> AccField;
 };
 
 struct Captured : Base,
-                  cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
+                  sycl::accessor<char, 1, sycl::access::mode::read> {
   int C;
 };
 
 int main() {
   Captured Obj;
-  cl::sycl::kernel_single_task<class kernel>(
+  sycl::kernel_single_task<class kernel>(
       [=]() {
         Obj.use();
       });
   return 0;
 }
 
 // Check kernel parameters
-// CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
-// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
+// CHECK: %[[RANGE_TYPE:"struct.*sycl::_V1::range"]]
+// CHECK: %[[ID_TYPE:"struct.*sycl::_V1::id"]]
 // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE6kernel
 // CHECK-SAME: i32 noundef [[ARG_A:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: i32 noundef [[ARG_B:%[a-zA-Z0-9_]+]],
@@ -71,10 +71,10 @@ int main() {
 // Check accessors initialization
 // CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 2
 // Default constructor call
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) {{[^,]*}} [[ACC_FIELD]])
+// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) {{[^,]*}} [[ACC_FIELD]])
 // CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, ptr addrspace(4) [[GEP]], i64 20
 // Default constructor call
-// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(ptr addrspace(4) {{[^,]*}} [[GEP1]])
+// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(ptr addrspace(4) {{[^,]*}} [[GEP1]])
 
 // CHECK C field initialization
 // CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, ptr addrspace(4) [[GEP]], i32 0, i32 2

clang/test/CodeGenSYCL/accessor_no_alias_property.cpp

@@ -5,23 +5,23 @@
 #include "Inputs/sycl.hpp"
 
 int main() {
-  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
-                     cl::sycl::access::target::global_buffer,
-                     cl::sycl::access::placeholder::false_t,
-                     cl::sycl::ext::oneapi::accessor_property_list<
-                         cl::sycl::ext::oneapi::property::no_alias::instance<true>>>
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                     sycl::access::target::global_buffer,
+                     sycl::access::placeholder::false_t,
+                     sycl::ext::oneapi::accessor_property_list<
+                         sycl::ext::oneapi::property::no_alias::instance<true>>>
       accessorA;
 
-  cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
-                     cl::sycl::access::target::global_buffer,
-                     cl::sycl::access::placeholder::false_t,
-                     cl::sycl::ext::oneapi::accessor_property_list<
-                         cl::sycl::ext::intel::property::buffer_location::instance<1>>>
+  sycl::accessor<int, 1, sycl::access::mode::read_write,
+                     sycl::access::target::global_buffer,
+                     sycl::access::placeholder::false_t,
+                     sycl::ext::oneapi::accessor_property_list<
+                         sycl::ext::intel::property::buffer_location::instance<1>>>
       accessorB;
 
   // Check that noalias parameter attribute is emitted when no_alias accessor property is used
   // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_accessorA, {{.*}})
-  cl::sycl::kernel_single_task<class kernel_function1>(
+  sycl::kernel_single_task<class kernel_function1>(
       [=]() {
         accessorA.use();
       });
@@ -30,7 +30,7 @@ int main() {
   // CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function2
   // CHECK-NOT: noalias
   // CHECK-SAME: {
-  cl::sycl::kernel_single_task<class kernel_function2>(
+  sycl::kernel_single_task<class kernel_function2>(
       [=]() {
         accessorB.use();
       });