Minimal case study of using SYCL from cpp2 #853
MiroPalmu
started this conversation in
Show and tell
Replies: 2 comments
-
For reference, these are the issues and comments related to the pain points:
|
Beta Was this translation helpful? Give feedback.
0 replies
-
Interesting to see some SYCL code written with cpp2! |
Beta Was this translation helpful? Give feedback.
0 replies
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
-
Motivation
Cpp1 is the way to program GPUs in scientific computing with many different programming models.
One of the interesting models is SYCL which aims to accomplish this with standard cpp1 but then interpret some functions as device code and use special compilers to compile it to correct target.
Intel describes SYCL as:
Newest spec can be found here.
Security goals of cpp2 might not be as desired in scientific computing as in other domains,
but simplification of language make it interesting for computational scientist,
who is more interested in just making algorithms work than in the technical details of the code.
(most of all I'm just interested to try cpp2)
Goals of case study
Cpp1 SYCL source
This is modified example from SYCL spec section 3.2
Cpp2 SYCL source
cppfront commit
0b333f3
is able to transpile following code to form that is compile-able with AdaptiveCpp.Pain points
Keep in mind that these points might be due to authors lack of knowledge.
Order of points is arbitrary and does not indicate severity.
Commit
4bd0c04
force mutability to anonymous functionsIn SYCL kernel functions are callable objects (most often lambdas) which get executed on device (gpu, etc.).
SYCL spec states that lambda kernel object can not be
mutable
.Stated commit adds
mutable
to anonymous functions, so this commit can not be used to write kernel functions as lambda.This is why previous commit was used.
Not being able to write default capture types
SYCL defines command group scope as the callable object that gets passed to
queue.submit
,in which
sycl::buffer
s are needed as reference and in kernel scopesycl::accessor
sare needed as values.
This leads to that being able to specify default capture type is very convenient,
which can be seen from SYCL spec which uses
[&]
and[=]
in every* example.In cpp2 we need to use
var&$*
andvar$
which is not as ergonomic as default captures.Matching role based parameters with cpp1 parameters
Seen in cpp1 code, lambda to
queue.submit(...)
has to take handler as reference.In all SYCL spec examples
all parameters to kernel functions are taken as values.
(This might not be necessary due to
in
role working in cpp2 code. I haven't found definite answer.)It might not be obvious which roles to choose for the parameters in cpp2
to match SYCL spec examples. This might be non-problem with enough experience with cpp2.
In kernel function case
copy
parameter role might be desired to be able to match cpp1 examples which take parameter as value.Having to write
sycl::accessor(..., sycl::property_list(sycl::no_init))
and notsycl::accessor(... , { sycl::no_init })
Structured binding are not supported in
for
loop.Beta Was this translation helpful? Give feedback.
All reactions