Fznamznon updated this revision to Diff 195128.
Fznamznon added a comment.
Applied comments from @bader
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D60455/new/
https://reviews.llvm.org/D60455
Files:
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/lib/Sema/SemaDeclAttr.cpp
clang/test/Misc/pragma-attribute-supported-attributes-list.test
clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
clang/test/SemaSYCL/device-attributes.cpp
Index: clang/test/SemaSYCL/device-attributes.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/device-attributes.cpp
@@ -0,0 +1,14 @@
+// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -verify %s
+
+[[clang::sycl_device]] int gv = 0;
+__attribute((sycl_device)) int gv1 = 0;
+[[clang::sycl_kernel]] int gv2 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
+__attribute((sycl_kernel)) int gv3 = 0; // expected-warning {{'sycl_kernel' attribute only applies to functions}}
+
+__attribute((sycl_kernel)) __attribute((sycl_device)) void foo();
+[[clang::sycl_kernel]] [[clang::sycl_device]] void foo1();
+
+__attribute((sycl_kernel(1))) void foo(); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+__attribute((sycl_device(1))) void foo1(); // expected-error {{'sycl_device' attribute takes no arguments}}
+[[clang::sycl_kernel(1)]] void foo2(); // expected-error {{'sycl_kernel' attribute takes no arguments}}
+[[clang::sycl_device(1)]] void foo3(); // expected-error {{'sycl_device' attribute takes no arguments}}
Index: clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
===================================================================
--- /dev/null
+++ clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
@@ -0,0 +1,18 @@
+// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -verify %s
+// Now pretend that we're compiling a C++ file. There should be warnings.
+// RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c++ %s
+
+#if not defined(__SYCL_DEVICE_ONLY__)
+// expected-warning@+8 {{'sycl_kernel' attribute ignored}}
+// expected-warning@+8 {{'sycl_device' attribute ignored}}
+// expected-warning@+8 {{'sycl_kernel' attribute ignored}}
+// expected-warning@+8 {{'sycl_device' attribute ignored}}
+#else
+// expected-no-diagnostics
+#endif
+
+__attribute((sycl_kernel)) void foo();
+__attribute((sycl_device)) void foo1();
+[[clang::sycl_kernel]] void foo2();
+[[clang::sycl_device]] void foo3();
+
Index: clang/test/Misc/pragma-attribute-supported-attributes-list.test
===================================================================
--- clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -124,6 +124,8 @@
// CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
// CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
// CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLDevice (SubjectMatchRule_function, SubjectMatchRule_variable)
+// CHECK-NEXT: SYCLKernel (SubjectMatchRule_function)
// CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
// CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
// CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member)
Index: clang/lib/Sema/SemaDeclAttr.cpp
===================================================================
--- clang/lib/Sema/SemaDeclAttr.cpp
+++ clang/lib/Sema/SemaDeclAttr.cpp
@@ -6755,6 +6755,12 @@
case ParsedAttr::AT_Flatten:
handleSimpleAttribute<FlattenAttr>(S, D, AL);
break;
+ case ParsedAttr::AT_SYCLDevice:
+ handleSimpleAttribute<SYCLDeviceAttr>(S, D, AL);
+ break;
+ case ParsedAttr::AT_SYCLKernel:
+ handleSimpleAttribute<SYCLKernelAttr>(S, D, AL);
+ break;
case ParsedAttr::AT_Format:
handleFormatAttr(S, D, AL);
break;
Index: clang/include/clang/Basic/AttrDocs.td
===================================================================
--- clang/include/clang/Basic/AttrDocs.td
+++ clang/include/clang/Basic/AttrDocs.td
@@ -253,6 +253,50 @@
}];
}
+def SYCLDeviceDocs : Documentation {
+ let Category = DocCatFunction;
+ let Content = [{
+The sycl_device attribute specifies function which is supposed to be compiled
+for the device and cannot be directly called by the host. Here is code example
+of the SYCL program, which demonstrates the need for this attribute:
+
+.. code-block:: c++
+
+ int foo(int x) { return ++x; }
+
+ using namespace cl::sycl;
+ queue Q;
+ buffer<int, 1> a(range<1>{1024});
+ Q.submit([&](handler& cgh) {
+ auto A = a.get_access<access::mode::write>(cgh);
+ cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
+ A[index] = index[0] * 2 + index[1] + foo(42);
+ });
+ }
+
+Code is passed to parallel_for is called "kernel function" and defines some
+entry point to device code i.e. will be called by host in run time. Compiler
+is supposed to traverse all symbols accessible from kernel functions and
+add them to the "device part" of the code i.e. mark these symbols with
+sycl_device attribute. In this code example foo compiler is supposed to mark
+"foo" function with sycl_device attrubute.
+The sycl_device attribute also can be used to mark functions which are called
+from the different translation units, i.e. compiler can't identify it w/o user's
+help.
+ }];
+}
+
+def SYCLKernelDocs : Documentation {
+ let Category = DocCatFunction;
+ let Content = [{
+The sycl_kernel attribute specifies sycl kernel function - function which is
+supposed to be compiled for device and can be directly called by the host. See
+also SYCLDeviceDocs above for code example.
+The sycl_kernel attribute includes functionality of the sycl_device attribute
+and provides additional host accessibility.
+ }];
+}
+
def C11NoReturnDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
Index: clang/include/clang/Basic/Attr.td
===================================================================
--- clang/include/clang/Basic/Attr.td
+++ clang/include/clang/Basic/Attr.td
@@ -291,6 +291,7 @@
def MicrosoftExt : LangOpt<"MicrosoftExt">;
def Borland : LangOpt<"Borland">;
def CUDA : LangOpt<"CUDA">;
+def SYCL : LangOpt<"SYCLIsDevice">;
def COnly : LangOpt<"CPlusPlus", 1>;
def CPlusPlus : LangOpt<"CPlusPlus">;
def OpenCL : LangOpt<"OpenCL">;
@@ -995,6 +996,22 @@
let Documentation = [Undocumented];
}
+def SYCLDevice : InheritableAttr {
+ let Spellings = [GNU<"sycl_device">, CXX11<"clang", "sycl_device">,
+ C2x<"clang", "sycl_device">];
+ let Subjects = SubjectList<[Function, Var]>;
+ let LangOpts = [SYCL];
+ let Documentation = [SYCLDeviceDocs];
+}
+
+def SYCLKernel : InheritableAttr {
+ let Spellings = [GNU<"sycl_kernel">, CXX11<"clang", "sycl_kernel">,
+ C2x<"clang", "sycl_kernel">];
+ let Subjects = SubjectList<[Function]>;
+ let LangOpts = [SYCL];
+ let Documentation = [SYCLKernelDocs];
+}
+
def C11NoReturn : InheritableAttr {
let Spellings = [Keyword<"_Noreturn">];
let Subjects = SubjectList<[Function], ErrorDiag>;
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
