[ESIMD] Fix sampler/stream parameter handling for ESIMD kernels (#5165)

sycl_explicit_simd attribute is crashing when used on a special class.

Author: zahiraam

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -11627,6 +11627,8 @@ def err_sycl_expected_finalize_method : Error<
 def ext_sycl_2020_attr_spelling : ExtWarn<
   "use of attribute %0 is a SYCL 2020 extension">,
   InGroup<Sycl2017Compat>;
+def err_sycl_esimd_not_supported_for_type : Error<
+  "type %0 is not supported in ESIMD context">;
 def err_sycl_taking_address_of_wrong_function : Error<
   "taking address of a function not marked with "
   "'intel::device_indirectly_callable' attribute is not allowed in SYCL device "

clang/lib/Sema/SemaSYCL.cpp

@@ -1566,6 +1566,7 @@ void KernelObjVisitor::visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
 // A type to check the validity of all of the argument types.
 class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   bool IsInvalid = false;
+  bool IsSIMD = false;
   DiagnosticsEngine &Diag;
   // Check whether the object should be disallowed from being copied to kernel.
   // Return true if not copyable, false if copyable.
@@ -1652,6 +1653,10 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     assert(Util::isSyclSpecialType(Ty) &&
            "Should only be called on sycl special class types.");
     const RecordDecl *RecD = Ty->getAsRecordDecl();
+    if (IsSIMD && !Util::isSyclType(Ty, "accessor", true /*Tmp*/))
+      return SemaRef.Diag(Loc.getBegin(),
+                          diag::err_sycl_esimd_not_supported_for_type)
+             << RecD;
     if (const ClassTemplateSpecializationDecl *CTSD =
             dyn_cast<ClassTemplateSpecializationDecl>(RecD)) {
       const TemplateArgumentList &TAL = CTSD->getTemplateArgs();
@@ -1666,8 +1671,9 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   }
 
 public:
-  SyclKernelFieldChecker(Sema &S)
-      : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {}
+  SyclKernelFieldChecker(Sema &S, bool isSIMD)
+      : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()),
+        IsSIMD(isSIMD) {}
   static constexpr const bool VisitNthArrayElement = false;
   bool isValid() { return !IsInvalid; }
 
@@ -2250,7 +2256,9 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
     const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
-        IsSIMD ? InitESIMDMethodName : InitMethodName;
+        (IsSIMD && Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
+            ? InitESIMDMethodName
+            : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
     assert(InitMethod && "The type must have the __init method");
     for (const ParmVarDecl *Param : InitMethod->parameters())
@@ -3549,11 +3557,12 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
   if (KernelObj->isInvalidDecl())
     return;
 
+  bool IsSIMDKernel = isESIMDKernelType(KernelObj);
+
   SyclKernelDecompMarker DecompMarker(*this);
-  SyclKernelFieldChecker FieldChecker(*this);
+  SyclKernelFieldChecker FieldChecker(*this, IsSIMDKernel);
   SyclKernelUnionChecker UnionChecker(*this);
 
-  bool IsSIMDKernel = isESIMDKernelType(KernelObj);
   SyclKernelArgsSizeChecker ArgsSizeChecker(*this, Args[0]->getExprLoc(),
                                             IsSIMDKernel);
 

clang/test/SemaSYCL/esimd-special-class.cpp

@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify -internal-isystem %S/Inputs -fsyntax-only %s
+
+// The test is to ensure that the use of sycl_explicit_simd attribute doesn't
+// crash when used with sampler or stream. Currently samplers/stream are not
+// supported in esimd.
+
+#include "sycl.hpp"
+using namespace cl::sycl;
+void test() {
+
+  queue q;
+
+  q.submit([&](handler &h) {
+    cl::sycl::sampler Smplr;
+    cl::sycl::stream Stream(1024, 128, h);
+    // expected-note@+1{{in instantiation of function template specialization}}
+    h.single_task<class SamplerTester>(
+        // expected-error@+1{{type 'sampler' is not supported in ESIMD context}}
+        [=]() [[intel::sycl_explicit_simd]] { Smplr.use(); });
+
+    // expected-note@+1{{in instantiation of function template specialization}}
+    h.single_task<class StreamTester>(
+        // expected-error@+1{{type 'stream' is not supported in ESIMD context}}
+        [=]() [[intel::sycl_explicit_simd]] { Stream.use(); });
+  });
+}