diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index fd700fbd83576..17538d32e47a4 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3719,6 +3719,12 @@ static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel, CalcEffectiveSubGroup(S.Context, S.getLangOpts(), FD)) return; + // No need to validate __spirv routines here since they + // are mapped to the equivalent SPIRV operations. + const IdentifierInfo *II = FD->getIdentifier(); + if (II && II->getName().startswith("__spirv_")) + return; + // Else we need to figure out why they don't match. SourceLocation FDAttrLoc = GetSubGroupLoc(FD); SourceLocation KernelAttrLoc = GetSubGroupLoc(SYCLKernel); diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index 7200b51695d2d..6ecbd802387cf 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -4,6 +4,11 @@ extern "C" int printf(const char* fmt, ...); +#ifdef __SYCL_DEVICE_ONLY__ +__attribute__((convergent)) extern SYCL_EXTERNAL void +__spirv_ControlBarrier(int, int, int) noexcept; +#endif + // Dummy runtime classes to model SYCL API. inline namespace cl { namespace sycl { @@ -399,10 +404,19 @@ kernel_parallel_for(const KernelType &KernelFunc) { KernelFunc(id()); } +// Dummy parallel_for_work_item function to mimic calls from +// parallel_for_work_group. +void parallel_for_work_item() { +#ifdef __SYCL_DEVICE_ONLY__ + __spirv_ControlBarrier(0, 0, 0); +#endif +} + template ATTR_SYCL_KERNEL void kernel_parallel_for_work_group(const KernelType &KernelFunc) { KernelFunc(group()); + parallel_for_work_item(); } class handler { diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp new file mode 100644 index 0000000000000..86f0cf44ebb6f --- /dev/null +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp @@ -0,0 +1,28 @@ +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -fdeclare-spirv-builtins %s -emit-llvm -o - | FileCheck %s + +// Test that when __spirv intrinsics are invoked from kernel functions +// that have a sub_group_size specified, that such invocations don't +// trigger the error diagnostic that the intrinsic routines must also +// marked with the same attribute. + +#include "Inputs/sycl.hpp" + +int main() { + sycl::queue q; + + q.submit([&](sycl::handler &cgh) { + auto kernel_ = [=](sycl::group<1> item) [[intel::sub_group_size(8)]] { + }; + + cgh.parallel_for_work_group( + cl::sycl::range<1>(), cl::sycl::range<1>(), kernel_); + }); + return 0; +} + +// CHECK: define dso_local spir_kernel void @{{.*}}main{{.*}}kernel_class() {{.*}} !intel_reqd_sub_group_size ![[SUBGROUPSIZE:[0-9]+]] +// CHECK: tail call spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}}) + +// CHECK: declare spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}}) + +// CHECK: ![[SUBGROUPSIZE]] = !{i32 8}