Add SYCL as a target #44
Conversation
| @@ -0,0 +1,240 @@ | |||
| __copyright__ = "Copyright (C) 2011-20 Andreas Kloeckner" | |||
There was a problem hiding this comment.
Please fix the copyright statement.
| ) | ||
|
|
||
|
|
||
| def dtype_to_cltype(dtype): |
There was a problem hiding this comment.
This looks like it's duplicated from CL. Please just import from there.
| yield "{}.submit({});".format( | ||
| self.queue, ", ".join(str(ad) for ad in self.args) | ||
| ) |
There was a problem hiding this comment.
Is this special syntax? This looks like a normal expression, and expressions are out of scope for cgen. (See, e.g. pymbolic.)
There was a problem hiding this comment.
@inducer We have a small thing to clarify. You are asking to expose an API that would generate this expression step by step right? An example that would be generated by this would be,
queue_.submit([&](sycl::handler &handler) {
handler.parallel_for(range_, [=](sycl::nd_item<3> item) {
a[item.get_group(0)] = item.get_group(0);
});
})
Let's assume we have a class cgen.Variable which constructs the queue_ part and the cgen.Call would call .submit on queue_ where cgen.Variable and cgen.Call are exposed as an API which helps generate the above expression.
Is this similar to the change that you expect?
There was a problem hiding this comment.
All of the things here are expressions (the call, the lambda, all of it). cgen doesn't currently model expressions, relying on pymbolic for that, though pymbolic better models Python than C++ ASTs. cgen is unlikely to become a full C++ AST, or I at least have no intention of taking it in that direction. Maybe you'll want to fork?
There was a problem hiding this comment.
Is it an acceptable change if we handle the expressions on the loopy side such that SYCL target will offer similar functionality as existing backends of cgen?
| yield "{}.parallel_for({});".format( | ||
| self.handler, ", ".join(str(ad) for ad in self.args) | ||
| ) |
There was a problem hiding this comment.
Is this special syntax? This looks like a normal expression, and expressions are out of scope for cgen. (See, e.g. pymbolic.)
| while len(dim) < 3: | ||
| dim = dim + (1,) | ||
| self.dim = dim | ||
| self.decl = 'extern "C" [[sycl::reqd_work_group_size({})]]'.format( |
There was a problem hiding this comment.
Is the extern "C" necessarily part of this? It would seem that this prohibits overloading.
| mapper_method = "map_sycl_global" | ||
|
|
||
|
|
||
| # TODO sycl Image |
| else: | ||
| return ["auto&"], Assign( | ||
| sub_decl, | ||
| f"*sycl::ext::oneapi::group_local_memory<{sub_tp[0]}>" |
There was a problem hiding this comment.
Is local memory not exposed in sycl? (without using an extension)
No description provided.