Test changes and improved this_item call detection.

Author: rdeodhar

clang/lib/Sema/SemaSYCL.cpp

@@ -80,10 +80,6 @@ class Util {
   /// stream class.
   static bool isSyclStreamType(const QualType &Ty);
 
-  /// Checks whether given clang type is a full specialization of the SYCL
-  /// item class.
-  static bool isSyclItemType(const QualType &Ty);
-
   /// Checks whether given clang type is a full specialization of the SYCL
   /// half class.
   static bool isSyclHalfType(const QualType &Ty);
@@ -103,10 +99,23 @@ class Util {
   /// \param Tmpl  whether the class is template instantiation or simple record
   static bool isSyclType(const QualType &Ty, StringRef Name, bool Tmpl = false);
 
+  /// Checks whether given function is a standard SYCL API function with given
+  /// name.
+  /// \param FD    the function being checked.
+  /// \param Name  the function name to be checked against.
+  static bool isSyclFunction(const FunctionDecl *FD, StringRef Name);
+
   /// Checks whether given clang type is a full specialization of the SYCL
   /// specialization constant class.
   static bool isSyclSpecConstantType(const QualType &Ty);
 
+  // Checks declaration context hierarchy.
+  /// \param DC     the context of the item to be checked.
+  /// \param Scopes the declaration scopes leading from the item context to the
+  ///               translation unit (excluding the latter)
+  static bool matchContext(const DeclContext *DC,
+                           ArrayRef<Util::DeclContextDesc> Scopes);
+
   /// Checks whether given clang type is declared in the given hierarchy of
   /// declaration contexts.
   /// \param Ty         the clang type being checked
@@ -2736,9 +2745,8 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
       if (!Visited.insert(FD).second)
         continue; // We've already seen this Decl
 
-      if (FD->isFunctionOrMethod() && FD->getIdentifier() &&
-          !FD->getName().empty() && "this_item" == FD->getName() &&
-          Util::isSyclItemType(FD->getReturnType())) {
+      // Check whether this call is to sycl::this_item().
+      if (Util::isSyclFunction(FD, "this_item")) {
         Header.setCallsThisItem(true);
         return;
       }
@@ -4014,10 +4022,6 @@ bool Util::isSyclStreamType(const QualType &Ty) {
   return isSyclType(Ty, "stream");
 }
 
-bool Util::isSyclItemType(const QualType &Ty) {
-  return isSyclType(Ty, "item", true /*Tmpl*/);
-}
-
 bool Util::isSyclHalfType(const QualType &Ty) {
   const StringRef &Name = "half";
   std::array<DeclContextDesc, 5> Scopes = {
@@ -4064,6 +4068,21 @@ bool Util::isSyclType(const QualType &Ty, StringRef Name, bool Tmpl) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclFunction(const FunctionDecl *FD, StringRef Name) {
+  if (!FD->isFunctionOrMethod() || !FD->getIdentifier() ||
+      FD->getName().empty() || Name != FD->getName())
+    return false;
+
+  const DeclContext *DC = FD->getDeclContext();
+  if (DC->isTranslationUnit())
+    return false;
+
+  std::array<DeclContextDesc, 2> Scopes = {
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
+      Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}};
+  return matchContext(DC, Scopes);
+}
+
 bool Util::isAccessorPropertyListType(const QualType &Ty) {
   const StringRef &Name = "accessor_property_list";
   std::array<DeclContextDesc, 4> Scopes = {
@@ -4074,21 +4093,15 @@ bool Util::isAccessorPropertyListType(const QualType &Ty) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
-bool Util::matchQualifiedTypeName(const QualType &Ty,
-                                  ArrayRef<Util::DeclContextDesc> Scopes) {
-  // The idea: check the declaration context chain starting from the type
+bool Util::matchContext(const DeclContext *Ctx,
+                        ArrayRef<Util::DeclContextDesc> Scopes) {
+  // The idea: check the declaration context chain starting from the item
   // itself. At each step check the context is of expected kind
   // (namespace) and name.
-  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
-
-  if (!RecTy)
-    return false; // only classes/structs supported
-  const auto *Ctx = cast<DeclContext>(RecTy);
   StringRef Name = "";
 
   for (const auto &Scope : llvm::reverse(Scopes)) {
     clang::Decl::Kind DK = Ctx->getDeclKind();
-
     if (DK != Scope.first)
       return false;
 
@@ -4102,11 +4115,21 @@ bool Util::matchQualifiedTypeName(const QualType &Ty,
       Name = cast<NamespaceDecl>(Ctx)->getName();
       break;
     default:
-      llvm_unreachable("matchQualifiedTypeName: decl kind not supported");
+      llvm_unreachable("matchContext: decl kind not supported");
     }
     if (Name != Scope.second)
       return false;
     Ctx = Ctx->getParent();
   }
   return Ctx->isTranslationUnit();
 }
+
+bool Util::matchQualifiedTypeName(const QualType &Ty,
+                                  ArrayRef<Util::DeclContextDesc> Scopes) {
+  const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl();
+
+  if (!RecTy)
+    return false; // only classes/structs supported
+  const auto *Ctx = cast<DeclContext>(RecTy);
+  return Util::matchContext(Ctx, Scopes);
+}

clang/test/CodeGenSYCL/parallel_for_this_item.cpp

@@ -80,13 +80,33 @@ using namespace cl::sycl;
 SYCL_EXTERNAL item<1> g() { return this_item<1>(); }
 SYCL_EXTERNAL item<1> f() { return g(); }
 
+// This is a similar-looking this_item function but not the real one.
+template <int Dims> item<Dims> this_item(int i) { return item<1>{i}; }
+
+// This is a method named this_item but not the real one.
+class C {
+public:
+  template <int Dims> item<Dims> this_item() { return item<1>{66}; };
+};
+
 int main() {
   queue myQueue;
   myQueue.submit([&](::handler &cgh) {
-    cgh.parallel_for<class GNU>(range<1>(1), [=](item<1> I) {});
+    // This kernel does not call sycl::this_item
+    cgh.parallel_for<class GNU>(range<1>(1),
+                                [=](item<1> I) { this_item<1>(55); });
+
+    // This kernel calls sycl::this_item
     cgh.parallel_for<class EMU>(range<1>(1),
                                 [=](::item<1> I) { this_item<1>(); });
-    cgh.parallel_for<class OWL>(range<1>(1), [=](id<1> I) {});
+
+    // This kernel does not call sycl::this_item
+    cgh.parallel_for<class OWL>(range<1>(1), [=](id<1> I) {
+      class C c;
+      c.this_item<1>();
+    });
+
+    // This kernel calls sycl::this_item
     cgh.parallel_for<class RAT>(range<1>(1), [=](id<1> I) { f(); });
   });
 

sycl/include/CL/sycl/handler.hpp

@@ -816,7 +816,7 @@ class __SYCL_EXPORT handler {
       };
 
       range<Dims> AdjustedRange = NumWorkItems;
-      AdjustedRange.set_range(NewValX);
+      AdjustedRange.set_range_dim0(NewValX);
 #ifdef __SYCL_DEVICE_ONLY__
       kernel_parallel_for<NameWT, TransformedArgType>(Wrapper);
 #else

sycl/include/CL/sycl/range.hpp

@@ -148,7 +148,7 @@ template <int dimensions = 1> class range : public detail::array<dimensions> {
   friend class detail::Builder;
 
   // Adjust the first dim of the range
-  void set_range(const size_t dim0) { this->common_array[0] = dim0; }
+  void set_range_dim0(const size_t dim0) { this->common_array[0] = dim0; }
 };
 
 #ifdef __cpp_deduction_guides