[SYCL] sampler implementation

This patch contains sampler class runtime implementation and compiler
support

Signed-off-by: Sindhu Chittireddy <[email protected]>

Author: schittir

clang/lib/Sema/SemaSYCL.cpp

@@ -63,6 +63,10 @@ class Util {
   /// accessor class.
   static bool isSyclAccessorType(const QualType &Ty);
 
+  /// Checks whether given clang type is a full specialization of the sycl
+  /// sampler class.
+  static bool isSyclSamplerType(const QualType &Ty);
+
   /// Checks whether given clang type is the sycl stream class.
   static bool isSyclStreamType(const QualType &Ty);
 
@@ -516,6 +520,59 @@ CreateSYCLKernelBody(Sema &S, FunctionDecl *KernelCallerFunc, DeclContext *DC) {
         CXXMemberCallExpr *Call = CXXMemberCallExpr::Create(
             S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation());
         BodyStmts.push_back(Call);
+      } else if (CRD && Util::isSyclSamplerType(FieldType)) {
+
+        // Sampler has only one TargetFuncParam, which should be set in
+        // __init method: _ValueType
+        const size_t NumParams = 1;
+        llvm::SmallVector<DeclRefExpr *, NumParams> ParamDREs(NumParams);
+        auto TFP = TargetFuncParam;
+        QualType ParamType = (*TFP)->getOriginalType();
+        ParamDREs[0] = DeclRefExpr::Create(
+            S.Context, NestedNameSpecifierLoc(), SourceLocation(), *TFP,
+            false, DeclarationNameInfo(), ParamType, VK_LValue);
+        DeclAccessPair FieldDAP = DeclAccessPair::make(Field, AS_none);
+
+        // kernel_obj.sampler
+        auto SamplerME = MemberExpr::Create(
+            S.Context, CloneRef, false, SourceLocation(),
+            NestedNameSpecifierLoc(), SourceLocation(), Field, FieldDAP,
+            DeclarationNameInfo(Field->getDeclName(), SourceLocation()),
+            nullptr, Field->getType(), VK_LValue, OK_Ordinary);
+
+        CXXMethodDecl *InitMethod = nullptr;
+        for (auto Method : CRD->methods()) {
+          if (Method->getNameInfo().getName().getAsString() == "__init") {
+            InitMethod = Method;
+            break;
+          }
+        }
+        assert(InitMethod && "The sampler must have the __init method");
+
+        // kernel_obj.sampler.__init
+        DeclAccessPair MethodDAP = DeclAccessPair::make(InitMethod, AS_none);
+        auto ME = MemberExpr::Create(
+            S.Context, SamplerME, false, SourceLocation(),
+            NestedNameSpecifierLoc(), SourceLocation(), InitMethod, MethodDAP,
+            InitMethod->getNameInfo(), nullptr, InitMethod->getType(),
+            VK_LValue, OK_Ordinary);
+
+        // Not referenced -> not emitted
+        S.MarkFunctionReferenced(SourceLocation(), InitMethod, true);
+
+        QualType ResultTy = InitMethod->getReturnType();
+        ExprValueKind VK = Expr::getValueKindForType(ResultTy);
+        ResultTy = ResultTy.getNonLValueExprType(S.Context);
+
+        // __init needs one parameter
+        auto ParamItr = InitMethod->param_begin();
+        // kernel_parameters
+        llvm::SmallVector<Expr *, NumParams> ParamStmts;
+        ParamStmts.push_back(getExprForPointer(
+            S, (*ParamItr)->getOriginalType(), ParamDREs[0]));
+        CXXMemberCallExpr *Call = CXXMemberCallExpr::Create(
+            S.Context, ME, ParamStmts, ResultTy, VK, SourceLocation());
+        BodyStmts.push_back(Call);
       } else if (CRD || FieldType->isScalarType()) {
         // If field have built-in or a structure/class type just initialize
         // this field with corresponding kernel argument using '=' binary
@@ -668,6 +725,15 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
           getFieldDeclByName(RecordDecl, {"impl", "Offset"});
       assert(OffsetFld && "The accessor.impl must contain the Offset field");
       CreateAndAddPrmDsc(OffsetFld, OffsetFld->getType());
+    } else if (Util::isSyclSamplerType(ArgTy)) {
+      // the parameter is a SYCL sampler object
+      const auto *RecordDecl = ArgTy->getAsCXXRecordDecl();
+      assert(RecordDecl && "sampler must be of a record type");
+
+      FieldDecl *ImplFld =
+          getFieldDeclByName(RecordDecl, {"impl", "m_Sampler"});
+      assert(ImplFld && "The sampler must contain impl field");
+      CreateAndAddPrmDsc(ImplFld, ImplFld->getType());
     } else if (Util::isSyclStreamType(ArgTy)) {
       // the parameter is a SYCL stream object
       llvm_unreachable("streams not supported yet");
@@ -724,6 +790,17 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
           AccTmplTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
       int Info = getAccessTarget(AccTmplTy) | (Dims << 11);
       H.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, Offset);
+    } else if (Util::isSyclSamplerType(ArgTy)) {
+      // The parameter is a SYCL sampler object
+      // It has only one descriptor, "m_Sampler"
+      const auto *SamplerTy = ArgTy->getAsCXXRecordDecl();
+      assert(SamplerTy && "sampler must be of a record type");
+      FieldDecl *ImplFld =
+          getFieldDeclByName(SamplerTy, {"impl", "m_Sampler"}, &Offset);
+      uint64_t Sz =
+          Ctx.getTypeSizeInChars(ImplFld->getType()).getQuantity();
+      H.addParamDesc(SYCLIntegrationHeader::kind_sampler,
+                     static_cast<unsigned>(Sz), static_cast<unsigned>(Offset));
     } else if (Util::isSyclStreamType(ArgTy)) {
       // the parameter is a SYCL stream object
       llvm_unreachable("streams not supported yet");
@@ -1125,6 +1202,14 @@ bool Util::isSyclAccessorType(const QualType &Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclSamplerType(const QualType &Ty) {
+  static std::array<DeclContextDesc, 3> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
+      Util::DeclContextDesc{clang::Decl::Kind::CXXRecord, "sampler"}};
+  return matchQualifiedTypeName(Ty, Scopes);
+}
+
 bool Util::isSyclStreamType(const QualType &Ty) {
   static std::array<DeclContextDesc, 3> Scopes = {
       Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -8,7 +8,22 @@
 
 // Dummy runtime classes to model SYCL API.
 namespace cl {
+namespace __spirv {
+class OpTypeSampler;
+}
+
 namespace sycl {
+struct sampler_impl {
+  __spirv::OpTypeSampler* m_Sampler;
+};
+
+class sampler {
+  struct sampler_impl impl;
+  void __init(__spirv::OpTypeSampler* Sampler) { impl.m_Sampler = Sampler; }
+
+public:
+  void use(void) const {}
+};
 
 namespace access {
 

clang/test/CodeGenSYCL/integration_header.cpp

@@ -24,14 +24,17 @@
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 4062, 4 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 7 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 16 },
 // CHECK-EMPTY:
 // CHECK-NEXT:   //--- _ZTSN16second_namespace13second_kernelIcEE
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 8 },
 // CHECK-EMPTY:
 // CHECK-NEXT:   //--- _ZTS12third_kernelILi1Ei5pointIZ4mainE1XEE
 // CHECK-NEXT:   { kernel_param_kind_t::kind_std_layout, 4, 0 },
 // CHECK-NEXT:   { kernel_param_kind_t::kind_accessor, 6112, 4 },
+// CHECK-NEXT:   { kernel_param_kind_t::kind_sampler, 8, 8 },
 // CHECK-EMPTY:
 // CHECK-NEXT: };
 //
@@ -65,6 +68,7 @@ int main() {
                      cl::sycl::access::placeholder::true_t>
       acc2;
   int i = 13;
+  cl::sycl::sampler smplr;
   // TODO: Uncomemnt when structures in kernel arguments are correctly processed
   //       by SYCL compiler
   /*  struct {
@@ -77,17 +81,20 @@ int main() {
 
       acc1.use();
       acc2.use();
+      smplr.use();
     }
   });
 
   kernel_single_task<class second_namespace::second_kernel<char>>([=]() {
     if (i == 13) {
       acc2.use();
+      smplr.use();
     }
   });
   kernel_single_task<class third_kernel<1, int,point<struct X>>>([=]() {
     if (i == 13) {
       acc2.use();
+      smplr.use();
     }
   });
 

clang/test/CodeGenSYCL/sampler.cpp

@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 -triple spir64-unknown-linux-sycldevice -std=c++11 -I %S/Inputs  -fsycl-is-device -disable-llvm-passes -emit-llvm -x c++ %s -o - | FileCheck --enable-var-scope %s
+// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%spirv.Sampler* [[SAMPLER_ARG:%[a-zA-Z0-9_]+]])
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[SAMPLER_ARG]].addr = alloca %spirv.Sampler*, align 8
+// CHECK-NEXT: [[ANON:%[0-9]+]] = alloca %class.anon, align 8
+// CHECK-NEXT: store %spirv.Sampler* [[SAMPLER_ARG]], %spirv.Sampler** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
+// CHECK-NEXT: [[BITCAST:%[0-9]+]] = bitcast %class.anon* [[ANON]] to i8*
+// CHECK-NEXT: call void @llvm.lifetime.start.p0i8(i64 8, i8* [[BITCAST]]) #4
+// CHECK-NEXT: [[GEP:%[0-9]+]]  = getelementptr inbounds %class.anon, %class.anon* [[ANON]], i32 0, i32 0
+// CHECK-NEXT: [[LOAD_SAMPLER_ARG:%[0-9]+]] = load %spirv.Sampler*, %spirv.Sampler** [[SAMPLER_ARG]].addr, align 8, !tbaa !9
+// CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.cl::sycl::sampler"* [[GEP]], %spirv.Sampler* [[LOAD_SAMPLER_ARG]])
+//
+#include "sycl.hpp"
+
+template <typename KernelName, typename KernelType>
+__attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  cl::sycl::sampler smplr;
+  kernel_single_task<class first_kernel>([=]() {
+    smplr.use();
+  });
+
+  return 0;
+}

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -8,6 +8,9 @@
 #endif
 
 namespace cl {
+namespace __spirv {
+class OpTypeSampler;
+}
 namespace sycl {
 namespace access {
 
@@ -65,11 +68,24 @@ class accessor {
   void use(void) const {}
   void use(void*) const {}
   _ImplT<dimensions> impl;
+
 private:
   void __init(__global dataT *Ptr, range<dimensions> AccessRange,
               range<dimensions> MemRange, id<dimensions> Offset) {}
 };
 
+struct sampler_impl {
+  __spirv::OpTypeSampler *m_Sampler;
+};
+
+class sampler {
+  struct sampler_impl impl;
+  void __init(__spirv::OpTypeSampler *Sampler) { impl.m_Sampler = Sampler; }
+
+public:
+  void use(void) const {}
+};
+
 } // namespace sycl
 } // namespace cl
 

clang/test/SemaSYCL/sampler.cpp

@@ -0,0 +1,33 @@
+// RUN: %clang -S -I %S/Inputs --sycl -Xclang -ast-dump %s | FileCheck %s
+
+#include <sycl.hpp>
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  cl::sycl::sampler Sampler;
+  kernel<class use_kernel_for_test>([=]() {
+    Sampler.use();
+  });
+  return 0;
+}
+
+// Check declaration of the test kernel
+// CHECK: FunctionDecl {{.*}}use_kernel_for_test 'void (__spirv::OpTypeSampler *)'
+//
+// Check parameters of the test kernel
+// CHECK: ParmVarDecl {{.*}} used _arg_m_Sampler '__spirv::OpTypeSampler *'
+//
+// Check that sampler field of the test kernel object is initialized using __init method
+// CHECK: CXXMemberCallExpr {{.*}} 'void'
+// CHECK-NEXT: MemberExpr {{.*}} 'void (__spirv::OpTypeSampler *)' lvalue .__init
+// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::sampler':'cl::sycl::sampler' lvalue
+// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}sampler.cpp{{.*}})'
+//
+// Check the parameters of __init method
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' <LValueToRValue>
+// CHECK-NEXT: ImplicitCastExpr {{.*}} '__spirv::OpTypeSampler *' lvalue <NoOp>
+// CHECK-NEXT: DeclRefExpr {{.*}} '__spirv::OpTypeSampler *' lvalue ParmVar {{.*}} '_arg_m_Sampler' '__spirv::OpTypeSampler

sycl/CMakeLists.txt

@@ -135,6 +135,7 @@ add_library("${SYCLLibrary}" SHARED
   "${sourceRootPath}/detail/program_manager/program_manager.cpp"
   "${sourceRootPath}/detail/queue_impl.cpp"
   "${sourceRootPath}/detail/os_util.cpp"
+  "${sourceRootPath}/detail/sampler_impl.cpp"
   "${sourceRootPath}/detail/scheduler/commands.cpp"
   "${sourceRootPath}/detail/scheduler/printers.cpp"
   "${sourceRootPath}/detail/scheduler/scheduler.cpp"
@@ -148,6 +149,7 @@ add_library("${SYCLLibrary}" SHARED
   "${sourceRootPath}/kernel.cpp"
   "${sourceRootPath}/platform.cpp"
   "${sourceRootPath}/queue.cpp"
+  "${sourceRootPath}/sampler.cpp"
   "${sourceRootPath}/spirv_ops.cpp"
 )
 

sycl/include/CL/__spirv/spirv_types.hpp

@@ -43,6 +43,9 @@ enum MemorySemantics {
 // Only in such cases the class is recognized as SPIRV type OpTypeEvent.
 class OpTypeEvent;
 
+// SPIRV type for sampler class
+class OpTypeSampler;
+
 enum GroupOperation { Reduce = 0, InclusiveScan = 1, ExclusiveScan = 2 };
 } // namespace __spirv
 } // namespace cl

sycl/include/CL/sycl.hpp

@@ -33,6 +33,7 @@
 #include <CL/sycl/program.hpp>
 #include <CL/sycl/queue.hpp>
 #include <CL/sycl/range.hpp>
+#include <CL/sycl/sampler.hpp>
 #include <CL/sycl/types.hpp>
 #include <CL/sycl/version.hpp>
 

sycl/include/CL/sycl/detail/sampler_impl.hpp

@@ -0,0 +1,64 @@
+//==----------------- sampler_impl.hpp - SYCL standard header file ---------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/__spirv/spirv_types.hpp>
+#include <CL/sycl/context.hpp>
+
+#include <unordered_map>
+
+namespace cl {
+namespace sycl {
+
+enum class addressing_mode : unsigned int;
+enum class filtering_mode : unsigned int;
+enum class coordinate_normalization_mode : unsigned int;
+
+namespace detail {
+class sampler_impl {
+public:
+#ifdef __SYCL_DEVICE_ONLY__
+  __spirv::OpTypeSampler *m_Sampler;
+  sampler_impl(__spirv::OpTypeSampler *Sampler) : m_Sampler(Sampler) {}
+#else
+  cl_sampler m_Sampler;
+  context m_SyclContext;
+  std::unordered_map<context, cl_sampler> m_contextToSampler;
+
+private:
+  coordinate_normalization_mode m_CoordNormMode;
+  addressing_mode m_AddrMode;
+  filtering_mode m_FiltMode;
+  bool m_ReleaseSampler;
+
+public:
+  sampler_impl(coordinate_normalization_mode normalizationMode,
+               addressing_mode addressingMode, filtering_mode filteringMode);
+
+  sampler_impl(cl_sampler clSampler, const context &syclContext);
+
+  addressing_mode get_addressing_mode() const;
+
+  filtering_mode get_filtering_mode() const;
+
+  coordinate_normalization_mode get_coordinate_normalization_mode() const;
+
+  cl_sampler getOrCreateSampler(const context &Context);
+#endif
+
+#ifdef __SYCL_DEVICE_ONLY__
+  ~sampler_impl() = default;
+#else
+  ~sampler_impl();
+#endif
+};
+
+} // namespace detail
+} // namespace sycl
+} // namespace cl