[SYCL] Incorrect diagnostic for __spirv calls (#6058)

When the reqd_sub_group_size attribute is applied on a
sycl kernel, we check that SYCL_EXTERNAL functions called
from such kernels also have the same sub_group_size. This
need not be enforced on __spirv intrinsics since they are
mapped tp the equivalent SPIR-V operations.

Author: premanandrao

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);

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<Dims>());
 }
 
+// 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 <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
 kernel_parallel_for_work_group(const KernelType &KernelFunc) {
   KernelFunc(group<Dims>());
+  parallel_for_work_item();
 }
 
 class handler {

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<class kernel_class>(
+        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}