[PATCH] D69938: [OpenCL] Use __generic addr space when generating internal representation of lambda

Anastasia Stulova via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Wed Nov 27 06:54:55 PST 2019


Anastasia updated this revision to Diff 231247.
Anastasia added a comment.

Added `getDefaultCXXMethodAddrSpace` helper to Sema


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D69938/new/

https://reviews.llvm.org/D69938

Files:
  clang/include/clang/Sema/Sema.h
  clang/lib/Sema/Sema.cpp
  clang/lib/Sema/SemaLambda.cpp
  clang/lib/Sema/SemaType.cpp
  clang/test/SemaOpenCLCXX/address-space-lambda.cl


Index: clang/test/SemaOpenCLCXX/address-space-lambda.cl
===================================================================
--- /dev/null
+++ clang/test/SemaOpenCLCXX/address-space-lambda.cl
@@ -0,0 +1,25 @@
+//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
+
+//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'int (int) const __generic'
+auto glambda = [](auto a) { return a; };
+
+__kernel void foo() {
+  int i;
+//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
+  auto  llambda = [&]() {i++;};
+  llambda();
+  glambda(1);
+  // Test lambda with default parameters
+//CHECK: CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
+  [&] {i++;} ();
+  __constant auto err = [&]() {}; //expected-note-re{{candidate function not viable: address space mismatch in 'this' argument ('__constant (lambda at {{.*}})'), parameter type must be 'const __generic (lambda at {{.*}})'}}
+  err();                          //expected-error-re{{no matching function for call to object of type '__constant (lambda at {{.*}})'}}
+  // FIXME: There is very limited addr space functionality
+  // we can test when taking lambda type from the object.
+  // The limitation is due to addr spaces being added to all
+  // objects in OpenCL. Once we add metaprogramming utility
+  // for removing address spaces from a type we can enhance
+  // testing here.
+  (*(__constant decltype(llambda) *)nullptr)(); //expected-error{{multiple address spaces specified for type}}
+  (*(decltype(llambda) *)nullptr)();
+}
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -4916,7 +4916,9 @@
                           .getScopeRep()
                           ->getKind() == NestedNameSpecifier::TypeSpec) ||
                  state.getDeclarator().getContext() ==
-                     DeclaratorContext::MemberContext;
+                     DeclaratorContext::MemberContext ||
+                 state.getDeclarator().getContext() ==
+                     DeclaratorContext::LambdaExprContext;
         };
 
         if (state.getSema().getLangOpts().OpenCLCPlusPlus && IsClassMember()) {
Index: clang/lib/Sema/SemaLambda.cpp
===================================================================
--- clang/lib/Sema/SemaLambda.cpp
+++ clang/lib/Sema/SemaLambda.cpp
@@ -887,6 +887,10 @@
         /*IsVariadic=*/false, /*IsCXXMethod=*/true));
     EPI.HasTrailingReturn = true;
     EPI.TypeQuals.addConst();
+    LangAS AS = getDefaultCXXMethodAddrSpace();
+    if (AS != LangAS::Default)
+      EPI.TypeQuals.addAddressSpace(getDefaultCXXMethodAddrSpace());
+
     // C++1y [expr.prim.lambda]:
     //   The lambda return type is 'auto', which is replaced by the
     //   trailing-return type if provided and/or deduced from 'return'
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -1279,6 +1279,12 @@
   return nullptr;
 }
 
+LangAS Sema::getDefaultCXXMethodAddrSpace() const {
+  if (getLangOpts().OpenCL)
+    return LangAS::opencl_generic;
+  return LangAS::Default;
+}
+
 void Sema::EmitCurrentDiagnostic(unsigned DiagID) {
   // FIXME: It doesn't make sense to me that DiagID is an incoming argument here
   // and yet we also use the current diag ID on the DiagnosticsEngine. This has
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -3724,6 +3724,9 @@
   /// Add the given method to the list of globally-known methods.
   void addMethodToGlobalList(ObjCMethodList *List, ObjCMethodDecl *Method);
 
+  /// Returns default addr space for method qualifiers.
+  LangAS getDefaultCXXMethodAddrSpace() const;
+
 private:
   /// AddMethodToGlobalPool - Add an instance or factory method to the global
   /// pool. See descriptoin of AddInstanceMethodToGlobalPool.


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D69938.231247.patch
Type: text/x-patch
Size: 4068 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20191127/b92c7716/attachment.bin>


More information about the cfe-commits mailing list