[SYCL] Specialization constants support in the Front End.

kbobrovs · kbobrovs · commit 42dfab40147e · 2020-03-27T09:59:59.000-07:00
1. Detect kernel lambda object captures corresponding to specialization
constants and (a) don't create kernel arguments for them (b) generate
specializations of the SpecConstantInfo structure into the integration
header.

2. Recognize the __sycl_fe_getStableUniqueTypeName intrinsic and replace
it with a string literal uniquely identifying the type of the typename
template parameter to this intrinsic.

3. FE-related changes in the runtime:
- new SpecConstantInfo templated struct for type-&gt;name translation for
  specialization constants used by integration header
- define the __sycl_fe_getStableUniqueTypeName intrinsic

Signed-off-by: Konstantin S Bobrovsky &lt;konstantin.s.bobrovsky@intel.com&gt;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
@@ -339,6 +339,9 @@ class SYCLIntegrationHeader {
   /// invocation descriptor has finished.
   void endKernel();
 
+  /// Registers a specialization constant to emit info for it into the header.
+  void addSpecConstant(StringRef IDName, QualType IDType);
+
 private:
   // Kernel actual parameter descriptor.
   struct KernelParamDesc {
@@ -407,6 +410,13 @@ class SYCLIntegrationHeader {
   /// SYCLIntegrationHeader::startKernel
   SmallVector<KernelDesc, 4> KernelDescs;
 
+  using SpecConstID = std::pair<QualType, std::string>;
+
+  /// Keeps specialization constants met in the translation unit. Maps spec
+  /// constant's ID type to generated unique name. Duplicates are removed at
+  /// integration header emission time.
+  llvm::SmallVector<SpecConstID, 4> SpecConsts;
+
   /// Used for emitting diagnostics.
   DiagnosticsEngine &Diag;
 
diff --git a/clang/lib/CodeGen/CGExpr.cpp b/clang/lib/CodeGen/CGExpr.cpp
@@ -17,6 +17,7 @@
 #include "CGObjCRuntime.h"
 #include "CGOpenMPRuntime.h"
 #include "CGRecordLayout.h"
+#include "CGSYCLRuntime.h"
 #include "CodeGenFunction.h"
 #include "CodeGenModule.h"
 #include "ConstantEmitter.h"
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
@@ -78,6 +78,10 @@ class Util {
   /// \param Tmpl  whether the class is template instantiation or simple record
   static bool isSyclType(const QualType &Ty, StringRef Name, bool Tmpl = false);
 
+  /// Checks whether given clang type is a full specialization of the SYCL
+  /// specialization constant class.
+  static bool isSyclSpecConstantType(const QualType &Ty);
+
   /// Checks whether given clang type is declared in the given hierarchy of
   /// declaration contexts.
   /// \param Ty         the clang type being checked
@@ -773,6 +777,14 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S,
                       getExprForSpecialSYCLObj(FldType, WrapperFld,
                                                WrapperFldCRD, Base,
                                                InitMethodName, BodyStmts);
+                    } else if (Util::isSyclSpecConstantType(FldType)) {
+                      // Specialization constants are "invisible" to the
+                      // kernel argument creation and device-side SYCL object
+                      // materialization infrastructure in this source.
+                      // It is OK not to really materialize them on the kernel
+                      // side, because their only use can be via
+                      // 'spec_const_obj.get()' method, which is translated to
+                      // an intrinsic and 'this' is really never used.
                     } else {
                       // Field is a structure or class so change the wrapper
                       // object and recursively search for accessor field.
@@ -816,6 +828,8 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S,
         InitExprs.push_back(MemberInit.get());
         getExprForSpecialSYCLObj(FieldType, Field, CRD, KernelObjCloneRef,
                                  InitMethodName, BodyStmts);
+      } else if (Util::isSyclSpecConstantType(FieldType)) {
+        // Just skip specialization constants - not part of signature.
       } else if (CRD || FieldType->isScalarType()) {
         // If field has built-in or a structure/class type just initialize
         // this field with corresponding kernel argument using copy
@@ -959,11 +973,13 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
               QualType FldType = WrapperFld->getType();
               if (FldType->isStructureOrClassType()) {
                 if (Util::isSyclAccessorType(FldType)) {
-                  // accessor field is found - create descriptor
+                  // Accessor field is found - create descriptor.
                   createSpecialSYCLObjParamDesc(WrapperFld, FldType);
+                } else if (Util::isSyclSpecConstantType(FldType)) {
+                  // Don't try recursive search below.
                 } else {
-                  // field is some class or struct - recursively check for
-                  // accessor fields
+                  // Field is some class or struct - recursively check for
+                  // accessor fields.
                   createParamDescForWrappedAccessors(WrapperFld, FldType);
                 }
               }
@@ -985,6 +1001,8 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
     QualType ArgTy = Fld->getType();
     if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) {
       createSpecialSYCLObjParamDesc(Fld, ArgTy);
+    } else if (Util::isSyclSpecConstantType(ArgTy)) {
+      // Specialization constants are not added as arguments.
     } else if (ArgTy->isStructureOrClassType()) {
       if (Context.getLangOpts().SYCLStdLayoutKernelParams) {
         if (!ArgTy->isStandardLayoutType()) {
@@ -1127,6 +1145,21 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
       uint64_t Sz = Ctx.getTypeSizeInChars(Fld->getType()).getQuantity();
       H.addParamDesc(SYCLIntegrationHeader::kind_pointer,
                      static_cast<unsigned>(Sz), static_cast<unsigned>(Offset));
+    } else if (Util::isSyclSpecConstantType(ArgTy)) {
+      // Add specialization constant ID to the header.
+      auto *TmplSpec =
+          cast<ClassTemplateSpecializationDecl>(ArgTy->getAsCXXRecordDecl());
+      const TemplateArgumentList *TemplateArgs =
+          &TmplSpec->getTemplateInstantiationArgs();
+      // Get specialization constant ID type, which is the second template
+      // argument.
+      QualType SpecConstIDTy = TypeName::getFullyQualifiedType(
+                                   TemplateArgs->get(1).getAsType(), Ctx, true)
+                                   .getCanonicalType();
+      const std::string SpecConstName = PredefinedExpr::ComputeName(
+          Ctx, PredefinedExpr::UniqueStableNameExpr, SpecConstIDTy);
+      H.addSpecConstant(SpecConstName, SpecConstIDTy);
+      // Spec constant lambda capture does not become a kernel argument.
     } else if (ArgTy->isStructureOrClassType() || ArgTy->isScalarType()) {
       // the parameter is an object of standard layout type or scalar;
       // the check for standard layout is done elsewhere
@@ -1658,6 +1691,13 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
   }
 }
 
+static std::string getCPPTypeString(QualType Ty) {
+  LangOptions LO;
+  PrintingPolicy P(LO);
+  P.SuppressTypedefs = true;
+  return eraseAnonNamespace(Ty.getAsString(P));
+}
+
 void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
@@ -1666,6 +1706,33 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";
 
   O << "\n";
+
+  if (SpecConsts.size() > 0) {
+    // Remove duplicates.
+    std::sort(SpecConsts.begin(), SpecConsts.end(),
+              [](const SpecConstID &SC1, const SpecConstID &SC2) {
+                // Sort by string IDs for stable spec consts order in the
+                // header.
+                return SC1.second.compare(SC2.second) < 0;
+              });
+    SpecConstID *End =
+        std::unique(SpecConsts.begin(), SpecConsts.end(),
+                    [](const SpecConstID &SC1, const SpecConstID &SC2) {
+                      // Here can do faster comparison of types.
+                      return SC1.first == SC2.first;
+                    });
+    O << "// Specialization constants IDs:\n";
+    for (const auto &P : llvm::make_range(SpecConsts.begin(), End)) {
+      std::string CPPName = getCPPTypeString(P.first);
+      O << "template <> struct sycl::detail::SpecConstantInfo<" << CPPName
+        << "> {\n";
+      O << "  static constexpr const char* getName() {\n";
+      O << "    return \"" << P.second << "\";\n";
+      O << "  }\n";
+      O << "};\n";
+    }
+  }
+
   if (!UnnamedLambdaSupport) {
     O << "// Forward declarations of templated kernel function types:\n";
 
@@ -1747,11 +1814,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
         O << "', '" << c;
       O << "'> {\n";
     } else {
-      LangOptions LO;
-      PrintingPolicy P(LO);
-      P.SuppressTypedefs = true;
-      O << "template <> struct KernelInfo<"
-        << eraseAnonNamespace(K.NameType.getAsString(P)) << "> {\n";
+      O << "template <> struct KernelInfo<" << getCPPTypeString(K.NameType)
+        << "> {\n";
     }
     O << "  DLL_LOCAL\n";
     O << "  static constexpr const char* getName() { return \"" << K.Name
@@ -1815,6 +1879,10 @@ void SYCLIntegrationHeader::endKernel() {
   // nop for now
 }
 
+void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) {
+  SpecConsts.emplace_back(std::make_pair(IDType, IDName.str()));
+}
+
 SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag,
                                              bool _UnnamedLambdaSupport)
     : Diag(_Diag), UnnamedLambdaSupport(_UnnamedLambdaSupport) {}
@@ -1835,6 +1903,16 @@ bool Util::isSyclStreamType(const QualType &Ty) {
   return isSyclType(Ty, "stream");
 }
 
+bool Util::isSyclSpecConstantType(const QualType &Ty) {
+  const StringRef &Name = "spec_constant";
+  std::array<DeclContextDesc, 4> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "experimental"},
+      Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, Name}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
   Decl::Kind ClassDeclKind =
       Tmpl ? Decl::Kind::ClassTemplateSpecialization : Decl::Kind::CXXRecord;
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -241,6 +241,23 @@ template <typename Type>
 struct get_kernel_name_t<auto_name, Type> {
   using name = Type;
 };
+
+namespace experimental {
+template <typename T, typename ID = T>
+class 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();
+  }
+};
+} // namespace experimental
+
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 template <typename KernelName = auto_name, typename KernelType>
 ATTR_SYCL_KERNEL void kernel_single_task(KernelType kernelFunc) {
diff --git a/clang/test/CodeGenSYCL/int_header_spec_const.cpp b/clang/test/CodeGenSYCL/int_header_spec_const.cpp
@@ -0,0 +1,64 @@
+// RUN: %clang -I %S/Inputs -fsycl-device-only -Xclang -fsycl-int-header=%t.h %s -c -o kernel.spv
+// RUN: FileCheck -input-file=%t.h %s
+
+#include "sycl.hpp"
+
+// This test verifies proper emission of specialization constants into the
+// integration header.
+
+class SpecializedKernel;
+class MyBoolConst;
+class MyInt8Const;
+class MyUInt8Const;
+class MyInt16Const;
+class MyUInt16Const;
+class MyInt32Const;
+class MyUInt32Const;
+
+class MyFloatConst;
+class MyDoubleConst;
+
+int main() {
+  // Create specialization constants.
+  cl::sycl::experimental::spec_constant<bool, MyBoolConst> i1(false);
+  cl::sycl::experimental::spec_constant<char, MyInt8Const> i8(0);
+  cl::sycl::experimental::spec_constant<unsigned char, MyUInt8Const> ui8(0);
+  cl::sycl::experimental::spec_constant<short, MyInt16Const> i16(0);
+  cl::sycl::experimental::spec_constant<unsigned short, MyUInt16Const> ui16(0);
+  cl::sycl::experimental::spec_constant<int, MyInt32Const> i32(0);
+  // Constant used twice, but there must be single entry in the int header,
+  // otherwise compilation error would be issued.
+  cl::sycl::experimental::spec_constant<int, MyInt32Const> i32_1(0);
+  cl::sycl::experimental::spec_constant<unsigned int, MyUInt32Const> ui32(0);
+  cl::sycl::experimental::spec_constant<float, MyFloatConst> f32(0);
+  cl::sycl::experimental::spec_constant<double, MyDoubleConst> f64(0);
+
+  double val;
+  double *ptr = &val; // to avoid "unused" warnings
+
+  cl::sycl::kernel_single_task<SpecializedKernel>([=]() {
+    *ptr = i1.get() +
+           // CHECK-DAG: template <> struct sycl::detail::SpecConstantInfo<class MyBoolConst> {
+           // CHECK-DAG-NEXT:   static constexpr const char* getName() {
+           // CHECK-DAG-NEXT:     return "_ZTS11MyBoolConst";
+           // CHECK-DAG-NEXT:   }
+           // CHECK-DAG-NEXT: };
+           i8.get() +
+           // CHECK-DAG: return "_ZTS11MyInt8Const";
+           ui8.get() +
+           // CHECK-DAG: return "_ZTS12MyUInt8Const";
+           i16.get() +
+           // CHECK-DAG: return "_ZTS12MyInt16Const";
+           ui16.get() +
+           // CHECK-DAG: return "_ZTS13MyUInt16Const";
+           i32.get() +
+           i32_1.get() +
+           // CHECK-DAG: return "_ZTS12MyInt32Const";
+           ui32.get() +
+           // CHECK-DAG: return "_ZTS13MyUInt32Const";
+           f32.get() +
+           // CHECK-DAG: return "_ZTS12MyFloatConst";
+           f64.get();
+    // CHECK-DAG: return "_ZTS13MyDoubleConst";
+  });
+}
diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp
@@ -4,7 +4,7 @@
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
-//===----------------------------------------------------------------------===////
+//===----------------------------------------------------------------------===//
 
 #pragma once
 
@@ -41,6 +41,11 @@ struct kernel_param_desc_t {
   int offset;
 };
 
+// Translates specialization constant type to its name.
+template <class Name> struct SpecConstantInfo {
+  static constexpr const char *getName() { return ""; }
+};
+
 #ifndef __SYCL_UNNAMED_LAMBDA__
 template <class KernelNameType> struct KernelInfo {
   static constexpr unsigned getNumParams() { return 0; }
diff --git a/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp b/sycl/include/CL/sycl/detail/sycl_fe_intrins.hpp
@@ -0,0 +1,24 @@
+//==------ sycl_fe_intrins.hpp --- SYCL Device Compiler's FE intrinsics ----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// C++ intrinsics recognized by the SYCL device compiler frontend
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#ifdef __SYCL_DEVICE_ONLY__
+
+// Returns a unique string identifying the template parameter type. Stable
+// across device compiler invocations.
+template <typename T> const char *__sycl_fe_getStableUniqueTypeName();
+
+// Get the value of the specialization constant with given name.
+// Post-link tool traces the ID to a string literal it points to and assigns
+// integer ID.
+template <typename T> T __sycl_getSpecConstantValue(const char *ID);
+
+#endif
