llvm

[SYCL][Doc] Add spec for "sycl" to kernel compiler

#11985

Merged

Why? That means the string can't include a header which in turn includes sycl.hpp.

That's a good point. We could relax this to say that the source string does not need to #include <sycl/sycl.hpp>, while still allowing it.

I think this would be better. But even saying that it isn't needed might have a disadvantage in the future. In the future we might want to split sycl.hpp into different parts to improve compilation speed. And then online compile wouldn't benefit from it because we would need to include all of the device side API to be backward compatible. I'm not sure the convenience of not having to include sycl.hpp is worth this potential downside in the future. It might be best only to say that sycl.hpp is automatically included in include_files so that the user doesn't have to do that part.

I think the user's source string should explicitly include #include <sycl/sycl.hpp>. I agree with @rolandschulz that we may want to change the headers in future, and it would be nice to have this flexibility.

I also think that it's weird to make an exception for the SYCL headers, because users may need to include other C++ headers (e.g., #include <cmath>) if they're making use of any std:: functionality inside of their kernels.

I think we should replace this with a note to the effect that the compiler behaves as-if the location of the SYCL headers and standard C++ headers are included in include_files. Do we also need to add something saying that these headers are guaranteed to be the same as those used when the host program was compiled? (Can we actually guarantee that?)

OK, I agree that this makes sense. Changed in d49df13.

Is this available somewhere?

No, it's not implemented yet. I expect that we will implement that extension at the same time as the new additions to "sycl_ext_oneapi_kernel_compiler" that are described in this PR.

No changes required here, but an observation: "an std::string" reads better for people who say "an S-T-D string", while "a std::string" reads better for people who say "a stood string". We should give some thought to what style we (and the SYCL specification) want to use. Saying "an object of type std::string" would sidestep it.

It appears that the consensus is to use "a std ..." (not "an std ..."), so I guess people have decided that the pronunciation is "stood" and not "s-t-d".

The C++ spec has 53 occurrences of "a std" and no occurrences of "an std".
The cppreference site has over 1300 occurrences of "a std" and only 71 occurrences of "an std".

I'll start using "a".

These occurrences fixed in #12374. Another occurrence fixed in this PR at 019977c.

In this design, foo doesn't accept sycl::id<> or sycl::item<> as a parameter, but instead is expected to use the free functions ( for example: sycl::ext::oneapi::experimental::this_id<1>(); ) to get the "index".

So, AFAICT at my initial stage of exploring this, if that's true, then there is no need for any argument for SYCL_EXT_ONEAPI_FUNCTION_PROPERTY. If it were just SYCL_EXT_ONEAPI_KERNEL_FUNCTION that might be enough to designate the following function as something that needs to be wrapped in a kernel call.

We actually had a long debate about this. We'd prefer to use the common syntax of compile-time properties, rather than inventing new macros. The macro SYCL_EXT_ONEAPI_FUNCTION_PROPERTY is part of the common syntax for compile-time properties.

FWIW, see 1b6f57e for the way it looked before we decided to use SYCL_EXT_ONEAPI_FUNCTION_PROPERTY.

I was thinking of something like
sycl::kernel iota = kb_exe.ext_oneapi_get_kernel( syclex::get_spirv_kernel_name("iota") );

Where get_spirv_kernel_name would return the mangled kernel name that matches. Of course, this doesn't handle different overloads of iota. A simple solution like that might work, with a fallback to something like iota_name"(void(*)(int*, int*))iota"}; when using overloads. Otherwise, it seems like we are needlessly complicating it for everyone who doesn't use overloads.

I was having trouble getting this extern "C" approach to work, mostly because we employ a lot of templates in the headers right now and I was trying to minimize the disruption. A demangling helper like syclex::get_kernel_spirv_name("iota") is a straightforward way of avoiding the problem.

I don't understand your comment about extern "C". Why does it matter that we use a lot of templating in our headers? How does this affect the user's ability to declare a function extern "C"?

Regarding, get_kernel_spirv_name("iota"), see the proposed function ext_oneapi_get_raw_kernel_name which provides similar functionality.

I have some reservations about this include_files struct.

We know for the first implementation that we are likely going to have to write out a temporary .cpp file and invoke clang++ -fsycl on it. So now we have to also write out the include files? What about #include "path/to/file.h" ? Will we then be parsing this and instantiating directories?

I presume we need to disambiguate #include "local/path.h" from #include <system/path.h> ?

Users can't link in libraries ( even if has a header ), and there is only the one TU. Supporting #include suggests the opposite. Maybe a note proscribing it is needed?

I do see the value in that if a user has a simple function that they want to both use in the host application and call from their kernel, this allows them to do so. And, if clang were to ever have an invocation API (like ocloc) then I guess this aligns with that.

One of the reasons we added the include_files property is for parity with CUDA, which has a similar feature in NVRTC. By adding the feature to SYCL, we make it easier to migrate code automatically using the SYCLomatic tool.

I was thinking that our implementation could create a temporary "include" directory and then write out the include_files header files into that directory. When we invoke clang, we can add a -I switch pointing at the temporary include directory.

Note that you don't need to parse any #include strings from the source code to get the file name paths. The include_files properties contains a list of (Name, Content) pairs. Each Name is the relative pathname of an include file, so you can just use this Name string verbatim.

Is this a Precondition, i.e., does using an invalid source language and format result in undefined behavior? Maybe we should move this under Throws instead?

This function succeeds even if some devices in `ctxt` do not support the source
language `lang`.
However, the `build` function fails unless _all_ of its devices support `lang`.
Therefore, applications should take care to omit devices that do not support
`lang` when calling `build`.

Should we maybe introduce an aspect to check a device supports a given source language?

Is this a Precondition, i.e., does using an invalid source language and format result in undefined behavior? Maybe we should move this under Throws instead?

Good idea. I added a Throws section in ce3bbbe.

Should we maybe introduce an aspect to check a device supports a given source language?

We don't need an aspect because device::ext_oneapi_can_compile provides this query.

I can see instances of "this property is only valid for X source language" in the future. This makes me think whether we want the source language as an argument or as a template argument... Having it as a template argument will enable constructs like:

• Further restricting properties being passed at compile-time
• Restricting input format (text or binary) at compile time

I might be missing something, but, what does having the source language as a runtime variable bring to the table? Can we just have users use some kind of dispatching to call create_kernel_bundle_from_source<lang> instead and improve compile-time checks?

I agree that making it a template argument would allow compile-time error checking here. However, I think it is valuable to allow applications to pass lang as a runtime value, and I don't want to require them to code a dispatch table in this case. Therefore, I'm more inclined to diagnose a runtime error if a source language does not support one of the properties.

Okey, sounds good to me. Now, should we add a Throws clause to create_kernel_bundle_from_source in case lang does not support any of the input properties or just state unsupported properties will be omitted?

Yup, I just ran out of time yesterday. Done now in 1c3f01b.

Does this require sycl-post-link?

I don't think I understand. Why would this involve sycl-post-link?

My reading from this is that the runtime will query a property set by sycl-post-link to check what SPIR-V fallback libraries are needed. Are you planning on reusing this functionality, reimplement it or is it not needed?

I expect the first implementation to be very simplistic. The SYCL runtime will simply write the input string to a temporary file and spawn clang++ -fsycl as a separate process. This will end up spawning sycl-post-link as a subprocess along with all of the other tools that the compiler driver invokes.

I guess the answer to your original question is "yes", we will use sycl-post-link.

Perfect. Thanks!

Should we add a Throws to handle the case in which input lists have a different size?

Even better IMO. Change this property so that it receives a vector of include structs with name and contents instead of two separate name and content.

Should we Constraint a list may only have a single instance of these include property? (or document here how they are to be combined)

Should we Throw if a name is duplicate?

Good suggestions. See 1d5cf11.

I think we are still missing what to do if more than one include property is passed. I am not familiar enough with the compile-time property list extension to know whether that's even allowed, tho. Rest of changes LGTM.

I think we are still missing what to do if more than one include property is passed.

This is a fundamental restriction on the PropertyListT. The sycl_ext_oneapi_properties extension does not allow a property list to contain more than one instance of any single property (even if the property values are different).

I assume a missing include is a compilation error and will be reported by build(?)

Yes. I don't think we need to call this out specifically. Attempting to #include a non-existent file is just one of many possible compilation errors.

Perfect. Totally agree 👍

Can we add a Precondition saying the kernel name is a valid qualified identifier? "", "int", "777_invalid_id" might be problematic, so I'd at least say that's UB.

Yes. See bb1e69f.

Great. Thanks!

Should we have a Precondition saying the kernel is a valid C++ qualified identifier?

I think this is handled by the Throws section below. Or maybe I don't understand your concern?

Yes, it is indeed covered. I was just thinking of a corner implementation case in which the queried name is not a valid identifier and that may cause issues, but, again, that should be handled by the implementation. My bad. LGTM.

"Pro tip": Using C++ in asciidoc might be problematic, as ++ are escape characters (everything between ++ and ++ won't be formatted). Add the following to the beginning of the file and use {cpp} instead of C++ to avoid surprises in the future:

:cpp: pass:[C++]

I ran into some instances of C++ rendering as C and it can be annoying...

Yes. Done in 002d47b.

Should we format as C++ all the code listings in the proposal?

Also yes. Done in 5f85662.

Question: What are the chances extensions to this extension define their own properties? If there is a chance, this constraint may not age well.

We don't make extensions to extensions. Instead, we would make a revision to this extension, which adds a new property. In that case, the wording would still be correct.

Excellent. Thanks!

Question: Could a source language allow both kinds of input in the future?

I don't think that is likely. If it does happen, we can revise this extension in a way that remains backward compatible.

Excellent. Thanks, Greg!

Not added in this PR, but I'd still drop "string" as the source code might be in different formats.

Makes sense. See d40c117.

Same: Could extensions add new properties or are all properties gonna be listed here?

Same answer as above.

NIT: Should we Precondition to is not null and outlives the build call?

Yes, this could be tightened. See cd36574.

Perfect. Thanks!

Just one question: why would we want to allow a nullptr being passed there? Wouldn't that make save_log a NOOP? Maybe passing std::string & makes more sense.

I was thinking that it might be useful to allow code like this:

auto props = properties{
  //...
  save_log{(enable_log) ? &log : nullptr},
  //...
};

Perfect. Oversaw that. Resolving this now 👍

I am wondering what would be the behavior in case the user query were ambiguous, e.g.:

void foo(int);
void foo(float);

Now, querying for "foo" would be ambiguous, as both void foo(int) and void foo(float) apply. Do we need to document this case?

We use normal C++ rules to disambiguate the two overloads. The formal requirement is that the string passed to registered_kernel_names is an expression for a function pointer. The way to do this in C++ when there are different overloads of the function is to use a cast. This is illustrated in the examples section "Disambiguating overloaded kernel functions" below.

This is really cool. However, what if the user does not instantiate the template? What if they just query the name and try to use it? Would the following code be an error?

sycl::kernel_bundle<sycl::bundle_state::executable> kb = syclex::build(kb_src,
  syclex::properties{syclex::registered_kernel_names{{"bartmpl"}});

sycl::kernel k = kb.ext_oneapi_get_kernel("bartmpl");

The section "Using the registered_kernel_names property" says:

Each string in the property must be the C++ expression for a pointer to a kernel function.

Your example violates this requirement, so I think the behavior is undefined. The string "bartmpl" is the name of a function template; it is not a valid expression for a pointer to a function.

I suspect our implementation will fail in the build call because it will try to compile "bartmpl" in a context where a function pointer is expected.

Okey, perfect. Thanks, Greg!

Double when here, as it is already present at the end of the previous line.

Thanks. Fixed in e69a5dd.