Rebased.

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.

In D81641#2088446, @jdoerfert wrote:

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

Generalized diagnostic between SYCL and OpenMP.

Fixed code style.

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

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

Fix is here D80829 .

In D74387#2063423, @ABataev wrote:

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

Included test cases from Johannes.

In D77918#2060446, @jdoerfert wrote:

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

Applied comments from Johannes.
Fixed failing tests.

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?

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

In D74387#2023200, @jdoerfert wrote:

In D74387#1992682, @Fznamznon wrote:

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?

Ping.

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) {}
}

In D74387#1974981, @jdoerfert wrote:

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?

In D74387#1970374, @jdoerfert wrote:

In D74387#1969891, @Fznamznon wrote:

In D74387#1967386, @jdoerfert wrote:

In D74387#1967289, @Fznamznon wrote:

In D74387#1965634, @jdoerfert wrote:

In D74387#1964483, @Fznamznon wrote:

In D74387#1950593, @jdoerfert wrote:

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?

In D74387#1967386, @jdoerfert wrote:

In D74387#1967289, @Fznamznon wrote:

In D74387#1965634, @jdoerfert wrote:

In D74387#1964483, @Fznamznon wrote:

In D74387#1950593, @jdoerfert wrote:

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.

In D74387#1965634, @jdoerfert wrote:

In D74387#1964483, @Fznamznon wrote:

In D74387#1950593, @jdoerfert wrote:

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 ...).

In D74387#1950593, @jdoerfert wrote:

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

Apply comments, rebase.

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

Ping.

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

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

@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.

In D74387#1869845, @rjmccall wrote:

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.

In D74387#1869819, @ABataev wrote:

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.

Thank you for this!

Updated tests using address space attributes added by D71005

LGTM with a couple of minor comments.

A couple of minor comments.

Minor fix

Added warning diagnostic for sycl_kernel attribute.

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

Updated sycl_kernel attribute documentation.

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

Ping.

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

Applied comments from @Anastasia

Applied comments from @Anastasia

• Added documentation for sycl_kernel function
• Added comments to Sema.h
• Added -std=c++11 to test run lines

Minor fix.

Added semantics for new attributes

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:

In D60455#1470318, @keryell wrote:

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

In D60455#1467279, @aaron.ballman wrote: