Address code-review comments; add non-deprecated group barrier interface

Author: premanandrao

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -176,16 +176,6 @@ 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>
@@ -421,30 +411,23 @@ 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>());
+// 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 {
 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

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -std=c++17 -internal-isystem %S/Inputs -fdeclare-spirv-builtins %s -emit-llvm -o - | FileCheck %s
+// 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
@@ -11,11 +11,11 @@ int main() {
   sycl::queue q;
 
   q.submit([&](sycl::handler &cgh) {
-    auto kernel_ = [=](sycl::nd_item<1> item) [[intel::sub_group_size(8)]] {
-      item.barrier(sycl::access::fence_space::local_space);
+    auto kernel_ = [=](sycl::group<1> item) [[intel::sub_group_size(8)]] {
     };
 
-    cgh.parallel_for<class kernel_class>(cl::sycl::nd_range<1>(), kernel_);
+    cgh.parallel_for_work_group<class kernel_class>(
+        cl::sycl::range<1>(), cl::sycl::range<1>(), kernel_);
   });
   return 0;
 }