diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 3a2e2bac9242a..9976ff94251a3 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/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; +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 " diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 760851818f448..dbe9a710093e7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1561,6 +1561,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. @@ -1647,6 +1648,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(RecD)) { const TemplateArgumentList &TAL = CTSD->getTemplateArgs(); @@ -1661,8 +1666,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; } @@ -2247,7 +2253,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()) @@ -3546,11 +3554,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); diff --git a/clang/test/SemaSYCL/esimd-special-class.cpp b/clang/test/SemaSYCL/esimd-special-class.cpp new file mode 100644 index 0000000000000..b58f1094f905a --- /dev/null +++ b/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( + // 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( + // expected-error@+1{{type 'stream' is not supported in ESIMD context}} + [=]() [[intel::sycl_explicit_simd]] { Stream.use(); }); + }); +}