- User Since
- Mar 5 2015, 8:05 AM (219 w, 3 d)
Fri, May 10
The overall structure looks very clear now! I just have a couple of nitpicks! Thank you!
Thu, May 9
Wed, May 8
- Improved comments
- Switched to SPIR kernel
Tue, May 7
Fri, May 3
Can this be closed now due to commit r358674?
Removed redundant C linkage checks!
Thu, May 2
Extended FunctionDecl exten C logic to take kernel function into account.
Added extra overload for isAddressSpaceSupersetOf.
Wed, May 1
I would suggest an IR test, especially since what you want to achieve is addrspacecast in IR output. We use AST dump only if there is no other good way to test something. In fact we already have a test with similar logic that tests various OpenCL address spaces: test/CodeGenOpenCLCXX/addrspace-of-this.cl. I think it should be reasonable enough to extend it.
Tue, Apr 30
Seems reasonable for OpenCL kernels. You might want to add an AST dump test to check that the visibility is being set correctly in case it's being printed in AST.
Sat, Apr 27
I think potentially reusing OpenCL features is desirable since the device code of SYCL is largely OpenCL. However I don't think we are clear enough about the overall device compilation flow of SYCL and I can easily suggest a number of different approaches including those that don't modify compiler at all. :) I am afraid until we have the big picture clear it will be hard to make any sensible decisions.
Fri, Apr 26
Wed, Apr 24
Apr 19 2019
SYCL is similar to OpenMP 5 for C++, where you use only C++ classes instead of #pragma. So it is quite C++-friendlier than OpenMP.
Apr 18 2019
We could still go for something like the following but accept that some a random failure might happen not necessarily on a commit that introduces it?
I am not sure that this is needed. Non-deterministic tests are really annoying.
Yes. I think this might be a wider topic to address if we happen to have more issues of this kind.
Btw, how is this tested? Do we have an existing test that exercises this already?
Apr 17 2019
I think what we are missing currently is a thorough analysis/comparison between SYCL device mode and OpenCL kernel language mode to understand what's the best implementation strategy. That would apply to many other features: kernel function restrictions, address spaces, vectors, special types, etc.
That would make definitely sense when we target OpenCL.
I still see no point in polluting our code base with extra code that just does the same thing. It will save us a lot of time to just work cooperatively on the same problem and even improve readability of the code. But of course this can only be done if there is no need to diverge the implementation significantly.
Yes. Probably that even when the target is not OpenCL, the general principles remain similar. Probably the same for CUDA & OpenMP 5 too...
Apr 16 2019
FYI, I switched to std::map in this review:
Not related to this patch but it might be good to start thinking about testing this functionality properly. In the past we haven't tested the header because it would take a lot of testing time. So I would suggest we keep a light minimal basic testing in regular clang tests as is, but then we either add special build target to tests extra header functionality or add special cmake option (i.e. LLVM_INCLUDE_OPENCL_BIFS_TESTS) that would enable such testing with regular check-clang. I assume the latter can even be used to avoid generating the extra tables from TableGen minimizing the impact on the clang binary size for the situations OpenCL isn't required in the installation.
Apr 15 2019
Our current approach is to add an attribute, 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.
Here is a link to the code in the SYCL runtime using sycl_kernel attribute: https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/handler.hpp#L267
I'm quite sure something similar should happen for other "single source" programming models like OpenMP/CUDA, except these attributes are exposed to the user and there is a specific requirement on attributes/pragma/keyword names.
Apr 9 2019
I was just wondering since SYCL is intended to be a single source C++ for OpenCL and this attribute is for internal use is it possible to just reuse existing OpenCL kernel attribute?
Apr 8 2019
LGTM! Great! Thanks!
Apr 4 2019
Apr 3 2019
Improved comment about initializers in __local addr space.
Apr 2 2019
- Use AggValueSlot in the constructor call generation to store/retrieve address space of 'this'.
- Fixed detecting the address space conversion while performing qualification conversion.
Apr 1 2019
Mar 29 2019
Mar 28 2019
Mar 27 2019
Extended test case and explained the intention of the test.
Mar 26 2019
Committed in r356987.
Mar 25 2019
Fixed typo in the comment.
Changed diagnostic wording to indicate that this feature is a Clang extension.
Mar 21 2019
Updated test to use root namespace (commented by Ronan).