[SYCL][ABI-break] Remove old spec constants extension and implementation (#9874)

`sycl_ext_oneapi_spec_constants` extension was deprecated a year ago in
b185269 and since it is dependent on
`program` class, which was removed from SYCL 2020 and was removed from
our implementation about 1.5 years ago in
e7cc7b0, there is no reason anymore to
keep related implementation bits.

This is technically an ABI break, because our compiler loses ability to
process device code compiled by older versions of the compiler.

Author: AlexeySachkov

clang/include/clang/Basic/Attr.td

@@ -1354,12 +1354,12 @@ def SYCLType: InheritableAttr {
   let Subjects = SubjectList<[CXXRecord, Enum], ErrorDiag>;
   let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
   let Args = [EnumArgument<"Type", "SYCLType",
-                           ["accessor", "local_accessor", "spec_constant",
+                           ["accessor", "local_accessor",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",
                             "stream", "sampler", "host_pipe"],
-                           ["accessor", "local_accessor", "spec_constant",
+                           ["accessor", "local_accessor",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
                             "private_memory", "aspect", "annotated_ptr", "annotated_arg",

clang/lib/Sema/SemaSYCL.cpp

@@ -1257,8 +1257,6 @@ class KernelObjVisitor {
                   QualType FieldTy, HandlerTys &... Handlers) {
     if (isSyclSpecialType(FieldTy, SemaRef))
       KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy);
-    else if (isSyclType(FieldTy, SYCLTypeAttr::spec_constant))
-      KF_FOR_EACH(handleSyclSpecConstantType, Field, FieldTy);
     else if (FieldTy->isStructureOrClassType()) {
       if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
         CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
@@ -1327,10 +1325,6 @@ class SyclKernelFieldHandlerBase {
   }
   virtual bool handleSyclSpecialType(FieldDecl *, QualType) { return true; }
 
-  virtual bool handleSyclSpecConstantType(FieldDecl *, QualType) {
-    return true;
-  }
-
   virtual bool handleStructType(FieldDecl *, QualType) { return true; }
   virtual bool handleUnionType(FieldDecl *, QualType) { return true; }
   virtual bool handleReferenceType(FieldDecl *, QualType) { return true; }
@@ -1826,11 +1820,6 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleSyclSpecConstantType(FieldDecl *, QualType) final {
-    CollectionStack.back() = true;
-    return true;
-  }
-
   bool handlePointerType(FieldDecl *, QualType) final {
     PointerStack.back() = true;
     return true;
@@ -3448,10 +3437,6 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
     return handleSpecialType(BS, Ty);
   }
 
-  bool handleSyclSpecConstantType(FieldDecl *FD, QualType Ty) final {
-    return handleSpecialType(FD, Ty);
-  }
-
   bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
     Expr *PointerRef =
         createPointerParamReferenceExpr(FieldTy, StructDepth != 0);
@@ -3717,21 +3702,6 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
     return true;
   }
 
-  bool handleSyclSpecConstantType(FieldDecl *FD, QualType FieldTy) final {
-    const TemplateArgumentList &TemplateArgs =
-        cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl())
-            ->getTemplateInstantiationArgs();
-    assert(TemplateArgs.size() == 2 &&
-           "Incorrect template args for spec constant type");
-    // Get specialization constant ID type, which is the second template
-    // argument.
-    QualType SpecConstIDTy = TemplateArgs.get(1).getAsType().getCanonicalType();
-    const std::string SpecConstName = SYCLUniqueStableNameExpr::ComputeName(
-        SemaRef.getASTContext(), SpecConstIDTy);
-    Header.addSpecConstant(SpecConstName, SpecConstIDTy);
-    return true;
-  }
-
   bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
     addParam(FD, FieldTy,
              ((StructDepth) ? SYCLIntegrationHeader::kind_std_layout

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -370,20 +370,6 @@ struct get_kernel_name_t<auto_name, Type> {
 namespace ext {
 namespace oneapi {
 namespace experimental {
-template <typename T, typename ID = T>
-class __SYCL_TYPE(spec_constant) spec_constant {
-public:
-  spec_constant() {}
-  spec_constant(T Cst) {}
-
-  T get() const { // explicit access.
-    return T();   // Dummy implementaion.
-  }
-  operator T() const { // implicit conversion.
-    return get();
-  }
-};
-
 #ifdef __SYCL_DEVICE_ONLY__
 #define __SYCL_CONSTANT_AS __attribute__((opencl_constant))
 #else

clang/test/CodeGenSYCL/int_header_spec_const.cpp

@@ -344,20 +344,6 @@ struct get_kernel_name_t<auto_name, Type> {
 namespace ext {
 namespace oneapi {
 namespace experimental {
-template <typename T, typename ID = T>
-class __SYCL_TYPE(spec_constant) spec_constant {
-public:
-  spec_constant() {}
-  spec_constant(T Cst) {}
-
-  T get() const { // explicit access.
-    return T();   // Dummy implementaion.
-  }
-  operator T() const { // implicit conversion.
-    return get();
-  }
-};
-
 #ifdef __SYCL_DEVICE_ONLY__
 #define __SYCL_CONSTANT_AS __attribute__((opencl_constant))
 #else

clang/test/Frontend/Inputs/sycl.hpp

@@ -395,16 +395,6 @@ template <typename T, access::address_space AS> class multi_ptr {
 namespace ext {
 namespace oneapi {
 namespace experimental {
-template <typename T, typename ID = T>
-class __SYCL_TYPE(spec_constant) spec_constant {
-public:
-  spec_constant() {}
-  explicit constexpr spec_constant(T defaultVal) : DefaultValue(defaultVal) {}
-
-private:
-  T DefaultValue;
-};
-
 template <typename T, typename... Props>
 class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) annotated_arg {
   T obj;

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -21,10 +21,6 @@ struct StructWithSampler {
   sycl::sampler sampl;
 };
 
-struct StructWithSpecConst {
-  sycl::ext::oneapi::experimental::spec_constant<int, class f1> SC;
-};
-
 sycl::handler H;
 
 struct StructWithStream {
@@ -136,20 +132,6 @@ int main() {
     // CHECK: FunctionDecl {{.*}}Sampl2{{.*}} 'void (sampler_t, StructNonDecomposed, int)'
   }
 
-  {
-    StructWithArray<StructWithSpecConst> t1;
-    myQueue.submit([&](sycl::handler &h) {
-      h.single_task<class SpecConst1>([=]() { return t1.i; });
-    });
-    // CHECK: FunctionDecl {{.*}}SpecConst{{.*}} 'void (StructNonDecomposed, int)'
-
-    DerivedStruct<StructWithSpecConst> t2;
-    myQueue.submit([&](sycl::handler &h) {
-      h.single_task<class SpecConst2>([=]() { return t2.i; });
-    });
-    // CHECK: FunctionDecl {{.*}}SpecConst2{{.*}} 'void (StructNonDecomposed, int)'
-  }
-
   {
     StructWithArray<StructWithStream> t1;
     myQueue.submit([&](sycl::handler &h) {

clang/test/SemaSYCL/decomposition.cpp

@@ -20,12 +20,6 @@ class [[__sycl_detail__::sycl_type(local_accessor)]] local_accessor {};
 // CHECK: CXXRecordDecl {{.*}} class local_accessor definition
 // CHECK: SYCLTypeAttr {{.*}} local_accessor
 
-class [[__sycl_detail__::sycl_type(spec_constant)]]
-[[__sycl_detail__::sycl_type(spec_constant)]] spec_constant;
-// CHECK: CXXRecordDecl {{.*}} spec_constant
-// CHECK: SYCLTypeAttr {{.*}} spec_constant
-// CHECK-NOT: SYCLTypeAttr {{.*}} spec_constant
-
 template <>
 class [[__sycl_detail__::sycl_type(local_accessor)]] local_accessor <int> {};
 // CHECK: ClassTemplateSpecializationDecl {{.*}} class local_accessor definition

clang/test/SemaSYCL/spec-const-kernel-arg.cpp

@@ -17,11 +17,6 @@ class [[__sycl_detail__::sycl_type()]] B {};
 // expected-error@+1{{'sycl_type' attribute argument 'NotValidType' is not supported}}
 class [[__sycl_detail__::sycl_type(NotValidType)]] C {};
 
-// expected-note@+1{{previous attribute is here}}
-class [[__sycl_detail__::sycl_type(spec_constant)]] spec_constant;
-// expected-error@+1{{attribute 'sycl_type' is already applied with different arguments}}
-class [[__sycl_detail__::sycl_type(accessor)]] spec_constant {};
-
 // expected-error@+2{{attribute 'sycl_type' is already applied with different arguments}}
 // expected-note@+1{{previous attribute is here}}
 class [[__sycl_detail__::sycl_type(group)]] [[__sycl_detail__::sycl_type(accessor)]] group {};