llvm-project/clang/test/SemaSYCL/kernel-attribute.cpp

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

45 lines
2.7 KiB
C++
Raw Normal View History

// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl -fsycl-is-device -verify %s
[SYCL] Add sycl_kernel attribute for accelerated code outlining SYCL is single source offload programming model relying on compiler to separate device code (i.e. offloaded to an accelerator) from the code executed on the host. Here is code example of the SYCL program to demonstrate compiler outlining work: ``` int foo(int x) { return ++x; } int bar(int x) { throw std::exception("CPU code only!"); } ... 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] + foo(42); }); } ... ``` SYCL device compiler must compile lambda expression passed to cl::sycl::handler::parallel_for method and function foo called from this lambda expression for an "accelerator". SYCL device compiler also must ignore bar function as it's not required for offloaded code execution. This patch adds the sycl_kernel attribute, which is used to mark code passed to cl::sycl::handler::parallel_for as "accelerated code". Attribute must be applied to function templates which parameters include at least "kernel name" and "kernel function object". These parameters will be used to establish an ABI between the host application and offloaded part. Reviewers: jlebar, keryell, Naghasan, ABataev, Anastasia, bader, aaron.ballman, rjmccall, rsmith Reviewed By: keryell, bader Subscribers: mgorny, OlegM, ArturGainullin, agozillon, aaron.ballman, ebevhan, Anastasia, cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D60455 Signed-off-by: Alexey Bader <alexey.bader@intel.com>
2019-11-06 22:35:50 +08:00
// Only function templates
[[clang::sycl_kernel]] int gv2 = 0; // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
__attribute__((sycl_kernel)) int gv3 = 0; // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
__attribute__((sycl_kernel)) void foo(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
[[clang::sycl_kernel]] void foo1(); // expected-warning {{'sycl_kernel' attribute only applies to function templates}}
// Attribute takes no arguments
template <typename T, typename A>
__attribute__((sycl_kernel(1))) void foo(T P); // expected-error {{'sycl_kernel' attribute takes no arguments}}
template <typename T, typename A, int I>
[[clang::sycl_kernel(1)]] void foo1(T P);// expected-error {{'sycl_kernel' attribute takes no arguments}}
// At least two template parameters
template <typename T>
__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
template <typename T>
[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{'sycl_kernel' attribute only applies to a function template with at least two template parameters}}
// First two template parameters cannot be non-type template parameters
template <typename T, int A>
__attribute__((sycl_kernel)) void foo(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
template <int A, typename T>
[[clang::sycl_kernel]] void foo1(T P); // expected-warning {{template parameter of a function template with the 'sycl_kernel' attribute cannot be a non-type template parameter}}
// Must return void
template <typename T, typename A>
__attribute__((sycl_kernel)) int foo(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
template <typename T, typename A>
[[clang::sycl_kernel]] int foo1(T P); // expected-warning {{function template with 'sycl_kernel' attribute must have a 'void' return type}}
// Must take at least one argument
template <typename T, typename A>
__attribute__((sycl_kernel)) void foo(); // expected-warning {{function template with 'sycl_kernel' attribute must have a single parameter}}
template <typename T, typename A>
[[clang::sycl_kernel]] void foo1(T t, A a); // expected-warning {{function template with 'sycl_kernel' attribute must have a single parameter}}
// No diagnostics
template <typename T, typename A>
__attribute__((sycl_kernel)) void foo(T P);
template <typename T, typename A, int I>
[[clang::sycl_kernel]] void foo1(T P);