diff --git a/clang/include/clang/Basic/DiagnosticFrontendKinds.td b/clang/include/clang/Basic/DiagnosticFrontendKinds.td index 87e5d15999acf..5eacfc7aa41e5 100644 --- a/clang/include/clang/Basic/DiagnosticFrontendKinds.td +++ b/clang/include/clang/Basic/DiagnosticFrontendKinds.td @@ -283,8 +283,8 @@ def err_avx_calling_convention : Error; def warn_sycl_device_has_aspect_mismatch : Warning<"function '%0' uses aspect '%1' not listed in its " - "'sycl::device_has' attribute">, BackendInfo, - InGroup; + "%select{'device_has' property|'sycl::device_has' attribute}2">, + BackendInfo, InGroup; def note_sycl_aspect_propagated_from_call : Note<"propagated from call to function '%0'">, BackendInfo; diff --git a/clang/lib/CodeGen/CodeGenAction.cpp b/clang/lib/CodeGen/CodeGenAction.cpp index 26234f026f1f8..b8863b1df177b 100644 --- a/clang/lib/CodeGen/CodeGenAction.cpp +++ b/clang/lib/CodeGen/CodeGenAction.cpp @@ -867,7 +867,8 @@ void BackendConsumer::AspectMismatchDiagHandler( assert(LocCookie.isValid() && "Invalid location for caller in aspect mismatch diagnostic"); Diags.Report(LocCookie, diag::warn_sycl_device_has_aspect_mismatch) - << llvm::demangle(D.getFunctionName().str()) << D.getAspect(); + << llvm::demangle(D.getFunctionName().str()) << D.getAspect() + << D.isFromDeviceHasAttribute(); for (const std::pair &CalleeInfo : D.getCallChain()) { LocCookie = SourceLocation::getFromRawEncoding(CalleeInfo.second); assert(LocCookie.isValid() && diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 504085a2a36e3..9db43f448a4ba 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -7819,7 +7819,8 @@ void Sema::CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) { for (const auto *Attr : std::vector{ D->getAttr(), D->getAttr(), - D->getAttr()}) + D->getAttr(), + D->getAttr()}) if (Attr) Diag(Attr->getLoc(), diag::warn_sycl_old_and_new_kernel_attributes) << Attr; diff --git a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp index 2b94876cf1644..90ad47a4ae46f 100644 --- a/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp +++ b/clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp @@ -1,8 +1,10 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify %s // Tests that add_ir_attributes_function causes a warning when appearing with // potentially conflicting SYCL attributes. +#include "sycl.hpp" + constexpr const char AttrName1[] = "Attr1"; constexpr const char AttrVal1[] = "Val1"; @@ -20,10 +22,13 @@ int main() { EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2)]] {}); EmptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); - EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); - EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); - EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); EmptyWrapper.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::device_has()]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + EmptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1)]] {}); @@ -32,11 +37,17 @@ int main() { // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_work_group_size(1,2,3)]] {}); // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} - NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); + NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1)]] {}); // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} - NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); + NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2)]] {}); // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} - NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); + NonemptyWrapper.kernel_single_task([]() [[sycl::work_group_size_hint(1,2,3)]] {}); // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} NonemptyWrapper.kernel_single_task([]() [[sycl::reqd_sub_group_size(1)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::device_has()]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu)]] {}); + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + NonemptyWrapper.kernel_single_task([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {}); } diff --git a/llvm/include/llvm/IR/DiagnosticInfo.h b/llvm/include/llvm/IR/DiagnosticInfo.h index 31e6f484bf2ac..122947f75fa84 100644 --- a/llvm/include/llvm/IR/DiagnosticInfo.h +++ b/llvm/include/llvm/IR/DiagnosticInfo.h @@ -1120,7 +1120,7 @@ class DiagnosticInfoDontCall : public DiagnosticInfo { void diagnoseAspectsMismatch(const Function *F, const SmallVector &CallChain, - StringRef Aspect); + StringRef Aspect, bool FromDeviceHasAttribute); // Diagnostic information for SYCL aspects usage mismatch. class DiagnosticInfoAspectsMismatch : public DiagnosticInfo { @@ -1128,15 +1128,16 @@ class DiagnosticInfoAspectsMismatch : public DiagnosticInfo { unsigned LocCookie; llvm::SmallVector, 8> CallChain; StringRef Aspect; + bool FromDeviceHasAttribute; public: DiagnosticInfoAspectsMismatch( StringRef FunctionName, unsigned LocCookie, const llvm::SmallVector, 8> &CallChain, - StringRef Aspect) + StringRef Aspect, bool FromDeviceHasAttribute) : DiagnosticInfo(DK_AspectMismatch, DiagnosticSeverity::DS_Warning), FunctionName(FunctionName), LocCookie(LocCookie), CallChain(CallChain), - Aspect(Aspect) {} + Aspect(Aspect), FromDeviceHasAttribute(FromDeviceHasAttribute) {} StringRef getFunctionName() const { return FunctionName; } unsigned getLocCookie() const { return LocCookie; } const llvm::SmallVector, 8> & @@ -1144,6 +1145,7 @@ class DiagnosticInfoAspectsMismatch : public DiagnosticInfo { return CallChain; } StringRef getAspect() const { return Aspect; } + bool isFromDeviceHasAttribute() const { return FromDeviceHasAttribute; } void print(DiagnosticPrinter &DP) const override; static bool classof(const DiagnosticInfo *DI) { return DI->getKind() == DK_AspectMismatch; diff --git a/llvm/lib/IR/DiagnosticInfo.cpp b/llvm/lib/IR/DiagnosticInfo.cpp index f5433c0221855..dc52a4041bb34 100644 --- a/llvm/lib/IR/DiagnosticInfo.cpp +++ b/llvm/lib/IR/DiagnosticInfo.cpp @@ -449,7 +449,8 @@ void DiagnosticInfoDontCall::print(DiagnosticPrinter &DP) const { void llvm::diagnoseAspectsMismatch(const Function *F, const SmallVector &CallChain, - StringRef Aspect) { + StringRef Aspect, + bool FromDeviceHasAttribute) { unsigned LocCookie = 0; if (MDNode *MD = F->getMetadata("srcloc")) LocCookie = @@ -466,7 +467,7 @@ void llvm::diagnoseAspectsMismatch(const Function *F, } DiagnosticInfoAspectsMismatch D(F->getName(), LocCookie, LoweredCallChain, - Aspect); + Aspect, FromDeviceHasAttribute); F->getContext().diagnose(D); } diff --git a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp index 205a2a2f71c94..687d3305b98c2 100644 --- a/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp +++ b/llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp @@ -246,29 +246,39 @@ AspectsSetTy getAspectsUsedByInstruction(const Instruction &I, using FunctionToAspectsMapTy = DenseMap; using CallGraphTy = DenseMap>; +// Finds the first function in a list that uses a given aspect. Returns nullptr +// if none of the functions satisfy the criteria. +Function *findFirstAspectUsageCallee( + const SmallPtrSetImpl &Callees, + const FunctionToAspectsMapTy &AspectsMap, int Aspect, + SmallPtrSetImpl *Visited = nullptr) { + for (Function *Callee : Callees) { + if (Visited && !Visited->insert(Callee).second) + continue; + + auto AspectIt = AspectsMap.find(Callee); + if (AspectIt != AspectsMap.end() && AspectIt->second.contains(Aspect)) + return Callee; + } + return nullptr; +} + // Constructs an aspect usage chain for a given aspect from the function to the // last callee in the first found chain. void constructAspectUsageChain(const Function *F, const FunctionToAspectsMapTy &AspectsMap, const CallGraphTy &CG, int Aspect, - SmallVector &CallChain, - SmallPtrSet &Visited) { + SmallVectorImpl &CallChain, + SmallPtrSetImpl &Visited) { const auto EdgeIt = CG.find(F); if (EdgeIt == CG.end()) return; - for (Function *Callee : EdgeIt->second) { - if (!Visited.insert(Callee).second) - continue; - - auto AspectIt = AspectsMap.find(Callee); - if (AspectIt == AspectsMap.end() || !AspectIt->second.contains(Aspect)) - continue; - - CallChain.push_back(Callee); - constructAspectUsageChain(Callee, AspectsMap, CG, Aspect, CallChain, - Visited); - break; + if (Function *AspectUsingCallee = findFirstAspectUsageCallee( + EdgeIt->second, AspectsMap, Aspect, &Visited)) { + CallChain.push_back(AspectUsingCallee); + constructAspectUsageChain(AspectUsingCallee, AspectsMap, CG, Aspect, + CallChain, Visited); } } @@ -313,22 +323,33 @@ void validateUsedAspectsForFunctions(const FunctionToAspectsMapTy &Map, continue; Function *F = It.first; - - // Entry points will have their declared aspects from their kernel call. - // To avoid double warnings, we skip them. - if (std::find(EntryPoints.begin(), EntryPoints.end(), F) != - EntryPoints.end()) - continue; - - const MDNode *DeviceHasMD = F->getMetadata("sycl_declared_aspects"); - if (!DeviceHasMD) - continue; - AspectsSetTy DeviceHasAspectSet; - for (size_t I = 0; I != DeviceHasMD->getNumOperands(); ++I) { - const auto *CAM = cast(DeviceHasMD->getOperand(I)); - const Constant *C = CAM->getValue(); - DeviceHasAspectSet.insert(cast(C)->getSExtValue()); + bool OriginatedFromAttribute = true; + if (const MDNode *DeviceHasMD = F->getMetadata("sycl_declared_aspects")) { + // Entry points will have their declared aspects from their kernel call. + // To avoid double warnings, we skip them. + if (is_contained(EntryPoints, F)) + continue; + for (const MDOperand &DeviceHasMDOp : DeviceHasMD->operands()) { + const auto *CAM = cast(DeviceHasMDOp); + const Constant *C = CAM->getValue(); + DeviceHasAspectSet.insert(cast(C)->getSExtValue()); + } + OriginatedFromAttribute = true; + } else if (F->hasFnAttribute("sycl-device-has")) { + Attribute DeviceHasAttr = F->getFnAttribute("sycl-device-has"); + SmallVector AspectValStrs; + DeviceHasAttr.getValueAsString().split( + AspectValStrs, ',', /*MaxSplit=*/-1, /*KeepEmpty=*/false); + for (StringRef AspectValStr : AspectValStrs) { + int AspectVal = -1; + assert(!AspectValStr.getAsInteger(10, AspectVal) && + "Aspect value in sycl-device-has is not an integer."); + DeviceHasAspectSet.insert(AspectVal); + } + OriginatedFromAttribute = false; + } else { + continue; } for (int Aspect : Aspects) { @@ -338,9 +359,19 @@ void validateUsedAspectsForFunctions(const FunctionToAspectsMapTy &Map, [=](auto AspectIt) { return Aspect == AspectIt.second; }); assert(AspectNameIt != AspectValues.end() && "Used aspect is not part of the existing aspects"); + // We may encounter an entry point when using the device_has property. + // In this case we act like the usage came from the first callee to + // avoid repeat warnings on the same line. + Function *AdjustedOriginF = + is_contained(EntryPoints, F) + ? findFirstAspectUsageCallee(CG.find(F)->second, Map, Aspect) + : F; + assert(AdjustedOriginF && + "Adjusted function pointer for aspect usage is null"); SmallVector CallChain = - getAspectUsageChain(F, Map, CG, Aspect); - diagnoseAspectsMismatch(F, CallChain, AspectNameIt->first); + getAspectUsageChain(AdjustedOriginF, Map, CG, Aspect); + diagnoseAspectsMismatch(AdjustedOriginF, CallChain, AspectNameIt->first, + OriginatedFromAttribute); } } } diff --git a/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp new file mode 100644 index 0000000000000..f24e089bed4b1 --- /dev/null +++ b/sycl/test/extensions/properties/properties_kernel_device_has_warning.cpp @@ -0,0 +1,181 @@ +// RUN: %clangxx -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s + +// Tests for warnings when propagated aspects do not match the aspects available +// in a function, as specified through the 'sycl::device_has' property. + +#include + +using namespace sycl; +using namespace sycl::ext::oneapi::experimental; + +// expected-note-re@+1 4 {{propagated from call to function '{{.*}}StructWithFP16::StructWithFP16({{.*}})'}} +struct [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] StructWithFP16 { + int a = 0; +}; + +// expected-note-re@+1 4 {{propagated from call to function '{{.*}}funcWithCPU(int)'}} +[[__sycl_detail__::__uses_aspects__(aspect::cpu)]] int funcWithCPU(int a) { + return 0; +} + +// expected-note-re@+1 4 {{propagated from call to function '{{.*}}funcUsingFP16(int, int, int)'}} +int funcUsingFP16(int a, int b, int c) { + StructWithFP16 s; + s.a = 1; + return s.a; +} + +// expected-note-re@+1 4 {{propagated from call to function '{{.*}}funcUsingFP16AndFP64(int, int)'}} +int funcUsingFP16AndFP64(int a, int b) { + double x = 3.0; + return funcUsingFP16(a, b, (int)x); +} + +// expected-note-re@+1 2 {{propagated from call to function '{{.*}}funcIndirectlyUsingFP16(int, int)'}} +int funcIndirectlyUsingFP16(int a, int b) { return funcUsingFP16(a, b, 1); } + +// expected-warning-re@+2 {{function '{{.*}}funcUsingFP16HasFP64(int)' uses aspect 'fp16' not listed in its 'device_has' property}} +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((device_has)) +int funcUsingFP16HasFP64(int a) { return funcIndirectlyUsingFP16(a, 1); } + +// expected-note-re@+1 4 {{propagated from call to function '{{.*}}funcUsingCPU(int, int, int)'}} +int funcUsingCPU(int a, int b, int c) { return funcWithCPU(a); } + +// expected-note-re@+1 4 {{propagated from call to function '{{.*}}funcUsingCPUAndFP64(int, int)'}} +int funcUsingCPUAndFP64(int a, int b) { + double x = 3.0; + return funcUsingCPU(a, b, (int)x); +} + +// expected-note-re@+1 2 {{propagated from call to function '{{.*}}funcIndirectlyUsingCPU(int, int)'}} +int funcIndirectlyUsingCPU(int a, int b) { return funcUsingCPU(a, b, 1); } + +// expected-warning-re@+2 {{function '{{.*}}funcUsingCPUHasFP64(int)' uses aspect 'cpu' not listed in its 'device_has' property}} +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((device_has)) +int funcUsingCPUHasFP64(int a) { return funcIndirectlyUsingCPU(a, 1); } + +int main() { + queue Q; + Q.submit([&](handler &CGH) { + CGH.single_task([=]() { int a = funcUsingFP16HasFP64(1); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingFP16(1, 2); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+3 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp16' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, [=]() { + if constexpr (false) { + int a = funcUsingFP16AndFP64(1, 2); + } + }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task( + properties{device_has}, + [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task( + properties{device_has}, + [=]() { int a = funcUsingFP16AndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{}, [=]() { int a = funcUsingCPUHasFP64(1); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcIndirectlyUsingCPU(1, 2); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+3 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + // expected-warning-re@+2 {{function '{{.*}}' uses aspect 'cpu' not listed in its 'device_has' property}} + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, [=]() { + if constexpr (false) { + int a = funcUsingCPUAndFP64(1, 2); + } + }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task(properties{device_has}, + [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task( + properties{device_has}, + [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + }); + Q.submit([&](handler &CGH) { + CGH.single_task( + properties{device_has}, + [=]() { int a = funcUsingCPUAndFP64(1, 2); }); + }); +} diff --git a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp index ed7705c484735..d451e319a2670 100644 --- a/sycl/test/extensions/properties/properties_kernel_negative_device.cpp +++ b/sycl/test/extensions/properties/properties_kernel_negative_device.cpp @@ -16,6 +16,20 @@ template struct KernelFunctorWithWGSizeWithAttr { } }; +struct KernelFunctorWithOnlyWGSizeHintAttr { + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + void operator() [[sycl::work_group_size_hint(32)]] () const {} +}; + +template struct KernelFunctorWithWGSizeHintWithAttr { + // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + void operator() [[sycl::work_group_size_hint(32)]] () const {} + auto get(sycl::ext::oneapi::experimental::properties_tag) { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size_hint}; + } +}; + struct KernelFunctorWithOnlySGSizeAttr { // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}} void operator() [[sycl::reqd_sub_group_size(32)]] () const {} @@ -30,6 +44,20 @@ template struct KernelFunctorWithSGSizeWithAttr { } }; +struct KernelFunctorWithOnlyDeviceHasAttr { + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + void operator() [[sycl::device_has(sycl::aspect::cpu)]] () const {} +}; + +template struct KernelFunctorWithDeviceHasWithAttr { + // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + void operator() [[sycl::device_has(sycl::aspect::cpu)]] () const {} + auto get(sycl::ext::oneapi::experimental::properties_tag) { + return sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::device_has}; + } +}; + void check_work_group_size() { sycl::queue Q; @@ -47,6 +75,23 @@ void check_work_group_size() { Q.single_task(KernelFunctorWithWGSizeWithAttr<1>{}); } +void check_work_group_size_hint() { + sycl::queue Q; + + // expected-warning@+4 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size_hint<1>}, + []() [[sycl::work_group_size_hint(32)]] {}); + + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::work_group_size_hint<1>}, + KernelFunctorWithOnlyWGSizeHintAttr{}); + + Q.single_task(KernelFunctorWithWGSizeHintWithAttr<1>{}); +} + void check_sub_group_size() { sycl::queue Q; @@ -64,8 +109,27 @@ void check_sub_group_size() { Q.single_task(KernelFunctorWithSGSizeWithAttr<1>{}); } +void check_device_has() { + sycl::queue Q; + + // expected-warning@+4 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}} + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::device_has}, + []() [[sycl::device_has(sycl::aspect::cpu)]] {}); + + Q.single_task( + sycl::ext::oneapi::experimental::properties{ + sycl::ext::oneapi::experimental::device_has}, + KernelFunctorWithOnlyDeviceHasAttr{}); + + Q.single_task(KernelFunctorWithDeviceHasWithAttr{}); +} + int main() { check_work_group_size(); + check_work_group_size_hint(); check_sub_group_size(); + check_device_has(); return 0; }