maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/SemaSYCL/kernel-attribute-on-non-syc...

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

15 lines
488 B
C++
Raw Normal View History

[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
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fsycl-is-device -verify %s
// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -x c++ %s
#ifndef __SYCL_DEVICE_ONLY__
// expected-warning@+7 {{'sycl_kernel' attribute ignored}}
// expected-warning@+8 {{'sycl_kernel' attribute ignored}}
#else
// expected-no-diagnostics
#endif
template <typename T, typename A, int B>
__attribute__((sycl_kernel)) void foo(T P);
template <typename T, typename A, int B>
[[clang::sycl_kernel]] void foo1(T P);