Page MenuHomePhabricator

[SYCL] Implement SYCL device code outlining
Needs ReviewPublic

Authored by Fznamznon on Apr 9 2019, 5:36 AM.

Details

Summary

SYCL is single source offload programming model so compiler should be
able to separate device code from host code.

Here is code example of the SYCL program demonstrates complier 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] * 2 + index[1] + foo(42);
      });
    }
...

SYCL compiler needs to compile lambda function passed to
cl::sycl::handler::parallel_for method and function foo called from this
lambda function. Compiler also must ignore bar function when we "device" part
of the single source code.

Current approach is to add an attribute SYCL kernel, which SYCL runtime will use
to mark code passed to cl::sycl::handler::parallel_for as "kernel functions".
Obviously runtime library can't mark foo as "device" code - this is a compiler
job: to traverse all symbols accessible from kernel functions and add them to
the "device part" of the code marking them with new attribute SYCL
device.

Event Timeline

There are a very large number of changes, so older changes are hidden. Show Older Changes
Fznamznon retitled this revision from [SYCL] Add support for SYCL device attributes to [SYCL] Implement SYCL device code outlining.May 21 2019, 8:03 AM
Fznamznon edited the summary of this revision. (Show Details)
Fznamznon updated this revision to Diff 200658.May 22 2019, 1:12 AM

Minor fix.

Anastasia added inline comments.May 22 2019, 10:04 AM
clang/include/clang/Basic/Attr.td
1022

Ok, I thought the earlier request was not to add undocumented attributes with the spelling?

Also did __kernel attribute not work at the end?

I can't quite get where the current disconnect comes from but I find it extremely unhelpful.

bader added inline comments.May 22 2019, 11:36 AM
clang/include/clang/Basic/Attr.td
1022

Hi @Anastasia, let me try to help.

Ok, I thought the earlier request was not to add undocumented attributes with the spelling?

Right. @Fznamznon, could you document sycl_kernel attribute, please?

Also did __kernel attribute not work at the end?

Maria left a comment with the summary of our experiment: https://reviews.llvm.org/D60455#1472705. There is a link to pull request, where @keryell and @agozillon expressed preference to have separate SYCL attributes. Let me copy their feedback here:

@keryell :

Thank you for the experiment.
That looks like a straight forward change.
The interesting part is that it does not expose any advantage from reusing OpenCL __kernel marker.... So I am not more convinced that it is the way to go, because we would use any other keyword or attribute and it would be the same...

@agozillon :

Just my two cents, I think a separation of concerns and having separate attributes will simplify things long-term.

While possibly similar just now, the SYCL specification is evolving and may end up targeting more than just OpenCL. So the semantics of the attributes may end up being quite different, even if at the moment the SYCL attribute is there mostly just to mark something for outlining.

If it doesn't then the case for refactoring and merging them in a future patch could be brought up again.

To summarize: we don't have good arguments to justify re-use of OpenCL __kernel keyword for SYCL mode requested by @aaron.ballman here https://reviews.llvm.org/D60455#1469150.

I can't quite get where the current disconnect comes from but I find it extremely unhelpful.

Let me know how I can help here.

Additional note. I've submitted initial version of SYCL compiler design document to the GItHub: https://github.com/intel/llvm/blob/sycl/sycl/doc/SYCL_compiler_and_runtime_design.md. Please, take a look and let me know if you have questions.

Fznamznon added inline comments.May 23 2019, 1:08 AM
clang/include/clang/Basic/Attr.td
1022

Ok, I thought the earlier request was not to add undocumented attributes with the spelling?

Right. @Fznamznon, could you document sycl_kernel attribute, please?

Do we really need add documentation for attribute which is not presented in SYCL spec and used for internal implementation details only because it has spelling?

Design question: since you are not aware what functions are to be run on a device while parsing them, at what point do you plan to diagnose the invalid behavior from the standard C++ i.e. using function pointers in kernel code is likely to cause issues?

Also do you need to outline the data structures too? For example classes used on device are not allowed to have virtual function.

clang/include/clang/Basic/Attr.td
1022
    Ok, I thought the earlier request was not to add undocumented attributes with the spelling?
 
Right. @Fznamznon, could you document sycl_kernel attribute, please?

Do we really need add documentation for attribute which is not presented in SYCL spec and used for internal implementation details only because it has spelling?

You are adding an attribute that is exposed to the programmers that use clang to compile their code, so unless you come up with some way to reject it in the non-toolchain mode it has to be documented. And for clang it will become "hidden" SYCL dialect so absolutely not different to __kernel.

Another aspect to consider is that clang used TypePrinter in diagnostics and even though printing entire function signature is rare it might appear in diagnostics and the programmer should have a way to understand what the "alien" construct is. This is where clang documentation will help.

1022

@keryell :

Thank you for the experiment.
That looks like a straight forward change.
The interesting part is that it does not expose any advantage from reusing OpenCL __kernel marker.... So I am not more convinced that it is the way to go, because we would use any other keyword or attribute and it would be the same...

I don't understand how this conclusions are made on incomplete implementation or even just an initial patch.

The kind of analysis I am missing at the moment is whether you would need to add similar logic for sycl_kernel as we have now for __kernel i.e. did anyone look at the occurrences of kernel handling in the code base to see if it's going to need the same logic or not:

include/clang/Basic/Attr.td:    : SubsetSubject<Function, [{S->hasAttr<OpenCLKernelAttr>()}],
include/clang/Parse/Parser.h:  void ParseOpenCLKernelAttributes(ParsedAttributes &attrs);
lib/AST/Decl.cpp:  if (hasAttr<OpenCLKernelAttr>())
lib/AST/Decl.cpp:  if (hasAttr<OpenCLKernelAttr>())
lib/AST/Decl.cpp:  if (hasAttr<OpenCLKernelAttr>())
lib/CodeGen/CGCall.cpp:  if (TargetDecl && TargetDecl->hasAttr<OpenCLKernelAttr>()) {
lib/CodeGen/CodeGenFunction.cpp:  if (!FD->hasAttr<OpenCLKernelAttr>())
lib/CodeGen/TargetInfo.cpp:    if (FD->hasAttr<OpenCLKernelAttr>()) {
lib/CodeGen/TargetInfo.cpp:    if (FD->hasAttr<OpenCLKernelAttr>()) {
lib/CodeGen/TargetInfo.cpp:  return D->hasAttr<OpenCLKernelAttr>() ||
lib/CodeGen/TargetInfo.cpp:  if (M.getLangOpts().OpenCL && FD->hasAttr<OpenCLKernelAttr>() &&
lib/Parse/ParseDecl.cpp:void Parser::ParseOpenCLKernelAttributes(ParsedAttributes &attrs) {
lib/Parse/ParseDecl.cpp:      ParseOpenCLKernelAttributes(DS.getAttributes());
lib/Sema/SemaDecl.cpp:        if (FD && !FD->hasAttr<OpenCLKernelAttr>()) {
lib/Sema/SemaDecl.cpp:        if (FD && FD->hasAttr<OpenCLKernelAttr>()) {
lib/Sema/SemaDecl.cpp:  if (getLangOpts().OpenCL && NewFD->hasAttr<OpenCLKernelAttr>()) {
lib/Sema/SemaDecl.cpp:        << FD->hasAttr<OpenCLKernelAttr>();
lib/Sema/SemaDecl.cpp:  if (FD->hasAttr<OpenCLKernelAttr>())
lib/Sema/SemaDeclAttr.cpp:    handleSimpleAttribute<OpenCLKernelAttr>(S, D, AL);
lib/Sema/SemaDeclAttr.cpp:  if (!D->hasAttr<OpenCLKernelAttr>()) {

I don't mind either way but I would like the decision to be based on the analysis of clang code base please!

@agozillon :

Just my two cents, I think a separation of concerns and having separate attributes will simplify things long-term.

This can potentially be a fair point!

While possibly similar just now, the SYCL specification is evolving and may end up targeting more than just OpenCL. So the semantics of the attributes may end up being quite different, even if at the moment the SYCL attribute is there mostly just to mark something for outlining.

This is really great! But unless you provide concrete information what the evolution is and what exactly you are trying to achieve and how it affect compiler design there is no way to review your patches.

Let me know how I can help here.

Additional note. I've submitted initial version of SYCL compiler design document to the GItHub: https://github.com/intel/llvm/blob/sycl/sycl/doc/SYCL_compiler_and_runtime_design.md. Please, take a look and let me know if you have questions.

Thanks for sharing! I will try to find time to look into this and provide my feedback if any.

clang/include/clang/Sema/Sema.h
11197

This deserves more explanation. I would suggest to just look at the code around to follow the style!

clang/lib/Sema/SemaSYCL.cpp
24

This is probably not something we can change at this point but I wish we could avoid complexities like this. :(

I think this is also preventing traditional linking of translation units. That is somewhat unfortunate.

It is good direction however to keep this logic in a separate dedicated compilation unit.

I would suggest to document it a bit more including any current limitations/assumption that you can mark under FIXME i.e. does your code handle lambdas yet, what if lambdas are used in function parameters, etc...

clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
2

Another confusion I have at the moment even though it doesn't belong to this patch - isn't SYCL based on C++11?

bader added a comment.May 27 2019, 7:16 AM

Design question: since you are not aware what functions are to be run on a device while parsing them, at what point do you plan to diagnose the invalid behavior from the standard C++ i.e. using function pointers in kernel code is likely to cause issues?

We are going to use DeviceDiagBuilder and related infrastructure implemented in Clang to diagnose device side code errors in OpenMP/CUDA modes.
More details are in the comments here:
https://clang.llvm.org/doxygen/classclang_1_1Sema_1_1DeviceDiagBuilder.html#details

Also do you need to outline the data structures too? For example classes used on device are not allowed to have virtual function.

Yes. This restriction is already implemented in our code base on GitHub.

clang/include/clang/Basic/Attr.td
1022

@Anastasia, I've looked at the occurrences of OpenCLKernelAttr attribute and it looks like the only part useful for SYCL is lib/CodeGen/CodeGenFunction.cpp, which emits OpenCL specific metadata required for SPIR-V translation.

Sema part is mostly not relevant for SYCL mode because SYCL API do not allow the cases currently detected by clang (e.g. constant address space variable declaration in OpenCL kernel scope, naming OpenCL kernel main, etc).
A couple of check that might be useful are:

  • void return type for kernel functions
  • kernel can't be static function

and some of the checks are harmful for proposed implementation (e.g. kernels can't be template functions).

@Anastasia, @keryell, @agozillon and @aaron.ballman need to agree if this sufficient to justify the re-use of OpenCL kernel attribute.
Let me know if you need any additional information to make a decision.

clang/lib/Sema/SemaSYCL.cpp
24

I think this is also preventing traditional linking of translation units.

Could you elaborate more on this topic, please?
What do you mean by "traditional linking of translation units" and what exactly "is preventing" it?
Do you compare with the linking of regular C++ code (i.e. which do not split into host and device code)?
If so, SYCL is different from this model and more similar to CUDA/OpenMP models, which also skip "linking" of irrelevant part (e.g. host code is not linked by the device compiler).
Mariya added Justin (@jlebar) and Alexey (@ABataev), who work on single-source programming models to make them aware and provide feedback if any.

Design question: since you are not aware what functions are to be run on a device while parsing them, at what point do you plan to diagnose the invalid behavior from the standard C++ i.e. using function pointers in kernel code is likely to cause issues?

We are going to use DeviceDiagBuilder and related infrastructure implemented in Clang to diagnose device side code errors in OpenMP/CUDA modes.
More details are in the comments here:
https://clang.llvm.org/doxygen/classclang_1_1Sema_1_1DeviceDiagBuilder.html#details

Just a thought, if you parse host code first and provide the device outlining information to the device compilation phase would you then be able to reuse more parsing functionality from OpenCL?

Also do you need to outline the data structures too? For example classes used on device are not allowed to have virtual function.

Yes. This restriction is already implemented in our code base on GitHub.

Cool, is it implemented in SemaSYCL.cpp too?

clang/include/clang/Basic/Attr.td
1022

Sema part is mostly not relevant for SYCL mode because SYCL API do not allow the cases currently detected by clang (e.g. constant address space variable declaration in OpenCL kernel scope, naming OpenCL kernel main, etc).

Would you mind pointing me to your impl of those?

A couple of check that might be useful are:

void return type for kernel functions
kernel can't be static function

and some of the checks are harmful for proposed implementation (e.g. kernels can't be template functions).

@Anastasia, @keryell, @agozillon and @aaron.ballman need to agree if this sufficient to justify the re-use of OpenCL kernel attribute.
Let me know if you need any additional information to make a decision.

Ok, if from ~20 occurrences in the source code you will be able to reuse only just 2 it doesn't seem like it's worth to share __kernel attribute.

clang/lib/Sema/SemaSYCL.cpp
24

Yes indeed, I mean linking of modules in C/C++ even though it doesn't necessarily mean linking of object files. So you don't plan to support SYCL_EXTERNAL in clang?

In CUDA the functions executed on device are annotated manually using __device__ hence separate translation units can specify external device function... although I don't know if CUDA implementation in clang support this.

I guess OpenMP is allowed to fall back to run on host?

Fznamznon added inline comments.Tue, May 28, 4:35 AM
clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
2

Sorry for confusion. The C++ features used in SYCL are a subset of the C++11 standard features.
I will add -std=c++11 key to run line to avoid such confusion in future.

Fznamznon updated this revision to Diff 201641.Tue, May 28, 5:07 AM

Applied comments from @Anastasia

  • Added documentation for sycl_kernel function
  • Added comments to Sema.h
  • Added -std=c++11 to test run lines
Anastasia added inline comments.Thu, May 30, 10:53 AM
clang/include/clang/Basic/Attr.td
1022

Undocumented -> SYCLKernelDocs

clang/include/clang/Basic/AttrDocs.td
269

The example doesn't demonstrate the use of the attribute.

It explains how it is used by the toolchain only!

May be @aaron.ballman can help here as I am not sure what the format should be.

clang/lib/Parse/ParseAST.cpp
171

Do you also need to prevent generation of non-device functions somehow?

clang/lib/Sema/SemaSYCL.cpp
24

Ping!

I would suggest to document it a bit more including any current limitations/assumption that you can mark under FIXME i.e. does your code handle lambdas yet, what if lambdas are used in function parameters, etc...

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
5535–5542

May be this should go into a helper function as it seems to be now a bigger chunk of code that is repeated?

Although, I am not very familiar with this code. You can try to get someone to review who has contributed to this more recently.

clang/test/CodeGenSYCL/device-functions.cpp
25

I can't see where the SPIR calling convention is currently set for SYCL?

clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
4

I don't think this comment is necessary.

Fznamznon added inline comments.Fri, May 31, 5:52 AM
clang/include/clang/Basic/Attr.td
1022

Oh, Thank you for that!

clang/lib/Parse/ParseAST.cpp
171

I think It's already prevented by change to CodeGenModule.cpp in this patch. CodeGen just ignores declarations without SYCL device attribute now.

clang/lib/Sema/SemaSYCL.cpp
24

Oh, sorry, I missed this comment when I updated patch last time.
Could you please advise in which form I can document it?

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
5535–5542

I think this chunk of code seems big because of big repeated comment.

clang/test/CodeGenSYCL/device-functions.cpp
25

If I understand correct it's set automatically on AST level because we use SPIR-based triple for device code. Only in case of C++ methods clang doesn't set SPIR calling convention. We did a modification in our codebase to get SPIR calling convention for C++ methods too (available here )

Anastasia added inline comments.Mon, Jun 3, 3:00 AM
clang/test/CodeGenSYCL/device-functions.cpp
25

Ok and what happens if some other target is used - not SPIR?

Fznamznon added inline comments.Mon, Jun 3, 3:28 AM
clang/test/CodeGenSYCL/device-functions.cpp
25

There will be no SPIR calling convention for device functions.

Anastasia added inline comments.Mon, Jun 3, 7:54 AM
clang/test/CodeGenSYCL/device-functions.cpp
25

Just FYI at some point we generalized SPIR calling convention to be used for kernels irrespective from target by default (see TargetCodeGenInfo::getOpenCLKernelCallingConv). Not sure if it might make sense to do for SYCL device functions too. I am not saying it belongs to this patch though.

Fznamznon updated this revision to Diff 203785.Mon, Jun 10, 2:24 AM

Applied comments from @Anastasia

  • Added link to documentation for sycl_device attribute
  • Removed redundant comment from test

@Anastasia, do you have additional comments?

@aaron.ballman , please let me know if you have additional comments/suggestions. If not, could you please accept this revision?

Most of the comments are about minor nits like grammar and coding conventions, but I did have some questions regarding what kinds of functions the sycl_kernel attribute gets applied to. Also, I'd like to see some additional tests that demonstrate the sycl device attribute is being implicitly created on the proper declarations as expected (can probably do that using -ast-dump and checking to see if the right functions have the attribute attached).

clang/include/clang/Basic/AttrDocs.td
259

is SYCL "kernel function" -> is a SYCL "kernel function"

260

SYCL -> A SYCL

261

Kernel is a -> A kernel is a

263–264

This doesn't really demonstrate the need for the attribute -- the attribute is never shown in the code example. I'd prefer an example that shows when and how a user would write this attribute.

278

called SYLC -> called a SYLC

280

use sycl_kernel -> use the sycl_kernel

281

as SYCL -> as a SYCL
Compiler is supposed to -> The compiler will

284

In this code example compiler is supposed to add "foo" function -> In this code example, the compiler will add the "foo" function

clang/include/clang/Sema/Sema.h
11197

Function -> function

11198

In SYCL, when we generate device code, we don't know

11199

we emit sycl kernels, so we add device

11204

adds the function declaration

11205

Should be named addSyclDeviceFunc() per coding standards. Similar for the other new functions.

11209

Don't repeat the function name in the comments, please. Also, rather than returning a concrete SmallVector<>, I think it would be more natural to return a SmallVectorImpl so that callers don't have to contend with the explicit size. There should also be a const overload for this function.

11212

Constructs a SYCL kernel that is compatible with OpenCL from the SYCL "kernel

11215–11216

Marks all functions accessible from SYCL kernels with the SYCL device attribute

clang/lib/CodeGen/CodeGenModule.cpp
2408

with the SYCL device attribute

2410–2413

These if statements can be combined.

2529

Missing a full stop at the end of the comment.

clang/lib/Sema/SemaSYCL.cpp
15

This include doesn't seem to be necessary?

24

e does not use our usual naming conventions.

42

Spurious whitespace can be removed.

45

Elide braces.

49

Don't use auto as the type is not spelled out in the initialization.

53

Elide braces.

69

with the SYCL device attribute

71

elt -> Elt per naming conventions

72

auto * since the type is spelled out in the initialization.

74

Elide braces

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
5538

for the SYCL kernel attribute

5540

Elide braces

5552

for the SYCL kernel attribute

5554

Elide braces

clang/test/SemaSYCL/device-attributes-on-non-sycl.cpp
5

#ifndef ?

12

I'd prefer to spell this with __attribute__, same in the other test

clang/test/SemaSYCL/device-attributes.cpp
4

I'd like to see some more tests covering less obvious scenarios. Can I add this attribute to a lambda? What about a member function? How does it work with virtual functions? That sort of thing.

Fznamznon added inline comments.Tue, Jun 18, 8:01 AM
clang/test/SemaSYCL/device-attributes.cpp
4

Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.

But I'm working on new patch which will put some requirements on function which is marked with sycl_kernel attribute.
This new patch will add generation of OpenCL kernel from function marked with sycl_kernel attribute. The main idea of this approach is described in this document (in this document generated kernel is called "kernel wrapper").
And to be able to generate OpenCL kernel using function marked with sycl_kernel attribute we put some requirements on this function, for example it must be a template function. You can find these requirements and example of proper function which can be marked with sycl_kernel in this comment .

aaron.ballman added inline comments.Tue, Jun 18, 3:15 PM
clang/test/SemaSYCL/device-attributes.cpp
4

Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.

So there are no concerns about code like:

struct Base {
  __attribute__((sycl_kernel)) virtual void foo();
  virtual void bar();
};

struct Derived : Base {
  void foo() override;
  __attribute__((sycl_kernel)) void bar() override;
};

void f(Base *B, Derived *D) {
  // Will all of these "do the right thing"?
  B->foo();
  B->bar();

  D->foo();
  D->bar();
}
Fznamznon updated this revision to Diff 205663.Wed, Jun 19, 1:42 PM

Appled part of comments from @aaron.ballman:

  • Fixed grammar and code style in all places except sycl_kernel docs
  • Added a lit test which checks that sycl_device attribute implicitly added to proper declarations
Fznamznon added inline comments.Wed, Jun 19, 1:45 PM
clang/include/clang/Basic/AttrDocs.td
263–264

I see. I will update documentation in the next version.

bader added inline comments.Wed, Jun 19, 2:00 PM
clang/test/SemaSYCL/device-attributes.cpp
4

Actually there is no restrictions for adding this attribute to any function to outline device code so I just checked the simplest variant.
But I'm working on new patch which will put some requirements on function which is marked with sycl_kernel attribute.

@aaron.ballman, sorry for confusing. The usage scenarios should have been articulated more accurately.
We have only four uses of this attribute in our implementation:
https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/handler.hpp#L538 (lines 538-605).
All four uses are applied to member functions of cl::sycl::handler class and all of them have similar prototype (which is mentioned by Mariya in the previous comment:

namespace cl { namespace sycl {
class handler {
  template <typename KernelName, typename KernelType/*, ...*/>
  __attribute__((sycl_kernel)) void sycl_kernel_function(KernelType KernelFuncObj) {
    KernelFuncObj();
  }
};
}}

Here is the list of SYCL device compiler expectations with regard to the function marked with sycl_kernel attribute.

  • Function template with at least one parameter is expected. The compiler generates OpenCL kernel and uses first template parameter as unique name to the generated OpenCL kernel. Host application uses this unique name to invoke the OpenCL kernel generated for the sycl_kernel_function specialized by this name and KernelType (which might be a lambda type).
  • Function must have at least one parameter. First parameter expected to be a function object type (named or unnamed i.e. lambda). Compiler uses function object type field to generate OpenCL kernel parameters.

Aaron, I hope it makes more sense now.

We don't plan in any use cases other than in SYCL standard library implementation mentioned above.
If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?
What is the best way to do this? Add more negative tests cases and make sure that clang generate error diagnostic messages?

keryell added inline comments.Wed, Jun 19, 6:55 PM
clang/test/SemaSYCL/device-attributes.cpp
4

If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?

But since it is an attribute to be used by SYCL run-time writers, I am not sure there is a lot of value in over-engineering the restrictions of its use. It diverts brain power from the real implementation & review and might even prevent innovation due to some creative use cases.

Fznamznon updated this revision to Diff 205813.Thu, Jun 20, 7:54 AM

Updated sycl_kernel attribute documentation.

Fznamznon updated this revision to Diff 205831.Thu, Jun 20, 8:58 AM

Fixed a couple coding style issues, renamed markDevice function with markSYCLDevice.

aaron.ballman added inline comments.Mon, Jun 24, 1:28 PM
clang/test/SemaSYCL/device-attributes.cpp
4

If I understand you concerns correctly, you want to be sure that clang prohibits other uses of this attribute, which are not intended. Right?

Effectively, yes. I'd like to ensure that situations where the attribute does not do what the user expects are diagnosed. A good rule of thumb that I use is to diagnose (as a warning) situations where the attribute will be silently ignored, and diagnose (as an error) situations where applying the attribute would cause really bad results (like miscompiles, security concerns, etc).

What is the best way to do this? Add more negative tests cases and make sure that clang generate error diagnostic messages?

That's a good approach, yes. Though for the situations you describe, I'd probably just warn rather than err because it seems like it's harmless to ignore the attribute so long as the user knows it's being ignored.

4

But since it is an attribute to be used by SYCL run-time writers, I am not sure there is a lot of value in over-engineering the restrictions of its use. It diverts brain power from the real implementation & review and might even prevent innovation due to some creative use cases.

I disagree. Part of the real implementation is ensuring the attribute is not accidentally misused. It's frustrating for users to have an attribute silently ignored because it's easy to mistake that situation for the attribute behaving as expected.