[SYCL] Incorrect diagnostic for __spirv calls

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,12 @@
 
 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 {
@@ -61,6 +67,13 @@ enum class address_space : int {
   constant_space,
   local_space
 };
+
+enum class fence_space {
+  local_space = 0,
+  global_space = 1,
+  global_and_local = 2
+};
+
 } // namespace access
 
 // Dummy aspect enum with limited enumerators
@@ -164,6 +177,16 @@ template <int dim> struct item {
   int Data;
 };
 
+template <int dimensions = 1> class nd_item {
+public:
+  void barrier(access::fence_space accessSpace =
+                   access::fence_space::global_and_local) const {
+#ifdef __SYCL_DEVICE_ONLY__
+    __spirv_ControlBarrier(0, 0, 0);
+#endif
+  }
+};
+
 namespace ext {
 namespace oneapi {
 template <typename... properties>
@@ -399,6 +422,12 @@ kernel_parallel_for(const KernelType &KernelFunc) {
   KernelFunc(id<Dims>());
 }
 
+template <typename KernelName, typename KernelType, int Dims>
+ATTR_SYCL_KERNEL void
+kernel_parallel_for_item(const KernelType &KernelFunc) {
+  KernelFunc(nd_item<Dims>());
+}
+
 template <typename KernelName, typename KernelType, int Dims>
 ATTR_SYCL_KERNEL void
 kernel_parallel_for_work_group(const KernelType &KernelFunc) {
@@ -407,6 +436,16 @@ kernel_parallel_for_work_group(const KernelType &KernelFunc) {
 
 class handler {
 public:
+  template <typename KernelName = auto_name, typename KernelType, int Dims>
+  void parallel_for(nd_range<Dims> numWorkItems, const KernelType &kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_parallel_for_item<NameT, KernelType, Dims>(kernelFunc);
+#else
+    kernelFunc(nd_item<Dims>());
+#endif
+  }
+
   template <typename KernelName = auto_name, typename KernelType, int Dims>
   void parallel_for(range<Dims> numWorkItems, const KernelType &kernelFunc) {
     using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;

clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp

@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -std=c++17 -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() {
+  const int local_size = 8;
+  const int global_size = 1*local_size;
+
+  cl::sycl::queue q;
+  
+  q.submit([&] (cl::sycl::handler &cgh) {
+    auto kernel_ = [=](cl::sycl::nd_item<1> item) [[intel::sub_group_size(8)]] {
+      item.barrier(sycl::access::fence_space::local_space);
+    };
+    
+    cgh.parallel_for<class kernel_class>(cl::sycl::nd_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}