[PATCH] D80932: [SYCL] Make default address space a superset of OpenCL address spaces.

[Alexey Bader via Phabricator via cfe-commits]

bader created this revision.
bader added reviewers: Anastasia, keryell, Naghasan, asavonic, Fznamznon.
Herald added subscribers: cfe-commits, ebevhan, yaxunl.
Herald added a project: clang.

This change enables conversion between OpenCL address spaces and default
address space similar to conversions between OpenCL generic and OpenCL
global/local/private address spaces.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D80932

Files:
  clang/include/clang/AST/Type.h
  clang/test/SemaOpenCLCXX/address-space-lambda.cl
  clang/test/SemaOpenCLCXX/pr-45472.cpp
  clang/test/SemaSYCL/address-space-parameter-conversions.cpp


Index: clang/test/SemaSYCL/address-space-parameter-conversions.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/address-space-parameter-conversions.cpp
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+
+void bar(int &Data) {}
+void bar2(int &Data) {}
+void bar(__attribute__((opencl_private)) int &Data) {}
+void foo(int *Data) {}
+void foo2(int *Data) {}
+void foo(__attribute__((opencl_private)) int *Data) {}
+
+template <typename T>
+void tmpl(T *t) {}
+
+void usages() {
+  __attribute__((opencl_global)) int *GLOB;
+  __attribute__((opencl_private)) int *PRIV;
+  __attribute__((opencl_local)) int *LOC;
+  int *NoAS;
+
+  bar(*GLOB);
+  bar2(*GLOB);
+
+  bar(*PRIV);
+  bar2(*PRIV);
+
+  bar(*NoAS);
+  bar2(*NoAS);
+
+  bar(*LOC);
+  bar2(*LOC);
+
+  foo(GLOB);
+  foo2(GLOB);
+  foo(PRIV);
+  foo2(PRIV);
+  foo(NoAS);
+  foo2(NoAS);
+  foo(LOC);
+  foo2(LOC);
+
+  tmpl(GLOB);
+  tmpl(PRIV);
+  tmpl(NoAS);
+  tmpl(LOC);
+
+  (void)static_cast<int *>(GLOB);
+  (void)static_cast<void *>(GLOB);
+  // FIXME: determine if we can warn on the below conversions.
+  int *i = GLOB;
+  void *v = GLOB;
+  (void)i;
+  (void)v;
+
+  // expected-error@+1{{address space is negative}}
+  __attribute__((address_space(-1))) int *TooLow;
+  // expected-error@+1{{unknown type name '__generic'}}
+  __generic int *IsGeneric;
+}
Index: clang/test/SemaOpenCLCXX/pr-45472.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaOpenCLCXX/pr-45472.cpp
@@ -0,0 +1,9 @@
+// RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
+// FIXME: expected error is not emitted due to wrong address space inference
+// see https://bugs.llvm.org/show_bug.cgi?id=45472 for more details.
+// XFAIL: *
+
+__kernel void test_qual() {
+  auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}}
+  priv3();                       //expected-error{{no matching function for call to object of type}}
+}
\ No newline at end of file
Index: clang/test/SemaOpenCLCXX/address-space-lambda.cl
===================================================================
--- clang/test/SemaOpenCLCXX/address-space-lambda.cl
+++ clang/test/SemaOpenCLCXX/address-space-lambda.cl
@@ -31,8 +31,6 @@
 //CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () {{.*}}const __generic'
   auto priv2 = []() __generic {};
   priv2();
-  auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}}
-  priv3(); //expected-error{{no matching function for call to object of type}}
 
   __constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}} //expected-note{{conversion candidate of type 'void (*)()'}}
   const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}
Index: clang/include/clang/AST/Type.h
===================================================================
--- clang/include/clang/AST/Type.h
+++ clang/include/clang/AST/Type.h
@@ -488,7 +488,12 @@
   /// Returns true if the address space in these qualifiers is equal to or
   /// a superset of the address space in the argument qualifiers.
   bool isAddressSpaceSupersetOf(Qualifiers other) const {
-    return isAddressSpaceSupersetOf(getAddressSpace(), other.getAddressSpace());
+    return isAddressSpaceSupersetOf(getAddressSpace(),
+                                    other.getAddressSpace()) ||
+           (!hasAddressSpace() &&
+            (other.getAddressSpace() == LangAS::opencl_private ||
+             other.getAddressSpace() == LangAS::opencl_local ||
+             other.getAddressSpace() == LangAS::opencl_global));
   }
 
   /// Determines if these qualifiers compatibly include another set.

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits