Page MenuHomePhabricator

Fznamznon (Mariya Podchishchaeva)
User

Projects

User does not belong to any projects.

User Details

User Since
Dec 18 2017, 12:01 AM (144 w, 5 d)

Recent Activity

Jun 17 2020

Fznamznon updated the diff for D81641: [SYCL][OpenMP] Implement thread-local storage restriction.

Rebased.

Jun 17 2020, 4:16 AM · Restricted Project

Jun 16 2020

Fznamznon added a comment to D81641: [SYCL][OpenMP] Implement thread-local storage restriction.

Seems that test/OpenMP/nvptx_target_codegen.cpp is completely not formatted. If I apply suggestion from pre-merge checks, this will look like a big unrelated to this patch change and it will contradict with the whole file style.

Jun 16 2020, 11:57 PM · Restricted Project
Fznamznon added a comment to D81641: [SYCL][OpenMP] Implement thread-local storage restriction.

OpenMP has the same restriction (no surprise I guess). Thanks for the ping!

I think we do not emit diagnosis right now: https://godbolt.org/z/srDkXZ
I think we also should diagnose this the same way, though it might be beyond the scope of this patch: https://godbolt.org/z/rRZFi4

Jun 16 2020, 10:26 AM · Restricted Project
Fznamznon retitled D81641: [SYCL][OpenMP] Implement thread-local storage restriction from [SYCL] Implement thread-local storage restriction to [SYCL][OpenMP] Implement thread-local storage restriction.
Jun 16 2020, 10:26 AM · Restricted Project
Fznamznon updated the diff for D81641: [SYCL][OpenMP] Implement thread-local storage restriction.

Generalized diagnostic between SYCL and OpenMP.

Jun 16 2020, 10:26 AM · Restricted Project

Jun 11 2020

Fznamznon added inline comments to D81641: [SYCL][OpenMP] Implement thread-local storage restriction.
Jun 11 2020, 7:41 AM · Restricted Project
Fznamznon updated the diff for D81641: [SYCL][OpenMP] Implement thread-local storage restriction.

Fixed code style.

Jun 11 2020, 7:41 AM · Restricted Project
Fznamznon updated subscribers of D81641: [SYCL][OpenMP] Implement thread-local storage restriction.

@jdoerfert , @ABataev , if OpenMP needs same diagnostic as well, I can generalize it between SYCL and OpenMP.

Jun 11 2020, 3:36 AM · Restricted Project
Fznamznon created D81641: [SYCL][OpenMP] Implement thread-local storage restriction.
Jun 11 2020, 3:36 AM · Restricted Project
Fznamznon added reviewers for D81641: [SYCL][OpenMP] Implement thread-local storage restriction: erichkeane, bader.
Jun 11 2020, 3:36 AM · Restricted Project

May 29 2020

Fznamznon added a comment to D80829: [OpenMP][SYCL] Do not crash on attempt to diagnose unsupported type use.

I don't have rights to merge. Could someone submit it, please?

May 29 2020, 12:33 PM · Restricted Project
Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Fix is here D80829 .

May 29 2020, 12:00 PM · Restricted Project
Fznamznon created D80829: [OpenMP][SYCL] Do not crash on attempt to diagnose unsupported type use.
May 29 2020, 12:00 PM · Restricted Project
Fznamznon added reviewers for D80829: [OpenMP][SYCL] Do not crash on attempt to diagnose unsupported type use: ABataev, bader.
May 29 2020, 12:00 PM · Restricted Project
Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Seems to me, this patch crashes llvm-project/openmp/libomptarget/test/mapping/declare_mapper_api.cpp.

May 29 2020, 11:26 AM · Restricted Project

May 28 2020

Fznamznon updated the diff for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Included test cases from Johannes.

May 28 2020, 10:55 AM · Restricted Project
Fznamznon added a comment to D77918: [OpenMP] Avoid crash in preparation for diagnose of unsupported type.

I think this will be fixed by D74387, we should include the tests somewhere though.

May 28 2020, 9:55 AM · Restricted Project
Fznamznon added inline comments to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.
May 28 2020, 8:06 AM · Restricted Project
Fznamznon updated the diff for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Applied comments from Johannes.
Fixed failing tests.

May 28 2020, 8:06 AM · Restricted Project

May 26 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

The tests are failing because calling function with unsupported type in arguments/return value is diagnosed as well, i.e. :

double math(float f, double d, long double ld) { ... } // `ld` is not used inside the `math` function
#pragma omp target map(r)
  { r += math(f, d, ld); } // error: 'math' requires 128 bit size 'long double' type support, but device 'nvptx64-nvidia-cuda' does not support it

Should we diagnose calls to such functions even if those arguments/return value aren't used?

May 26 2020, 3:44 AM · Restricted Project

May 25 2020

Fznamznon retitled D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage from [SYCL] Defer __float128 type usage diagnostics to [OpenMP][SYCL] Improve diagnosing of unsupported types usage.
May 25 2020, 12:51 PM · Restricted Project
Fznamznon updated the diff for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Re-implemented diagnostic itself, now only usages of declarations
with unsupported types are diagnosed.
Generalized approach between OpenMP and SYCL.

May 25 2020, 12:51 PM · Restricted Project

May 8 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Okay, seems like OpenMP needs unsupported types diagnosing as well. I'm trying to adapt this patch for OpenMP, but it doesn't work out of the box because it diagnoses memcpy like operations, so with the current patch the code like this will cause diagnostics:

 struct T {
   char a;
   __float128 f;
   char c;
   T() : a(12), f(15) {}
}

#pragma omp declare target
T a = T();
#pragma omp end declare target

It happens because member initialization in the constructor is still usage of f field which is marked unavailable because of type. I'm not sure that it is possible to understand how the unavailable declaration is used in the place where diagnostic about usage of unavailable declaration is actually emitted, so I will probably need some other place/solution for it.

@jdoerfert , could you please help to understand how the diagnostic should work for OpenMP cases? Or you probably have some spec/requirements for it?
Understanding what exactly is needed will help with the implementation, I guess.

I missed this update, sorry.

I don't think we have a spec wording for this, it is up to the implementations.

In the example, a diagnostic is actually fine (IMHO). You cannot assign 15 to the __float128 on the device. It doesn't work. The following code however should go through without diagnostic:

struct T {
   char a;
   __float128 f;
   char c;
   T() : a(12), c(15) {}
}

and it should translate to

struct T {
   char a;
   alignas(host_float128_alignment) char[16] __unavailable_f;
   char c;
   T() : a(12), c(15) {}
}

Do you have other questions or examples we should discuss?

May 8 2020, 4:47 AM · Restricted Project

May 6 2020

Herald added a reviewer for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage: aaron.ballman.

Ping.

May 6 2020, 10:13 AM · Restricted Project

Apr 20 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Okay, seems like OpenMP needs unsupported types diagnosing as well. I'm trying to adapt this patch for OpenMP, but it doesn't work out of the box because it diagnoses memcpy like operations, so with the current patch the code like this will cause diagnostics:

 struct T {
   char a;
   __float128 f;
   char c;
   T() : a(12), f(15) {}
}
Apr 20 2020, 10:16 AM · Restricted Project

Apr 14 2020

Fznamznon added inline comments to D77918: [OpenMP] Avoid crash in preparation for diagnose of unsupported type.
Apr 14 2020, 10:11 AM · Restricted Project

Apr 13 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

As I mentioned before. As long as the type is not "used" you can treat it as a sequence of bytes just as well. So we can lower __float128 to char [16] with the right alignment. SPIRV will never see unsupported types and the code works because we never access it as float128 anyway. WDYT?

Apr 13 2020, 11:17 AM · Restricted Project

Apr 9 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

This is needed for OpenMP as well. Does it make sense to include it in this patch or in another one?

I thought OpenMP already has diagnostics for unsupported types (at least looking into this commit https://github.com/llvm/llvm-project/commit/123ad1969171d0b22d0c5d0ec23468586c4d8fa7). Am I wrong?
The diagnostic which I'm implementing here is stricter than existing OpenMP diagnostic, the main goal is do not emit unsupported type at all. Does OpenMP need such restriction as well?

OpenMP handling needs to be reverted/redone:

  1. If no aux triple is available it just crashes.
  2. If the unavailable type is not used in one of the pattern matched expressions it crashes (usually during instruction selection but not always). Try a call with long double arguments for example.

I'm not sure this patch fits the bill but what I was thinking we need is roughly:
If you have a expression with operands or function definition with return/argument types which are not supported on the target, mark the definition as unavailable with the type note you have.
We should especially allow members to have unavailable types if the member is not accessed. Memcpy like operations (=mapping) are OK though. I think this should be the same for OpenMP and Sycl (and HIP, and ...).

Why we should allow members to have unavailable types if the member is not accessed? I don't think that we always can do it, especially for SYCL. Even if the member is not accessed directly, the whole struct with unavailable type inside will get into resulting LLVM IR module anyway, this can be a problem, I guess.

On the host you know how large the type is so you can replace it in the device module with a placeholder of the appropriate size. You want to do this (in OpenMP for sure) because things you map might have constitutes you don't want to access on the device but you can also not (easily) split out of your mapped type.

Okay, I see. Am I right that OpenMP already has such thing implemented, but only for functions return types? I suppose, for SYCL, we might need to replace unsupported type in device module everywhere...
BTW, one more question, we also have a diagnostic which is emitted on attempt to declare a variable with unsupported type inside the device code for this __float128 type and other ones (https://github.com/intel/llvm/pull/1465/files). Does OpenMP (and probably HIP, CUDA etc) need such diagnostic as well?

I'm not sure we want this and I'm not sure why you would. To me, it seems user hostile to disallow unsupported types categorically. We also know from our codes that people have unsupported types in structs that they would rather not refactor. Given that there is not really a need for this anyway, why should we make them? Arguably you cannot "use" unsupported types but an error like that makes sense to people. So as long as you don't use the unsupported type as an operand in an expression you should be fine.

We have some detection for this in clang for OpenMP but it is not sufficient. We also should generalize this (IMHO) and stop duplicating logic between HIP/CUDA/OpenMP/SYCL/... That said, we cannot error out because the types are present but only if they are used. I would hope you would reconsider and do the same. Arguably, mapping/declaring a unsupported type explicitly could be diagnosed (with a warning) but as part of a struct I would advice against.

Maybe I just don't understand. Could you elaborate why you think sycl has to forbid them categorically?

Apr 9 2020, 6:28 AM · Restricted Project

Apr 8 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

This is needed for OpenMP as well. Does it make sense to include it in this patch or in another one?

I thought OpenMP already has diagnostics for unsupported types (at least looking into this commit https://github.com/llvm/llvm-project/commit/123ad1969171d0b22d0c5d0ec23468586c4d8fa7). Am I wrong?
The diagnostic which I'm implementing here is stricter than existing OpenMP diagnostic, the main goal is do not emit unsupported type at all. Does OpenMP need such restriction as well?

OpenMP handling needs to be reverted/redone:

  1. If no aux triple is available it just crashes.
  2. If the unavailable type is not used in one of the pattern matched expressions it crashes (usually during instruction selection but not always). Try a call with long double arguments for example.

I'm not sure this patch fits the bill but what I was thinking we need is roughly:
If you have a expression with operands or function definition with return/argument types which are not supported on the target, mark the definition as unavailable with the type note you have.
We should especially allow members to have unavailable types if the member is not accessed. Memcpy like operations (=mapping) are OK though. I think this should be the same for OpenMP and Sycl (and HIP, and ...).

Why we should allow members to have unavailable types if the member is not accessed? I don't think that we always can do it, especially for SYCL. Even if the member is not accessed directly, the whole struct with unavailable type inside will get into resulting LLVM IR module anyway, this can be a problem, I guess.

On the host you know how large the type is so you can replace it in the device module with a placeholder of the appropriate size. You want to do this (in OpenMP for sure) because things you map might have constitutes you don't want to access on the device but you can also not (easily) split out of your mapped type.

Apr 8 2020, 10:18 AM · Restricted Project

Apr 7 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

This is needed for OpenMP as well. Does it make sense to include it in this patch or in another one?

I thought OpenMP already has diagnostics for unsupported types (at least looking into this commit https://github.com/llvm/llvm-project/commit/123ad1969171d0b22d0c5d0ec23468586c4d8fa7). Am I wrong?
The diagnostic which I'm implementing here is stricter than existing OpenMP diagnostic, the main goal is do not emit unsupported type at all. Does OpenMP need such restriction as well?

OpenMP handling needs to be reverted/redone:

  1. If no aux triple is available it just crashes.
  2. If the unavailable type is not used in one of the pattern matched expressions it crashes (usually during instruction selection but not always). Try a call with long double arguments for example.

I'm not sure this patch fits the bill but what I was thinking we need is roughly:
If you have a expression with operands or function definition with return/argument types which are not supported on the target, mark the definition as unavailable with the type note you have.
We should especially allow members to have unavailable types if the member is not accessed. Memcpy like operations (=mapping) are OK though. I think this should be the same for OpenMP and Sycl (and HIP, and ...).

Apr 7 2020, 10:21 AM · Restricted Project

Apr 6 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

This is needed for OpenMP as well. Does it make sense to include it in this patch or in another one?

Apr 6 2020, 10:18 AM · Restricted Project

Mar 31 2020

Fznamznon added inline comments to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.
Mar 31 2020, 9:22 AM · Restricted Project
Fznamznon updated the diff for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Apply comments, rebase.

Mar 31 2020, 9:22 AM · Restricted Project

Mar 30 2020

Fznamznon updated the diff for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Rebased to fresh version. Applied fixes after https://reviews.llvm.org/D70172

Mar 30 2020, 9:43 AM · Restricted Project
Fznamznon added inline comments to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.
Mar 30 2020, 9:10 AM · Restricted Project

Mar 27 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Ping.

Mar 27 2020, 6:29 AM · Restricted Project

Mar 23 2020

Fznamznon updated the diff for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Fix the test by adding the target with __float128 support and make sure that
no diagnostic are emitted.

Mar 23 2020, 4:20 AM · Restricted Project

Mar 20 2020

Fznamznon added a reviewer for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage: bader.
Mar 20 2020, 1:02 PM · Restricted Project
Fznamznon updated the summary of D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.
Mar 20 2020, 1:01 PM · Restricted Project
Fznamznon updated the diff for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

Added diagnosing of __float128 type usage.
See the summary of revision for details.

Mar 20 2020, 1:01 PM · Restricted Project
Fznamznon retitled D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage from [SYCL] Do not diagnose use of __float128 to [SYCL] Defer __float128 type usage diagnostics.
Mar 20 2020, 1:01 PM · Restricted Project

Feb 27 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

@rjmccall, Thank you very much for so detailed response, It really helps. I started working on implementation and I have a couple of questions/problems with this particular appoach.

Feb 27 2020, 5:38 AM · Restricted Project

Feb 17 2020

Fznamznon added inline comments to D70172: [CUDA][HIP][OpenMP] Emit deferred diagnostics by a post-parsing AST travese.
Feb 17 2020, 8:19 AM · Restricted Project

Feb 13 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

The right approach here is probably what we do in ObjC ARC when we see types that are illegal in ARC: in system headers, we allow the code but add a special UnavailableAttr to the declaration so that it can't be directly used.

That is straightforward enough that I think you should just do it instead of leaving this as technical debt.

I haven't considered something like this, because I'm not familiar with ObjC at all... I will give it a try, thanks.

Feb 13 2020, 9:01 AM · Restricted Project

Feb 11 2020

Fznamznon added a comment to D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.

I would add a check for the use of unsupported types in kernels. They should not be allowed to be used if target does not support it.

Feb 11 2020, 9:35 AM · Restricted Project
Fznamznon added reviewers for D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage: rsmith, rjmccall, ABataev.
Feb 11 2020, 12:24 AM · Restricted Project
Fznamznon created D74387: [OpenMP][SYCL] Improve diagnosing of unsupported types usage.
Feb 11 2020, 12:12 AM · Restricted Project

Jan 21 2020

Fznamznon accepted D73104: [Attr][Doc][NFC] Fix code snippet formatting for attribute documentation.

Thank you for this!

Jan 21 2020, 10:10 AM · Restricted Project

Dec 13 2019

Fznamznon added inline comments to D71016: [SYCL] Implement OpenCL kernel function generation.
Dec 13 2019, 6:34 AM · Restricted Project

Dec 10 2019

Fznamznon updated the diff for D71016: [SYCL] Implement OpenCL kernel function generation.

Updated tests using address space attributes added by D71005

Dec 10 2019, 1:19 AM · Restricted Project

Dec 4 2019

Fznamznon added a reviewer for D71016: [SYCL] Implement OpenCL kernel function generation: keryell.
Dec 4 2019, 9:51 AM · Restricted Project
Fznamznon created D71016: [SYCL] Implement OpenCL kernel function generation.
Dec 4 2019, 7:30 AM · Restricted Project
Fznamznon added reviewers for D71016: [SYCL] Implement OpenCL kernel function generation: bader, Naghasan, ABataev.
Dec 4 2019, 7:30 AM · Restricted Project

Dec 2 2019

Fznamznon accepted D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

LGTM with a couple of minor comments.

Dec 2 2019, 4:21 AM · Restricted Project

Nov 30 2019

Fznamznon added a comment to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

A couple of minor comments.

Nov 30 2019, 9:59 AM · Restricted Project

Jun 27 2019

Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Minor fix

Jun 27 2019, 8:52 AM · Restricted Project
Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Added warning diagnostic for sycl_kernel attribute.

Jun 27 2019, 7:18 AM · Restricted Project

Jun 24 2019

Fznamznon added reviewers for D63710: [SYCL] Re-use OpenCL sampler in SYCL device mode: bader, Anastasia.
Jun 24 2019, 5:33 AM · Restricted Project
Fznamznon created D63710: [SYCL] Re-use OpenCL sampler in SYCL device mode.
Jun 24 2019, 5:33 AM · Restricted Project

Jun 20 2019

Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

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

Jun 20 2019, 8:59 AM · Restricted Project
Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Updated sycl_kernel attribute documentation.

Jun 20 2019, 7:54 AM · Restricted Project

Jun 19 2019

Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
Jun 19 2019, 1:46 PM · Restricted Project
Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

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
Jun 19 2019, 1:42 PM · Restricted Project

Jun 18 2019

Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
Jun 18 2019, 8:01 AM · Restricted Project
Fznamznon added a comment to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Ping.

Jun 18 2019, 4:45 AM · Restricted Project

Jun 11 2019

Fznamznon added a comment to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

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

Jun 11 2019, 5:43 AM · Restricted Project

Jun 10 2019

Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Applied comments from @Anastasia

Jun 10 2019, 2:25 AM · Restricted Project

Jun 3 2019

Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
Jun 3 2019, 3:32 AM · Restricted Project

May 31 2019

Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
May 31 2019, 5:52 AM · Restricted Project

May 30 2019

vladimirlaz <vladimir.lazarev@intel.com> committed rG9f46e66220d4: [SYCL][NFC] Remove return from void function (authored by Fznamznon).
[SYCL][NFC] Remove return from void function
May 30 2019, 8:08 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG39e2c74a39cd: [SYCL] Allow host access for interoperability buffers (authored by Fznamznon).
[SYCL] Allow host access for interoperability buffers
May 30 2019, 8:08 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rGc9f21fad6162: [SYCL] Implement buffer constructor with iterators in accordance with spec (authored by Fznamznon).
[SYCL] Implement buffer constructor with iterators in accordance with spec
May 30 2019, 8:08 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG82fead6e0563: [SYCL] Implement basic sub-buffers support (authored by Fznamznon).
[SYCL] Implement basic sub-buffers support
May 30 2019, 8:08 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rGe8ac887f8642: [SYCL] Implement broadcasting vec::operator= (authored by Fznamznon).
[SYCL] Implement broadcasting vec::operator=
May 30 2019, 8:07 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG3e11f59cf86c: [SYCL] Initialize buffer range and size in interop constructor (authored by Fznamznon).
[SYCL] Initialize buffer range and size in interop constructor
May 30 2019, 8:07 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG1e514c82b139: [SYCL] Use std::enable_if properly in buffer contructors (authored by Fznamznon).
[SYCL] Use std::enable_if properly in buffer contructors
May 30 2019, 8:07 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG474c9343ad41: [SYCL] Implement buffer::set_write_back (authored by Fznamznon).
[SYCL] Implement buffer::set_write_back
May 30 2019, 8:07 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG0b4e9e9447e0: [SYCL] Implement buffer::has_property and buffer::get_property (authored by Fznamznon).
[SYCL] Implement buffer::has_property and buffer::get_property
May 30 2019, 8:07 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG0b9673fda707: [SYCL] Implement buffer::get_access with range and offset (authored by Fznamznon).
[SYCL] Implement buffer::get_access with range and offset
May 30 2019, 8:07 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG55b8dcea61a5: [SYCL] Add support for custom allocators in buffer (authored by Fznamznon).
[SYCL] Add support for custom allocators in buffer
May 30 2019, 8:06 AM
Vladimir Lazarev <vladimir.lazarev@intel.com> committed rG0baa4475cf85: [SYCL] Fix accessor to subset of buffer indexing (runtime part) (authored by Fznamznon).
[SYCL] Fix accessor to subset of buffer indexing (runtime part)
May 30 2019, 8:06 AM
Vladimir Lazarev <vladimir.lazarev@intel.com> committed rG7edf29c68344: [SYCL] Fix accessor to subset of buffer indexing (compiler part) (authored by Fznamznon).
[SYCL] Fix accessor to subset of buffer indexing (compiler part)
May 30 2019, 8:05 AM
Vladimir Lazarev <vladimir.lazarev@intel.com> committed rGf0d1636d09d3: [SYCL] Switch to use one definition of SYCL runtime classes in SYCL tests (authored by Fznamznon).
[SYCL] Switch to use one definition of SYCL runtime classes in SYCL tests
May 30 2019, 8:04 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG9fa265405a65: [SYCL] Do not always create new OpenCL buffers to copy from device to host (authored by Fznamznon).
[SYCL] Do not always create new OpenCL buffers to copy from device to host
May 30 2019, 8:03 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG5124d7c12978: [SYCL] Implement cl::sycl::buffer::reinterpret (authored by Fznamznon).
[SYCL] Implement cl::sycl::buffer::reinterpret
May 30 2019, 8:03 AM
vladimirlaz <vladimir.lazarev@intel.com> committed rG3ce5091397c3: [SYCL] Added constexpr specifier to standard functions declarations in… (authored by Fznamznon).
[SYCL] Added constexpr specifier to standard functions declarations in…
May 30 2019, 8:01 AM

May 28 2019

Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Applied comments from @Anastasia

  • Added documentation for sycl_kernel function
  • Added comments to Sema.h
  • Added -std=c++11 to test run lines
May 28 2019, 5:08 AM · Restricted Project
Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
May 28 2019, 4:36 AM · Restricted Project

May 23 2019

Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
May 23 2019, 1:08 AM · Restricted Project

May 22 2019

Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Minor fix.

May 22 2019, 1:12 AM · Restricted Project

May 21 2019

Fznamznon retitled D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining from [SYCL] Add support for SYCL device attributes to [SYCL] Implement SYCL device code outlining.
May 21 2019, 8:07 AM · Restricted Project
Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Added semantics for new attributes

May 21 2019, 7:59 AM · Restricted Project

Apr 19 2019

Fznamznon added a comment to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

I tried to reuse OpenCL kernel attribute with "__kernel" keyword in our current SYCL implementation. PR with this try is here - https://github.com/intel/llvm/pull/97
Now It looks feasible but with a couple notes:
From SYCL specification "SYCL is designed to be as close to standard C++ as possible. Standard C++ compiler can compile the SYCL programs and they will run correctly on host CPU." So SYCL doesn't provide non-standard kernel keyword which is provided by OpenCL. Due this fact it's not possible to add kernel keyword as in OpenCL, it will prevent compilation of following valid SYCL code:

Apr 19 2019, 8:01 AM · Restricted Project

Apr 17 2019

Fznamznon added a comment to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

It would be better to rename clang/test/SemaSYCL/device-attrubutes.cpp to clang/test/SemaSYCL/device-attributes.cpp

Apr 17 2019, 8:16 AM · Restricted Project

Apr 16 2019

Fznamznon added a comment to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
Apr 16 2019, 7:08 AM · Restricted Project
Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
Apr 16 2019, 6:24 AM · Restricted Project
Fznamznon added inline comments to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.
Apr 16 2019, 5:34 AM · Restricted Project
Fznamznon added a comment to D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Just to understand how this will work. I would imagine you can have a device function definition preceding its use. When the function is being parsed it's not known yet whether it will be called from the device or not, so it won't be possible to set the language mode correctly and hence provide the right diagnostics. So is the plan to launch a separate parsing phase then just to extract the call graph and annotate the device functions?

I'm not an expert in clang terminology but I will try briefly explain our current implementation approach.
In SYCL all kernel functions should be template functions so these functions have a deferred instantiation. If we found that we instantiated a sycl kernel function - we add it to a special array with sycl device functions (you can see the corresponding code here and here, actually AddSyclKernel adds declaration to the special array inside the Sema). After performing
pending instantiations we run a recursive AST visitor for each SYCL kernel to mark all device functions and add them to a special array with SYCL device functions (here we start traverse AST from MarkDevice function, code of MarkDevice is here).
To get a correct set of SYCL device functions in produced module we added a check for all declarations inside the CodeGen on sycl_device attribute existence - so it will ignore declarations without sycl_device attribute if we are compiling for SYCL device (code is here). But with this check it's possible situation when function was parsed and ignored by the CodeGen before we added sycl_device attribute to it so we added yet another parsing action inside the clang::ParseAST to generate code for all SYCL device functions from the special array (code is here).

Apr 16 2019, 5:06 AM · Restricted Project

Apr 15 2019

Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Applied comment from @bader.

Apr 15 2019, 4:39 AM · Restricted Project
Fznamznon updated the diff for D60455: [SYCL] Add sycl_kernel attribute for accelerated code outlining.

Applied comments from @bader

Apr 15 2019, 4:05 AM · Restricted Project