[OpenMP][SYCL] Improve diagnosing of unsupported types usage

[Fznamznon]

Summary

Diagnostic is emitted if some declaration of unsupported type
declaration is used inside device code.
Memcpy operations for structs containing member with unsupported type
are allowed. Fixed crash on attempt to emit diagnostic outside of the
functions.

The approach is generalized between SYCL and OpenMP.
CUDA/OMP deferred diagnostic interface is going to be used for SYCL device.

[ABataev]

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.

[rjmccall]

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.

[Fznamznon]

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.

Yeah, I think so. We tried to make it using deferred diagnostics. Unfortunately there isn't a single 'create declaration' type place that we could diagnose this. This resulted to a lot of changes around Sema and probably unhanded cases. For example we need to diagnose each appearance of __float128 type in device code. And __float128 can appear in device code through so many ways, for example through auto types variable declaration and initializing it with __float128 value captured from the host, like this:

 // HOST CODE
  __float128 B = 1; // No errors
...
// DEVICE CODE
  kernel<class some_kernel>([=]() {
          auto C = B; }); // Problem, C will actually have __float128 type!

And for example, we can't just trigger on __float128 type appearance in some code, like the diagnosic which I'm disabling does, because I believe that some unevaluated contexts shouldn't trigger errors, because they don't bring the unsupported type to the device code:

template<typename t> void foo(){};
__float128 nonemittedfunc();

// DEVICE CODE
foo<__float128>(); // This shouldn't bring errors
std::conditional_t<SomeI < 1, decltype(nonemittedfunc()), int> SomeVar; // This shouldn't bring errors

The whole patch with test cases is available here https://github.com/intel/llvm/pull/1040 .
We decided to disable this until we figure out the way how to properly diagnose this.

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.

[Fznamznon]

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.

Hi @rjmccall , I assume, I took a look at this.
Let's imagine, I will try to diagnose __float128 type using already implemented functionality. It seems like I need to call something like

S.DelayedDiagnostics.add(                                        
    sema::DelayedDiagnostic::makeForbiddenType(loc,              
        diag::err_type_unsupported, type, "__float128"));
`

right?
I suppose, then this diagnostic will be saved and emitted inside function named handleDelayedForbiddenType.
Here it checks that this forbidden type is actually allowed and emits a diagnostic if it's not.
The first problem that handleDelayedForbiddenType is called too early. We don't know in this place whether we are in SYCL device code or not. Because basically entry point to SYCL device code is a template function with sycl_kernel attribute, and every function it calls become a device function. So we only know where is device code only after templates instantiation, it happens a bit later after handleDelayedForbiddenType call.

It seems that the second problem is the same problem which prevented me from implementing diagnosing of __float128 type through CUDA/OMP deferred diagnostics (I mentioned my attempt in the last comment https://reviews.llvm.org/D74387#1870014). I still need to find best place for diagnostic issuing. It seems that there are so many places where type can actually be introduced to resulting LLVM IR module, and in some of them I need to check some additional conditions to do not prevent __float128 usage when it actually doesn't introduce forbidden type to resulting LLVM IR module.

Please, correct me if I don't understand something or said something wrong.
I would appreciate if you had some advices.

Thanks a lot.

[rjmccall]

In D74387#1874612, @Fznamznon wrote:

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.

Hi @rjmccall , I assume, I took a look at this.
Let's imagine, I will try to diagnose __float128 type using already implemented functionality. It seems like I need to call something like

S.DelayedDiagnostics.add(                                        
    sema::DelayedDiagnostic::makeForbiddenType(loc,              
        diag::err_type_unsupported, type, "__float128"));
`

right?
I suppose, then this diagnostic will be saved and emitted inside function named handleDelayedForbiddenType.
Here it checks that this forbidden type is actually allowed and emits a diagnostic if it's not.

This isn't quite right. The goal here is to delay the diagnostic *twice*. The first delay is between the point where we parse/process the type (i.e. SemaType) and the point where we've fully processed the declaration that the type is part of (i.e. SemaDecl). That's the point where we call handleDelayedForbiddenType, and you're right that it's too early to know whether the declaration is really device code. However, you're missing something important about how handleDelayedForbiddenType works: it's never really trying to suppress the diagnostic completely, but just to delay it for certain declarations until the point that the declaration is actually used, under the hope that in fact it will never be used and everything will work out. For ARC, we chose to delay for all declarations in system headers, under the assumption that (1) system headers will never introduce functions that have to be emitted eagerly, and (2) we always want to warn people about problematic code in their own headers. Those choices don't really fit SYCL's use case, and you should change the logic in isForbiddenTypeAllowed to delay your diagnostic for essentially all declarations (except kernels?), since effectively all non-kernel code in device mode is lazily emitted. But if you do that, it should combine well with CUDA/OMP deferred diagnostics:

• If foo uses __float128 (whether in its signature or internally), that is invalid in device mode, but the diagnostic will be delayed by the forbidden-type mechanism, meaning that it will become an unavailable attribute on foo.
• If bar uses foo, that use is invalid in device mode (because of the unavailable attribute), but the diagnostic will be delayed via the standard CUDA/OMP mechanism because we don't know yet whether bar should be emitted as a device function.
• If kernel uses bar, that will trigger the emission of the delayed diagnostics of bar, including the use-of-unavailable-function diagnostic where it uses foo. It should be straightforward to specialize this diagnostic so that it reports the error by actually diagnosing the use of __float128 at the original location (which is recorded in the unavailable attribute) and then just adding a note about how foo is used by bar.

It seems that the second problem is the same problem which prevented me from implementing diagnosing of __float128 type through CUDA/OMP deferred diagnostics (I mentioned my attempt in the last comment https://reviews.llvm.org/D74387#1870014). I still need to find best place for diagnostic issuing. It seems that there are so many places where type can actually be introduced to resulting LLVM IR module, and in some of them I need to check some additional conditions to do not prevent __float128 usage when it actually doesn't introduce forbidden type to resulting LLVM IR module.

The key thing here is that all uses should be associated with some top-level declaration that's either eagerly-emitted in device mode or not.

[Fznamznon]

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

• If foo uses __float128 (whether in its signature or internally), that is invalid in device mode, but the diagnostic will be delayed by the forbidden-type mechanism, meaning that it will become an unavailable attribute on foo.

So, for example if some variable is declared with __float128 type, we are adding to parent function Unavaliable attribute, right?

• If bar uses foo, that use is invalid in device mode (because of the unavailable attribute), but the diagnostic will be delayed via the standard CUDA/OMP mechanism because we don't know yet whether bar should be emitted as a device function.
• If kernel uses bar, that will trigger the emission of the delayed diagnostics of bar, including the use-of-unavailable-function diagnostic where it uses foo. It should be straightforward to specialize this diagnostic so that it reports the error by actually diagnosing the use of __float128 at the original location (which is recorded in the unavailable attribute) and then just adding a note about how foo is used by bar.

Consider following example (this is absolutely valid SYCL code, except __float128 usage):

// Host code:
__float128 A;
// Everything what lambda passed to `sycl_kernel` calls becomes device code. Capturing of host variables means that these variables will be passed to device by value, so using of A in this lambda is invalid.
sycl_kernel<class kernel_name>([=]() {auto B = A});

In this case we add unavailable attribute to parent function for variable A declaration. But this function is not called from device code. Please correct me if I'm wrong but it seems that we need to diagnose not only functions, but also usages of any declarations with unavailable attribute including variable declarations, right?

In addition, there are a couple of problems with this approach, for example with unevaluated sizeof context, i.e. code like this:

sycl_kernel<class kernel_name>([=]() {int A = sizeof(__float128);});

is diagnosed too, I think this is not correct.

I can upload what I have now to this review if it will help better (or maybe we will understand that I'm doing something wrong).

I'm also thinking about a bit another approach:

• If some declaration uses __float128 it will become an unavailable attribute on this declaration. That means we don't always add unavailable attribute to the function which uses __float128 internally.
• In the place where clang actually emits use-of-unavailable-declaration diagnostics (somewhere in DoEmitAvailabilityWarning function, defined in SemaAvailability.cpp) for SYCL, we make these diagnostics deferred using CUDA/OMP deferred diagnostics mechanism (using SYCL-specific analog of function like diagIfOpenMPDeviceCode/CUDADiagIfDeviceCode).

But, for example, this won't emit diagnostics for simple variable declarations in device code which has __float128 type, but is not used anywhere else.

I'm also curious about OpenMP handling of this "unsupported type" problem. @ABataev , Am I right that in OpenMP such diagnostics are emitted only if forbidden type is used in some arithmetical operations? Is it enough to prevent problems on various GPU devices which don't support this type?

[rjmccall]

In D74387#1895264, @Fznamznon wrote:

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

• If foo uses __float128 (whether in its signature or internally), that is invalid in device mode, but the diagnostic will be delayed by the forbidden-type mechanism, meaning that it will become an unavailable attribute on foo.

So, for example if some variable is declared with __float128 type, we are adding to parent function Unavaliable attribute, right?

That's how it's supposed to work. I can't guarantee that it will actually always work that way, because I'm sure you'll be pushing on this code in some new ways.

• If bar uses foo, that use is invalid in device mode (because of the unavailable attribute), but the diagnostic will be delayed via the standard CUDA/OMP mechanism because we don't know yet whether bar should be emitted as a device function.
• If kernel uses bar, that will trigger the emission of the delayed diagnostics of bar, including the use-of-unavailable-function diagnostic where it uses foo. It should be straightforward to specialize this diagnostic so that it reports the error by actually diagnosing the use of __float128 at the original location (which is recorded in the unavailable attribute) and then just adding a note about how foo is used by bar.

Consider following example (this is absolutely valid SYCL code, except __float128 usage):

// Host code:
__float128 A;
// Everything what lambda passed to `sycl_kernel` calls becomes device code. Capturing of host variables means that these variables will be passed to device by value, so using of A in this lambda is invalid.
sycl_kernel<class kernel_name>([=]() {auto B = A});

In this case we add unavailable attribute to parent function for variable A declaration. But this function is not called from device code. Please correct me if I'm wrong but it seems that we need to diagnose not only functions, but also usages of any declarations with unavailable attribute including variable declarations, right?

Right. The diagnosis side of that should already happen — unavailable diagnostics apply to uses of any kind of declaration, not just functions or variables. Current delayed diagnostics should be enough to make the unavailable attribute get applied to the global variable A in your example, since it's a use from the declarator. If SYCL supports C++-style dynamic global initializers, you'll probably need to add code so that uses of __float128 within a global initializer get associated with the global, which currently won't happen because the initializer isn't "in scope". But there are at least two other patches underway that are dealing with similar issues: https://reviews.llvm.org/D71227 and https://reviews.llvm.org/D70172.

In addition, there are a couple of problems with this approach, for example with unevaluated sizeof context, i.e. code like this:

sycl_kernel<class kernel_name>([=]() {int A = sizeof(__float128);});

is diagnosed too, I think this is not correct.

Okay, that's a much thornier problem if you want to allow that. What you're talking about is essentially the difference between an evaluated and an unevaluated context, but we don't generally track that for uses of *types*. It's much easier to set things up so that you only complain about uses of *values* like the global variable A if they're in evaluated expressions, but types are never really "evaluated" in the same way that expressions are, so it's complicated.

I think that's a very separable question, so I would recommend you focus on the first problem right now, and then if you really care about allowing sizeof(__float128), we can approach that later.

I can upload what I have now to this review if it will help better (or maybe we will understand that I'm doing something wrong).

I'm also thinking about a bit another approach:

• If some declaration uses __float128 it will become an unavailable attribute on this declaration. That means we don't always add unavailable attribute to the function which uses __float128 internally.
• In the place where clang actually emits use-of-unavailable-declaration diagnostics (somewhere in DoEmitAvailabilityWarning function, defined in SemaAvailability.cpp) for SYCL, we make these diagnostics deferred using CUDA/OMP deferred diagnostics mechanism (using SYCL-specific analog of function like diagIfOpenMPDeviceCode/CUDADiagIfDeviceCode).

Sure, but you'll have to write a custom walk of the body looking for uses of the type that you don't like; that seems like a lot of work to get right, and it'll tend to fail "open", i.e. allowing things you don't want to allow, whereas this approach will tend to fail "closed", i.e. tending towards being conservatively correct.

[Fznamznon]

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

[Fznamznon]

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

[Fznamznon]

Ping.

[rjmccall]

clang/include/clang/Sema/Sema.h

12417

Will this collect notes associated with the diagnostic correctly?

clang/lib/Sema/SemaAvailability.cpp

479 ↗

(On Diff #251982)

All of the other cases are setting this to a note, not an error, so I suspect this will read wrong.

534 ↗

(On Diff #251982)

Are you sure you want to be applying this to all of the possible diagnostics here, rather than just for SYCLForbiddenType unavailable attributes?

clang/lib/Sema/SemaDecl.cpp

18124

So you want to emit it for the definition in addition to emitting it for specific specializations?

clang/lib/Sema/SemaDeclAttr.cpp

7771 ↗

(On Diff #251982)

I wonder if it's reasonable to treat all forbidden types the same here or if we want different functions for the ARC and SYCL use cases.

[Fznamznon]

clang/include/clang/Sema/Sema.h

12417

Could you please make your question a bit more concrete?
This function is supposed to work in the same way as Sema::CUDADiagIfDeviceCode and Sema::diagIfOpenMPDeviceCode . It emits given diagnostic if the current context is known as "device code" and makes this diagnostic deferred otherwise. It uses the DeviceDiagBuilder which was implemented earlier. This DeviceDiagBuilder also tries to emit callstack notes for the given diagnostics. Do you mean these callstack notes or something else?

clang/lib/Sema/SemaAvailability.cpp

479 ↗

(On Diff #251982)

Yes, this is not a note. For such samples:

int main() {
  __float128 CapturedToDevice = 1;
  kernel<class variables>([=]() {
    decltype(CapturedToDevice) D;
  });
}

It looks like this:

float128.cpp:63:14: error: 'CapturedToDevice' is unavailable
    decltype(CapturedToDevice) D;
             ^
float128.cpp:59:14: error: '__float128' is not supported on this target   /// This emitted instead of note 
  __float128 CapturedToDevice = 1;
             ^

I had feeling that it should probably be a note. But there is no implemented note for unsupported types. I think I can add a new one if it will make it better. Should I?

534 ↗

(On Diff #251982)

I suppose it is reasonable if we want to reuse unavaliable attribute for other SYCL use cases. Plus, In SYCL we don't know where is device code until we instantiate templates, it happens late, so we have to defer any diagnostic while compiling for device, otherwise we can point to host code where much more is allowed.

clang/lib/Sema/SemaDecl.cpp

18124

Somehow diagnostics are emitted only for the definitions.
Without this change diagnostics aren't emitted at all.

clang/lib/Sema/SemaDeclAttr.cpp

7771 ↗

(On Diff #251982)

I think it could be reasonable if we will have forbidden type cases for SYCL sometime. For now, I don't see the purpose in a separate function for SYCL.

[Fznamznon]

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

[jdoerfert]

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

[rjmccall]

clang/include/clang/Sema/Sema.h

12417

Logically, notes that are emitted after a warning or error are considered to be part of that diagnostic. A custom DiagBuilder that only redirects the main diagnostic but allows the notes to still be emitted will effectively cause those notes to misleadingly follow whatever previous diagnostic might have been emitted.

I call this out specifically because some of the places where you're using this still seem to try to emit notes afterwards, at least in some cases. It's possible that CUDADiagIfDeviceCode happens to not be used in such a way. Really I'm not sure this conditional DiagBuilder approach was a good idea the first time, and I think we should probably reconsider rather than duplicating it.

clang/lib/Sema/SemaAvailability.cpp

479 ↗

(On Diff #251982)

Yeah, this should be a note, like "note: variable is unavailable because it uses a type '__float128' that is not supported on this target". You should add that.

534 ↗

(On Diff #251982)

My point is actually the reverse of that. This code path is also used for normal unavailable attributes, not just the special ones you're synthesizing. Diagnostics from the use of explicitly-unavailable declarations shouldn't get any special treatment here, no more than you'd give special treatment to a diagnostic arising from an attempt to assign a pointer into a float. In the logic above where you recognize IR_SYCLForbiddenType, I think you should just check whether you should transitively defer the diagnostic and, if so, do so and then bail out of this function early. That might mean you don't need the custom DiagBuilder, too.

clang/lib/Sema/SemaDecl.cpp

18124

Hmm. We might be marking the template pattern invalid; that could result in all sorts of diagnostics being suppressed. We definitely shouldn't be marking things invalid without emitting an eager diagnostic.

[Fznamznon]

Apply comments, rebase.

[Fznamznon]

clang/include/clang/Sema/Sema.h

12417

I think if there are some notes associated with the main diagnostic and we want to make this diagnostic deferred by using SYCLDiagIfDeviceCode, we have to use this function SYCLDiagIfDeviceCode for notes as well. In my changes I didn't do so because I didn't expect notes emitted after new diagnostic.
In our SYCL implementation we find function like SYCLDiagIfDeviceCode pretty useful because we don't know where is device code until templates are instantiated. We need some mechanism to defer diagnostics pointing to unsupported features used in device code.
Do you have better approach in mind?

clang/lib/Sema/SemaAvailability.cpp

479 ↗

(On Diff #251982)

Okay, done.

534 ↗

(On Diff #251982)

Okay, I understand. I was under impression that unavailable attributes can appear only for ObjC ARC, so It is safe to defer everything in SYCL, so I moved calls of SYCLDiagIfDeviceCode as you requested.
It's a bit unclear how to avoid custom DiagBuilder here.

[Fznamznon]

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?

[jdoerfert]

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

[Fznamznon]

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.

[jdoerfert]

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.

[Fznamznon]

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?

[jdoerfert]

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?

[Fznamznon]

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?

Roughly speaking, SYCL is a wrapper over OpenCL. SYCL device compiler should be able to produce device code module in a form acceptable by OpenCL backends. For this purpose we use SPIR-V intermediate language (https://www.khronos.org/registry/spir-v/specs/unified1/SPIRV.html). We transform LLVM IR emitted by clang (in SYCL device mode) into SPIR-V, then feed it to OpenCL backends. To be able to do it, produced SPIR-V must be valid and do not require additional features/capabilities comparing with SPIR-V produced from pure OpenCL, otherwise OpenCL backends just don't work with it. Nor OpenCL neither SPIRV doesn't support __float128 type, for example. From SPIR-V spec:

Scalar floating-point types can be parameterized only as 32 bit, plus any additional sizes enabled by capabilities (i.e. 16 and 64 for some devices).

Right now It is not possible to produce valid SPIR-V from LLVM IR containing unsupported types. We use official Khronos SPIRV translator (https://github.com/KhronosGroup/SPIRV-LLVM-Translator). SPIR-V translator relies on clang to prohibit unsupported features, so they are not expected in LLVM IR. That is why we might need completely prohibit (or maybe we need to replace it in resulting LLVM module completely if it is possible?) now.

I'm also curious why OpenMP can just allow presence of unsupported type in the resulting module? Doesn't it produce any problems while compiling code by device-specific back-end for some specific device which don't support such type?

I also think that we need to generalize approaches between OpenMP/SYCL/CUDA/HIP. We can start with generalized diagnostic which points to using of unsupported type at least, then we can add additional restriction for SYCL (or other programming models) if we need one. @bader , @erichkeane please comment if you don't agree.

[jdoerfert]

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?

[Fznamznon]

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?

Yes, it can work for SYCL without additional diagnostics if it is possible to replace __float128 with char [16] everywhere (including struct definitions and so on) in the resulting LLVM IR module.

[Fznamznon]

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.

[Fznamznon]

Ping.

[jdoerfert]

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?

[Fznamznon]

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?

I'm not sure that I've discovered all examples and problems, but I have a couple of ones. I started with adapting current implementation for OpenMP and right now I'm analyzing corresponding OpenMP test fails (i.e. clang/test/OpenMP/nvptx_unsupported_type_messages.cpp and clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp). There are a lot of differences between the old approach and new one, which I'm working on. The new diagnostic triggers more errors than the old one, so I'd like to understand in which concrete cases we shouldn't emit diagnostic. For example you mentioned that memcopy-like operations should be ok in device code.
Right now the current implementation of the diagnostic also emits errors for sample like this:

struct T {
  char a;
  __float128 f;
  char c;
};

#pragma omp declare target
T a;
T b = a; // The diagnostic is triggered here, because implicit copy constructor uses unavailable field
#pragma omp end declare target

Should we emit errors in such case too?

[jdoerfert]

In D74387#2026844, @Fznamznon wrote:

I'm not sure that I've discovered all examples and problems, but I have a couple of ones. I started with adapting current implementation for OpenMP and right now I'm analyzing corresponding OpenMP test fails (i.e. clang/test/OpenMP/nvptx_unsupported_type_messages.cpp and clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp). There are a lot of differences between the old approach and new one, which I'm working on. The new diagnostic triggers more errors than the old one, so I'd like to understand in which concrete cases we shouldn't emit diagnostic. For example you mentioned that memcopy-like operations should be ok in device code.

You can reach me here or via email to discuss more. We also can do it over openmp-dev if you like :)

Right now the current implementation of the diagnostic also emits errors for sample like this:

struct T {
  char a;
  __float128 f;
  char c;
};

#pragma omp declare target
T a;
T b = a; // The diagnostic is triggered here, because implicit copy constructor uses unavailable field
#pragma omp end declare target

Should we emit errors in such case too?

Preferably, I would allow the above case, or in general trivial copies of unavailable basic types. What I would like to happen is that we memcpy the unavailable field.
I guess if we could "simply" replace the unavailable types in the device code right away with the byte array replacement, most things should fall into place. Basically,
we could provide even provide the replacements in a header that we include automatically:

clang/lib/Headers/OpenMP/typedefs.h:

#ifdef __DEFINE_FLOAT128__
typedef char __float128[__FLOAT128_SIZE__] alignas(__FLOAT128_ALIGNMENT__);

#undef __DEFINE_FLOAT128__
#undef __FLOAT128_SIZE__
#undef __FLOAT128_ALIGNMENT__
#endif

Now copy constructors (and other "OK" uses) should work fine. If people use the member in anything that actually doesn't work on a char array, they get a (probably ugly) error.
Preferably we would intercept the diagnose messages or at least issue a note if we see __float128 used, maybe along the lines of:

Note: The target does not support operations on `__float128` types. Values of this type are consequently represented by character arrays of appropriate size and alignment.

[Fznamznon]

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

[Fznamznon]

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?

[jdoerfert]

In D74387#2053742, @Fznamznon wrote:

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

Great, thanks a lot!

In D74387#2054337, @Fznamznon wrote:

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?

Yes, please! The test case (which I added) is broken and would result in a crash when you actually ask for PTX and not IR: https://godbolt.org/z/vL5Biw
This is exactly what we need to diagnose :)

---

I think the code looks good and this looks like a really nice way to fix this properly.

I inlined some questions. We might need to add some test coverage (if we haven't already), e.g., for the memcpy case. For example in OpenMP an object X with such types should be ok in a map(tofrom:X) clause.

clang/lib/Sema/Sema.cpp

1727

Nit: Move below CheckType to avoid shadowing and confusion with the arg there.

clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp

21

Why is this not diagnosed? I mean we cannot assign 15 on the device, can we? Or does it work because it is a constant (and we basically just memcpy something)?

If it's the latter, do we have a test in the negative version that makes sure T(int i) : a(i), f(i) {} is flagged?

81

Just checking, we verify in the other test this would result in an error, right?

[Fznamznon]

Applied comments from Johannes.
Fixed failing tests.

[Fznamznon]

clang/lib/Sema/Sema.cpp

1727

Done, thanks

clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp

21

Unfortunately, nor this case neither T(int i) : a(i), f(i) {} is not diagnosed. This happens because DiagnoseUseOfDecl call seems missing for member initializers, not because there is memcpy. So, for example, such case is diagnosed:

struct B {
  __float128 a;
};
#pragma omp declare target
void foo() {
  B var = {1}; // error: 'a' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it                
}

DiagnoseUseOfDecl function is called in so many cases and I guess it is meant to be called on each usage of each declaration, that is why I think the correct fix is add call to DiagnoseUseOfDecl somewhere near building of member initializers . This change even doesn't break my local check-clang LIT tests run, but I'm not really sure that such change is in scope of this patch, because DiagnoseUseOfDecl contains a lot of other diagnostics as well.

81

Yes, I added such test case in nvptx_unsupported_type_messages.cpp .

[jdoerfert]

This change even doesn't break my local check-clang LIT tests run, but I'm not really sure that such change is in scope of this patch, because DiagnoseUseOfDecl contains a lot of other diagnostics as well.

Fair. Let's do the following. We go with this as it is a clear improvement and almost complete. If you could provide a follow up to call DiagnoseUseOfDecl, or maybe just to the part we actually need, for member initialization, we can ask Clang folks to take a look. I think there is a reasonable way forward.

LGTM. Thanks a lot for implementing this in a generic and sharable way!

[jdoerfert]

Can you include these:

long double qa, qb;
decltype(qa + qb) qc;
double qd[sizeof(-(-(qc * 2)))];

[Fznamznon]

Included test cases from Johannes.

[ABataev]

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

[Fznamznon]

In D74387#2063423, @ABataev wrote:

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

It seems this patch caused asking size of dependent type, AST context doesn't seem expecting it. I'll provide follow up fix shortly. Sorry for inconvenience.

[Fznamznon]

Fix is here D80829 .

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_expected_private_copy_for_allocate : Error<
   "the referenced item is not found in any private clause on the same directive">;
 def err_omp_stmt_depends_on_loop_counter : Error<
   "the loop %select{initializer|condition}0 expression depends on the current loop control variable">;
 def err_omp_invariant_or_linear_dependency : Error<
   "expected loop invariant expression or '<invariant1> * %0 + <invariant2>' kind of expression">;
 def err_omp_wrong_dependency_iterator_type : Error<
   "expected an integer or a pointer type of the outer loop counter '%0' for non-rectangular nests">;
-def err_omp_unsupported_type : Error <
-  "host requires %0 bit size %1 type support, but device '%2' does not support it">;
+def err_device_unsupported_type : Error <
+  "%0 requires %1 bit size %2 type support, but device '%3' does not support it">;
 def err_omp_lambda_capture_in_declare_target_not_to : Error<
   "variable captured in declare target region must appear in a to clause">;
 def err_omp_device_type_mismatch : Error<
   "'device_type(%0)' does not match previously specified 'device_type(%1)' for the same declaration">;
 def err_omp_wrong_device_function_call : Error<
   "function with 'device_type(%0)' is not available on %select{device|host}1">;
 def note_omp_marked_device_type_here : Note<"marked as 'device_type(%0)' here">;
 def warn_omp_declare_target_after_first_use : Warning<

clang/include/clang/Sema/Sema.h

private:
   int getNumberOfConstructScopes(unsigned Level) const;
 
   /// Push new OpenMP function region for non-capturing function.
   void pushOpenMPFunctionRegion();
 
   /// Pop OpenMP function region for non-capturing function.
   void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI);
 
-  /// Check if the expression is allowed to be used in expressions for the
-  /// OpenMP devices.
-  void checkOpenMPDeviceExpr(const Expr *E);
-
   /// Checks if a type or a declaration is disabled due to the owning extension
   /// being disabled, and emits diagnostic messages if it is disabled.
   /// \param D type or declaration to be checked.
   /// \param DiagLoc source location for the diagnostic message.
   /// \param DiagInfo information to be emitted for the diagnostic message.
   /// \param SrcRange source range of the declaration.
   /// \param Map maps type or declaration to the extensions.
   /// \param Selector selects diagnostic message: 0 for type and 1 for
public:
   ///  // Variable-length arrays are not allowed in NVPTX device code.
   ///  if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported))
   ///    return ExprError();
   ///  // Otherwise, continue parsing as normal.
   DeviceDiagBuilder diagIfOpenMPHostCode(SourceLocation Loc, unsigned DiagID);
 
   DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID);
 
+  /// Check if the expression is allowed to be used in expressions for the
+  /// offloading devices.
+  void checkDeviceDecl(const ValueDecl *D, SourceLocation Loc);
+
   enum CUDAFunctionTarget {
     CFT_Device,
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };
 
public:
   /// Describes the reason a calling convention specification was ignored, used
   /// for diagnostics.
   enum class CallingConventionIgnoredReason {
     ForThisTarget = 0,
     VariadicFunction,
     ConstructorDestructor,
     BuiltinFunction
   };
+  /// Creates a DeviceDiagBuilder that emits the diagnostic if the current
+  /// context is "used as device code".
+  ///
+  /// - If CurLexicalContext is a kernel function or it is known that the
+  ///   function will be emitted for the device, emits the diagnostics
+  ///   immediately.
+  /// - If CurLexicalContext is a function and we are compiling
+  ///   for the device, but we don't know that this function will be codegen'ed
+  ///   for devive yet, creates a diagnostic which is emitted if and when we
+  ///   realize that the function will be codegen'ed.
+  ///
+  /// Example usage:
+  ///
+  /// Diagnose __float128 type usage only from SYCL device code if the current
+  /// target doesn't support it
+  /// if (!S.Context.getTargetInfo().hasFloat128Type() &&
+  ///     S.getLangOpts().SYCLIsDevice)
+  ///   SYCLDiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128";
+  DeviceDiagBuilder SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID);
+
+  /// Check whether we're allowed to call Callee from the current context.
+  ///
+  /// - If the call is never allowed in a semantically-correct program
+  ///   emits an error and returns false.
+  ///
+  /// - If the call is allowed in semantically-correct programs, but only if
+  ///   it's never codegen'ed, creates a deferred diagnostic to be emitted if
+  ///   and when the caller is codegen'ed, and returns true.
+  ///
+  /// - Otherwise, returns true without emitting any diagnostics.
+  ///
+  /// Adds Callee to DeviceCallGraph if we don't know if its caller will be
+  /// codegen'ed yet.
+  bool checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee);
 };
 
 /// RAII object that enters a new expression evaluation context.
 class EnterExpressionEvaluationContext {
   Sema &Actions;
   bool Entered = true;
 
 public:

clang/lib/Sema/CMakeLists.txt

add_clang_library(clangSema
   SemaModule.cpp
   SemaObjCProperty.cpp
   SemaOpenMP.cpp
   SemaOverload.cpp
   SemaPseudoObject.cpp
   SemaStmt.cpp
   SemaStmtAsm.cpp
   SemaStmtAttr.cpp
+  SemaSYCL.cpp
   SemaTemplate.cpp
   SemaTemplateDeduction.cpp
   SemaTemplateInstantiate.cpp
   SemaTemplateInstantiateDecl.cpp
   SemaTemplateVariadic.cpp
   SemaType.cpp
   TypeLocBuilder.cpp
 

clang/lib/Sema/Sema.cpp

 
 Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) {
   if (LangOpts.OpenMP)
     return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID)
                                    : diagIfOpenMPHostCode(Loc, DiagID);
   if (getLangOpts().CUDA)
     return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID)
                                       : CUDADiagIfHostCode(Loc, DiagID);
+
+  if (getLangOpts().SYCLIsDevice)
+    return SYCLDiagIfDeviceCode(Loc, DiagID);
+
   return DeviceDiagBuilder(DeviceDiagBuilder::K_Immediate, Loc, DiagID,
                            getCurFunctionDecl(), *this);
 }
 
+void Sema::checkDeviceDecl(const ValueDecl *D, SourceLocation Loc) {
+  if (isUnevaluatedContext())
+    return;
+
+  Decl *C = cast<Decl>(getCurLexicalContext());
+
+  // Memcpy operations for structs containing a member with unsupported type
+  // are ok, though.
+  if (const auto *MD = dyn_cast<CXXMethodDecl>(C)) {
+    if ((MD->isCopyAssignmentOperator() || MD->isMoveAssignmentOperator()) &&
+        MD->isTrivial())
+      return;
+
+    if (const auto *Ctor = dyn_cast<CXXConstructorDecl>(MD))
+      if (Ctor->isCopyOrMoveConstructor() && Ctor->isTrivial())
+        return;
+  }
+
+  auto CheckType = [&](QualType Ty) {
+    if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
+        ((Ty->isFloat128Type() ||
+          (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&
+         !Context.getTargetInfo().hasFloat128Type()) ||
+        (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
+         !Context.getTargetInfo().hasInt128Type())) {
+      targetDiag(Loc, diag::err_device_unsupported_type)
+          << D << static_cast<unsigned>(Context.getTypeSize(Ty)) << Ty
+          << Context.getTargetInfo().getTriple().str();
+      targetDiag(D->getLocation(), diag::note_defined_here) << D;
+    }
+  };
+
+  QualType Ty = D->getType();
+  CheckType(Ty);
+
+  if (const auto *FPTy = dyn_cast<FunctionProtoType>(Ty)) {
+    for (const auto &ParamTy : FPTy->param_types())
+      CheckType(ParamTy);
+    CheckType(FPTy->getReturnType());
+  }
+}
+
 /// Looks through the macro-expansion chain for the given
 /// location, looking for a macro expansion with the given name.
 /// If one is found, returns true and sets the location to that
 /// expansion loc.
 bool Sema::findMacroSpelling(SourceLocation &locref, StringRef name) {
   SourceLocation loc = locref;
   if (!loc.isMacroID()) return false;
 

clang/lib/Sema/SemaDecl.cpp

Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
   PopFunctionScopeInfo(ActivePolicy, dcl);
   // If any errors have occurred, clear out any temporaries that may have
   // been leftover. This ensures that these temporaries won't be picked up for
   // deletion in some later function.
   if (getDiagnostics().hasErrorOccurred()) {
     DiscardCleanupsInEvaluationContext();
   }
 
-  if (LangOpts.OpenMP || LangOpts.CUDA) {
+  if (LangOpts.OpenMP || LangOpts.CUDA || LangOpts.SYCLIsDevice) {
     auto ES = getEmissionStatus(FD);
     if (ES == Sema::FunctionEmissionStatus::Emitted ||
         ES == Sema::FunctionEmissionStatus::Unknown)
       DeclsToCheckForDeferredDiags.push_back(FD);
   }
 
   return dcl;
 }

 }
 
 Decl *Sema::getObjCDeclContext() const {
   return (dyn_cast_or_null<ObjCContainerDecl>(CurContext));
 }
 
 Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD,
                                                      bool Final) {
+  // SYCL functions can be template, so we check if they have appropriate
+  // attribute prior to checking if it is a template.
+  if (LangOpts.SYCLIsDevice && FD->hasAttr<SYCLKernelAttr>())
+    return FunctionEmissionStatus::Emitted;
+
   // Templates are emitted when they're instantiated.
   if (FD->isDependentContext())
     return FunctionEmissionStatus::TemplateDiscarded;
 
   FunctionEmissionStatus OMPES = FunctionEmissionStatus::Unknown;
   if (LangOpts.OpenMPIsDevice) {
     Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy =
         OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl());

clang/lib/Sema/SemaDeclCXX.cpp

Sema::BuildCXXConstructExpr(SourceLocation ConstructLoc, QualType DeclInitType,
                             SourceRange ParenRange) {
   assert(declaresSameEntity(
              Constructor->getParent(),
              DeclInitType->getBaseElementTypeUnsafe()->getAsCXXRecordDecl()) &&
          "given constructor for wrong type");
   MarkFunctionReferenced(ConstructLoc, Constructor);
   if (getLangOpts().CUDA && !CheckCUDACall(ConstructLoc, Constructor))
     return ExprError();
+  if (getLangOpts().SYCLIsDevice &&
+      !checkSYCLDeviceFunction(ConstructLoc, Constructor))
+    return ExprError();
 
   return CheckForImmediateInvocation(
       CXXConstructExpr::Create(
           Context, DeclInitType, ConstructLoc, Constructor, Elidable, ExprArgs,
           HadMultipleCandidates, IsListInitialization,
           IsStdInitListInitialization, RequiresZeroInit,
           static_cast<CXXConstructExpr::ConstructionKind>(ConstructKind),
           ParenRange),

clang/lib/Sema/SemaExpr.cpp

if (FunctionDecl *FD = dyn_cast<FunctionDecl>(D)) {
     // If the function has a deduced return type, and we can't deduce it,
     // then we can't use it either.
     if (getLangOpts().CPlusPlus14 && FD->getReturnType()->isUndeducedType() &&
         DeduceReturnType(FD, Loc))
       return true;
 
     if (getLangOpts().CUDA && !CheckCUDACall(Loc, FD))
       return true;
+
+    if (getLangOpts().SYCLIsDevice && !checkSYCLDeviceFunction(Loc, FD))
+      return true;
   }
 
   if (auto *MD = dyn_cast<CXXMethodDecl>(D)) {
     // Lambdas are only default-constructible or assignable in C++2a onwards.
     if (MD->getParent()->isLambda() &&
         ((isa<CXXConstructorDecl>(MD) &&
           cast<CXXConstructorDecl>(MD)->isDefaultConstructor()) ||
          MD->isCopyAssignmentOperator() || MD->isMoveAssignmentOperator())) {
bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
 
   DiagnoseAvailabilityOfDecl(D, Locs, UnknownObjCClass, ObjCPropertyAccess,
                              AvoidPartialAvailabilityChecks, ClassReceiver);
 
   DiagnoseUnusedOfDecl(*this, D, Loc);
 
   diagnoseUseOfInternalDeclInInlineFunction(*this, D, Loc);
 
+  if (LangOpts.SYCLIsDevice || (LangOpts.OpenMP && LangOpts.OpenMPIsDevice))
+    if (const auto *VD = dyn_cast<ValueDecl>(D))
+      checkDeviceDecl(VD, Loc);
+
   if (isa<ParmVarDecl>(D) && isa<RequiresExprBodyDecl>(D->getDeclContext()) &&
       !isUnevaluatedContext()) {
     // C++ [expr.prim.req.nested] p3
     //   A local parameter shall only appear as an unevaluated operand
     //   (Clause 8) within the constraint-expression.
     Diag(Loc, diag::err_requires_expr_parameter_referenced_in_evaluated_context)
         << D;
     Diag(D->getLocation(), diag::note_entity_declared_at) << D;
if (LHSTy->isImageType() || RHSTy->isImageType() ||
         LHSTy->isSamplerT() || RHSTy->isSamplerT() ||
         LHSTy->isPipeType() || RHSTy->isPipeType() ||
         LHSTy->isBlockPointerType() || RHSTy->isBlockPointerType()) {
       ResultTy = InvalidOperands(OpLoc, LHS, RHS);
       return ExprError();
     }
   }
 
-  // Diagnose operations on the unsupported types for OpenMP device compilation.
-  if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice) {
-    if (Opc != BO_Assign && Opc != BO_Comma) {
-      checkOpenMPDeviceExpr(LHSExpr);
-      checkOpenMPDeviceExpr(RHSExpr);
-    }
-  }
-
   switch (Opc) {
   case BO_Assign:
     ResultTy = CheckAssignmentOperands(LHS.get(), RHS, OpLoc, QualType());
     if (getLangOpts().CPlusPlus &&
         LHS.get()->getObjectKind() != OK_ObjCProperty) {
       VK = LHS.get()->getValueKind();
       OK = LHS.get()->getObjectKind();
     }
if (getLangOpts().OpenCL) {
     // only with a builtin functions and therefore should be disallowed here.
         (Ty->isImageType() || Ty->isSamplerT() || Ty->isPipeType()
         || Ty->isBlockPointerType())) {
       return ExprError(Diag(OpLoc, diag::err_typecheck_unary_expr)
                        << InputExpr->getType()
                        << Input.get()->getSourceRange());
     }
   }
-  // Diagnose operations on the unsupported types for OpenMP device compilation.
-  if (getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice) {
-    if (UnaryOperator::isIncrementDecrementOp(Opc) ||
-        UnaryOperator::isArithmeticOp(Opc))
-      checkOpenMPDeviceExpr(InputExpr);
-  }
 
   switch (Opc) {
   case UO_PreInc:
   case UO_PreDec:
   case UO_PostInc:
   case UO_PostDec:
     resultType = CheckIncrementDecrementOperand(*this, Input.get(), VK, OK,
                                                 OpLoc,
void Sema::MarkFunctionReferenced(SourceLocation Loc, FunctionDecl *Func,
   if (NeedDefinition &&
       (Func->getTemplateSpecializationKind() != TSK_Undeclared ||
        Func->getMemberSpecializationInfo()))
     checkSpecializationVisibility(Loc, Func);
 
   if (getLangOpts().CUDA)
     CheckCUDACall(Loc, Func);
 
+  if (getLangOpts().SYCLIsDevice)
+    checkSYCLDeviceFunction(Loc, Func);
+
   // If we need a definition, try to create one.
   if (NeedDefinition && !Func->getBody()) {
     runWithSufficientStackSpace(Loc, [&] {
       if (CXXConstructorDecl *Constructor =
               dyn_cast<CXXConstructorDecl>(Func)) {
         Constructor = cast<CXXConstructorDecl>(Constructor->getFirstDecl());
         if (Constructor->isDefaulted() && !Constructor->isDeleted()) {
           if (Constructor->isDefaultConstructor()) {

clang/lib/Sema/SemaOpenMP.cpp

enum class FunctionEmissionStatus {
   Unknown,
 };
 } // anonymous namespace
 
 Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc,
                                                      unsigned DiagID) {
   assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice &&
          "Expected OpenMP device compilation.");
-  FunctionEmissionStatus FES = getEmissionStatus(getCurFunctionDecl());
+
+  FunctionDecl *FD = getCurFunctionDecl();
   DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop;
+  if (FD) {
+    FunctionEmissionStatus FES = getEmissionStatus(FD);
     switch (FES) {
     case FunctionEmissionStatus::Emitted:
       Kind = DeviceDiagBuilder::K_Immediate;
       break;
     case FunctionEmissionStatus::Unknown:
-    Kind = isOpenMPDeviceDelayedContext(*this) ? DeviceDiagBuilder::K_Deferred
-                                               : DeviceDiagBuilder::K_Immediate;
+      Kind = isOpenMPDeviceDelayedContext(*this)
+                 ? DeviceDiagBuilder::K_Deferred
       break;
     case FunctionEmissionStatus::TemplateDiscarded:
     case FunctionEmissionStatus::OMPDiscarded:
       Kind = DeviceDiagBuilder::K_Nop;
       break;
     case FunctionEmissionStatus::CUDADiscarded:
       llvm_unreachable("CUDADiscarded unexpected in OpenMP device compilation");
       break;
     }
+  }
 
   return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
 Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
                                                    unsigned DiagID) {
   assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice &&
          "Expected OpenMP host compilation.");
Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc,
   case FunctionEmissionStatus::CUDADiscarded:
     Kind = DeviceDiagBuilder::K_Nop;
     break;
   }
 
   return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this);
 }
 
-void Sema::checkOpenMPDeviceExpr(const Expr *E) {
-  assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice &&
-         "OpenMP device compilation mode is expected.");
-  QualType Ty = E->getType();
-  if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
-      ((Ty->isFloat128Type() ||
-        (Ty->isRealFloatingType() && Context.getTypeSize(Ty) == 128)) &&
-       !Context.getTargetInfo().hasFloat128Type()) ||
-      (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
-       !Context.getTargetInfo().hasInt128Type()))
-    targetDiag(E->getExprLoc(), diag::err_omp_unsupported_type)
-        << static_cast<unsigned>(Context.getTypeSize(Ty)) << Ty
-        << Context.getTargetInfo().getTriple().str() << E->getSourceRange();
-}
-
 static OpenMPDefaultmapClauseKind
 getVariableCategoryFromDecl(const LangOptions &LO, const ValueDecl *VD) {
   if (LO.OpenMP <= 45) {
     if (VD->getType().getNonReferenceType()->isScalarType())
       return OMPC_DEFAULTMAP_scalar;
     return OMPC_DEFAULTMAP_aggregate;
   }
   if (VD->getType().getNonReferenceType()->isAnyPointerType())

clang/lib/Sema/SemaSYCL.cpp

+//===- SemaSYCL.cpp - Semantic Analysis for SYCL constructs ---------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+// This implements Semantic Analysis for SYCL constructs.
+//===----------------------------------------------------------------------===//
+
+#include "clang/Sema/Sema.h"
+#include "clang/Sema/SemaDiagnostic.h"
+
+using namespace clang;
+
+// -----------------------------------------------------------------------------
+// SYCL device specific diagnostics implementation
+// -----------------------------------------------------------------------------
+
+Sema::DeviceDiagBuilder Sema::SYCLDiagIfDeviceCode(SourceLocation Loc,
+                                                   unsigned DiagID) {
+  assert(getLangOpts().SYCLIsDevice &&
+         "Should only be called during SYCL compilation");
+  FunctionDecl *FD = dyn_cast<FunctionDecl>(getCurLexicalContext());
+  DeviceDiagBuilder::Kind DiagKind = [this, FD] {
+    if (!FD)
+      return DeviceDiagBuilder::K_Nop;
+    if (getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted)
+      return DeviceDiagBuilder::K_ImmediateWithCallStack;
+    return DeviceDiagBuilder::K_Deferred;
+  }();
+  return DeviceDiagBuilder(DiagKind, Loc, DiagID, FD, *this);
+}
+
+bool Sema::checkSYCLDeviceFunction(SourceLocation Loc, FunctionDecl *Callee) {
+  assert(getLangOpts().SYCLIsDevice &&
+         "Should only be called during SYCL compilation");
+  assert(Callee && "Callee may not be null.");
+
+  // Errors in unevaluated context don't need to be generated,
+  // so we can safely skip them.
+  if (isUnevaluatedContext() || isConstantEvaluated())
+    return true;
+
+  DeviceDiagBuilder::Kind DiagKind = DeviceDiagBuilder::K_Nop;
+
+  return DiagKind != DeviceDiagBuilder::K_Immediate &&
+         DiagKind != DeviceDiagBuilder::K_ImmediateWithCallStack;
+}

clang/lib/Sema/SemaType.cpp

static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
   case DeclSpec::TST_double:
     if (DS.getTypeSpecWidth() == DeclSpec::TSW_long)
       Result = Context.LongDoubleTy;
     else
       Result = Context.DoubleTy;
     break;
   case DeclSpec::TST_float128:
     if (!S.Context.getTargetInfo().hasFloat128Type() &&
+        !S.getLangOpts().SYCLIsDevice &&
         !(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
       S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
         << "__float128";
     Result = Context.Float128Ty;
     break;
   case DeclSpec::TST_bool: Result = Context.BoolTy; break; // _Bool or bool
     break;
   case DeclSpec::TST_decimal32:    // _Decimal32

clang/test/Headers/nvptx_device_math_sin.c

 // REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -x c -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=SLOW
 // RUN: %clang_cc1 -x c -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast
 // RUN: %clang_cc1 -x c -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=FAST
 // expected-no-diagnostics
 
 #include <math.h>
 
-double math(float f, double d, long double ld) {
+double math(float f, double d) {
   double r = 0;
 // SLOW:  call float @__nv_sinf(float
 // FAST:  call fast float @__nv_fast_sinf(float
   r += sinf(f);
 // SLOW:  call double @__nv_sin(double
 // FAST:  call fast double @__nv_sin(double
   r += sin(d);
   return r;
 }
 
 long double foo(float f, double d, long double ld) {
   double r = ld;
-  r += math(f, d, ld);
+  r += math(f, d);
 #pragma omp target map(r)
-  { r += math(f, d, ld); }
+  { r += math(f, d); }
   return r;
 }

clang/test/Headers/nvptx_device_math_sin.cpp

 // REQUIRES: nvptx-registered-target
 // RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
 // RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix=SLOW
 // RUN: %clang_cc1 -x c++ -internal-isystem %S/Inputs/include -fopenmp -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc -ffast-math -ffp-contract=fast
 // RUN: %clang_cc1 -x c++ -include __clang_openmp_device_functions.h -internal-isystem %S/../../lib/Headers/openmp_wrappers -internal-isystem %S/Inputs/include -fopenmp -triple nvptx64-nvidia-cuda -aux-triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -ffast-math -ffp-contract=fast | FileCheck %s --check-prefix=FAST
 // expected-no-diagnostics
 
 #include <cmath>
 
-double math(float f, double d, long double ld) {
+double math(float f, double d) {
   double r = 0;
 // SLOW:  call float @__nv_sinf(float
 // FAST:  call fast float @__nv_fast_sinf(float
   r += sin(f);
 // SLOW:  call double @__nv_sin(double
 // FAST:  call fast double @__nv_sin(double
   r += sin(d);
   return r;
 }
 
 long double foo(float f, double d, long double ld) {
   double r = ld;
-  r += math(f, d, ld);
+  r += math(f, d);
 #pragma omp target map(r)
-  { r += math(f, d, ld); }
+  { r += math(f, d); }
   return r;
 }

clang/test/OpenMP/nvptx_unsupported_type_codegen.cpp

 #else
 typedef long double BIGTYPE;
 #endif
 
 struct T {
   char a;
   BIGTYPE f;
   char c;
   T() : a(12), f(15) {}
   T &operator+(T &b) { f += b.a; return *this;}
 };
 
 struct T1 {
   char a;
   __int128 f;
   __int128 f1;
   char c;

 }
 // CHECK: define{{ hidden | }}void @{{.+}}baz1{{.+}}()
 void baz1() {
 // CHECK: call [[T1]] @{{.+}}bar1{{.+}}()
   T1 t = bar1();
 }
 #pragma omp end declare target
 
-BIGTYPE foo(BIGTYPE f) {
-#pragma omp target map(f)
-  f = 1;
-  return f;
-}
-
-// CHECK: define weak void @__omp_offloading_{{.+}}foo{{.+}}_l75([[BIGTYPE:.+]]*
-// CHECK: store [[BIGTYPE]] {{0xL00000000000000003FFF000000000000|0xM3FF00000000000000000000000000000}}, [[BIGTYPE]]* %

clang/test/OpenMP/nvptx_unsupported_type_messages.cpp

 // Test target codegen - host bc file has to be created first.
 // RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple x86_64-unknown-linux -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
 // RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-host.bc
 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -aux-triple powerpc64le-unknown-linux-gnu -fopenmp-targets=nvptx64-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-host.bc -fsyntax-only
 
 struct T {
   char a;
 #ifndef _ARCH_PPC
+  // expected-note@+1 {{'f' defined here}}
   __float128 f;
 #else
+  // expected-note@+1 {{'f' defined here}}
   long double f;
 #endif
   char c;
   T() : a(12), f(15) {}
 #ifndef _ARCH_PPC
-// expected-error@+4 {{host requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+5 {{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
 #else
-// expected-error@+2 {{host requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+// expected-error@+3 {{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
 #endif
-  T &operator+(T &b) { f += b.a; return *this;}
+  T &operator+(T &b) {
+    f += b.a;
+    return *this;
+  }
 };
 
 struct T1 {
   char a;
   __int128 f;
   __int128 f1;
   char c;
   T1() : a(12), f(15) {}
-  T1 &operator/(T1 &b) { f /= b.a; return *this;}
+  T1 &operator/(T1 &b) {
+    f /= b.a;
+    return *this;
+  }
 };
 
+#ifndef _ARCH_PPC
+// expected-note@+1 {{'boo' defined here}}
+void boo(__float128 A) { return; }
+#else
+// expected-note@+1 {{'boo' defined here}}
+void boo(long double A) { return; }
+#endif
 #pragma omp declare target
 T a = T();
 T f = a;
 void foo(T a = T()) {
   a = a + f; // expected-note {{called by 'foo'}}
+#ifndef _ARCH_PPC
+// expected-error@+4 {{'boo' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+#else
+// expected-error@+2 {{'boo' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+#endif
+  boo(0);
   return;
 }
 T bar() {
   return T();
 }
+
 void baz() {
   T t = bar();
 }
 T1 a1 = T1();
 T1 f1 = a1;
 void foo1(T1 a = T1()) {
   a = a / f1;
   return;
 }
 T1 bar1() {
   return T1();
 }
 void baz1() {
   T1 t = bar1();
 }
 #pragma omp end declare target
+
+#ifndef _ARCH_PPC
+// expected-note@+1 3{{'f' defined here}}
+__float128 foo1(__float128 f) {
+#pragma omp target map(f)
+  // expected-error@+1 3{{'f' requires 128 bit size '__float128' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+  f = 1;
+  return f;
+}
+#else
+// expected-note@+1 3{{'f' defined here}}
+long double foo1(long double f) {
+#pragma omp target map(f)
+  // expected-error@+1 3{{'f' requires 128 bit size 'long double' type support, but device 'nvptx64-unknown-unknown' does not support it}}
+  f = 1;
+  return f;
+}
+#endif
+
+T foo3() {
+  T S;
+#pragma omp target map(S)
+  S.a = 1;
+  return S;
+}
+
+// Allow all sorts of stuff on host
+#ifndef _ARCH_PPC
+__float128 q, b;
+__float128 c = q + b;
+#else
+long double q, b;
+long double c = q + b;
+#endif
+
+void hostFoo() {
+  boo(c - b);
+}
+
+long double qa, qb;
+decltype(qa + qb) qc;
+double qd[sizeof(-(-(qc * 2)))];

clang/test/SemaSYCL/float128.cpp

+// RUN: %clang_cc1 -triple spir64 -fsycl -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -fsycl -fsycl-is-device -fsyntax-only %s
+
+typedef __float128 BIGTY;
+
+template <class T>
+class Z {
+public:
+  // expected-note@+1 {{'field' defined here}}
+  T field;
+  // expected-note@+1 2{{'field1' defined here}}
+  __float128 field1;
+  using BIGTYPE = __float128;
+  // expected-note@+1 {{'bigfield' defined here}}
+  BIGTYPE bigfield;
+};
+
+void host_ok(void) {
+  __float128 A;
+  int B = sizeof(__float128);
+  Z<__float128> C;
+  C.field1 = A;
+}
+
+void usage() {
+  // expected-note@+1 3{{'A' defined here}}
+  __float128 A;
+  Z<__float128> C;
+  // expected-error@+2 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+  // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+  C.field1 = A;
+  // expected-error@+1 {{'bigfield' requires 128 bit size 'Z::BIGTYPE' (aka '__float128') type support, but device 'spir64' does not support it}}
+  C.bigfield += 1.0;
+
+  // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+  auto foo1 = [=]() {
+    __float128 AA;
+    // expected-note@+2 {{'BB' defined here}}
+    // expected-error@+1 {{'A' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    auto BB = A;
+    // expected-error@+1 {{'BB' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    BB += 1;
+  };
+
+  // expected-note@+1 {{called by 'usage'}}
+  foo1();
+}
+
+template <typename t>
+void foo2(){};
+
+// expected-note@+3 {{'P' defined here}}
+// expected-error@+2 {{'P' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+// expected-note@+1 2{{'foo' defined here}}
+__float128 foo(__float128 P) { return P; }
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
+  // expected-note@+1 5{{called by 'kernel}}
+  kernelFunc();
+}
+
+int main() {
+  // expected-note@+1 {{'CapturedToDevice' defined here}}
+  __float128 CapturedToDevice = 1;
+  host_ok();
+  kernel<class variables>([=]() {
+    decltype(CapturedToDevice) D;
+    // expected-error@+1 {{'CapturedToDevice' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    auto C = CapturedToDevice;
+    Z<__float128> S;
+    // expected-error@+1 {{'field1' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    S.field1 += 1;
+    // expected-error@+1 {{'field' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    S.field = 1;
+  });
+
+  kernel<class functions>([=]() {
+    // expected-note@+1 2{{called by 'operator()'}}
+    usage();
+    // expected-note@+1 {{'BBBB' defined here}}
+    BIGTY BBBB;
+    // expected-note@+3 {{called by 'operator()'}}
+    // expected-error@+2 2{{'foo' requires 128 bit size '__float128' type support, but device 'spir64' does not support it}}
+    // expected-error@+1 {{'BBBB' requires 128 bit size 'BIGTY' (aka '__float128') type support, but device 'spir64' does not support it}}
+    auto A = foo(BBBB);
+  });
+
+  kernel<class ok>([=]() {
+    Z<__float128> S;
+    foo2<__float128>();
+    auto A = sizeof(CapturedToDevice);
+  });
+
+  return 0;
+}