[SYCL] Add handling for wrapped sampler (#1942)

This adds proper initialization for sampler object that is wrapped by a
struct. This is temporary change because it contradicts the OpenCL and
SPIR-V spec, since the struct with sampler opaque type field appears as
kernel argument. We need it now to fix crashes in OpenCL backends.

Author: Fznamznon

clang/lib/Sema/SemaSYCL.cpp

@@ -784,12 +784,14 @@ static void VisitField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
                        Handlers &... handlers) {
   if (Util::isSyclAccessorType(ItemTy))
     KF_FOR_EACH(handleSyclAccessorType, Item, ItemTy);
-  if (Util::isSyclStreamType(ItemTy))
+  else if (Util::isSyclStreamType(ItemTy))
     KF_FOR_EACH(handleSyclStreamType, Item, ItemTy);
-  if (ItemTy->isStructureOrClassType())
+  else if (Util::isSyclSamplerType(ItemTy))
+    KF_FOR_EACH(handleSyclSamplerType, Item, ItemTy);
+  else if (ItemTy->isStructureOrClassType())
     VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(),
                          handlers...);
-  if (ItemTy->isArrayType())
+  else if (ItemTy->isArrayType())
     VisitArrayElements(Item, ItemTy, handlers...);
 }
 
@@ -892,6 +894,9 @@ template <typename Derived> class SyclKernelFieldHandler {
     return true;
   }
   virtual bool handleSyclAccessorType(FieldDecl *, QualType) { return true; }
+  virtual bool handleSyclSamplerType(const CXXBaseSpecifier &, QualType) {
+    return true;
+  }
   virtual bool handleSyclSamplerType(FieldDecl *, QualType) { return true; }
   virtual bool handleSyclSpecConstantType(FieldDecl *, QualType) {
     return true;
@@ -1204,6 +1209,7 @@ class SyclKernelDeclCreator
     return ArrayRef<ParmVarDecl *>(std::begin(Params) + LastParamIndex,
                                    std::end(Params));
   }
+  using SyclKernelFieldHandler::handleSyclSamplerType;
 };
 
 class SyclKernelBodyCreator
@@ -1458,6 +1464,7 @@ class SyclKernelBodyCreator
   }
 
   using SyclKernelFieldHandler::enterStruct;
+  using SyclKernelFieldHandler::handleSyclSamplerType;
   using SyclKernelFieldHandler::leaveStruct;
 };
 
@@ -1607,6 +1614,7 @@ class SyclKernelIntHeaderCreator
     CurOffset -= Layout.getBaseClassOffset(BS.getType()->getAsCXXRecordDecl())
                      .getQuantity();
   }
+  using SyclKernelFieldHandler::handleSyclSamplerType;
 };
 } // namespace
 

clang/test/CodeGenSYCL/sampler.cpp

@@ -11,8 +11,20 @@
 // CHECK-NEXT: [[GEPCAST:%[0-9]+]] = addrspacecast %"class{{.*}}.cl::sycl::sampler"* [[GEP]] to %"class{{.*}}.cl::sycl::sampler" addrspace(4)*
 // CHECK-NEXT: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* [[GEPCAST]], %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG]])
 //
+
+// CHECK: define spir_kernel void @{{[a-zA-Z0-9_]+}}(%struct{{.*}}sampler_wrapper{{.*}} %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED:%[a-zA-Z0-9_]+]])
+// CHECK: [[SAMPLER_ARG_WRAPPED]].addr = alloca %opencl.sampler_t addrspace(2)*, align 8
+// CHECK: store %opencl.sampler_t addrspace(2)* [[SAMPLER_ARG_WRAPPED]], %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8
+// CHECK: [[LOAD_SAMPLER_ARG_WRAPPED:%[0-9]+]] = load %opencl.sampler_t addrspace(2)*, %opencl.sampler_t addrspace(2)** [[SAMPLER_ARG_WRAPPED]].addr, align 8
+// CHECK: call spir_func void @{{[a-zA-Z0-9_]+}}(%"class.{{.*}}.cl::sycl::sampler" addrspace(4)* {{.*}}, %opencl.sampler_t addrspace(2)* [[LOAD_SAMPLER_ARG_WRAPPED]])
+//
 #include "sycl.hpp"
 
+struct sampler_wrapper {
+  cl::sycl::sampler smpl;
+  int a;
+};
+
 template <typename KernelName, typename KernelType>
 __attribute__((sycl_kernel)) void kernel_single_task(KernelType kernelFunc) {
   kernelFunc();
@@ -24,5 +36,10 @@ int main() {
     smplr.use();
   });
 
+  sampler_wrapper wrappedSampler = {smplr, 1};
+  kernel_single_task<class second_kernel>([=]() {
+    wrappedSampler.smpl.use();
+  });
+
   return 0;
 }

sycl/test/basic_tests/sampler/sampler.cpp

@@ -22,6 +22,15 @@ namespace sycl {
 using namespace cl::sycl;
 }
 
+struct SamplerWrapper {
+  SamplerWrapper(sycl::coordinate_normalization_mode Norm,
+                 sycl::addressing_mode Addr, sycl::filtering_mode Filter)
+      : Smpl(Norm, Addr, Filter), A(0) {}
+
+  sycl::sampler Smpl;
+  int A;
+};
+
 int main() {
   // Check constructor from enums
   sycl::sampler A(sycl::coordinate_normalization_mode::unnormalized,
@@ -88,13 +97,18 @@ int main() {
   assert(C == A);
   assert(Hasher(C) != Hasher(B));
 
+  SamplerWrapper WrappedSmplr(
+      sycl::coordinate_normalization_mode::normalized,
+      sycl::addressing_mode::repeat, sycl::filtering_mode::linear);
+
   // Device sampler.
   {
     sycl::queue Queue;
     Queue.submit([&](sycl::handler &cgh) {
       cgh.single_task<class kernel>([=]() {
         sycl::sampler C = A;
         sycl::sampler D(C);
+        sycl::sampler E(WrappedSmplr.Smpl);
       });
     });
   }