Page MenuHomePhabricator

[OpenCL] Enable address spaces for references in C++

Authored by Anastasia on Oct 26 2018, 7:55 AM.



I first enabled AS deduction for references that allowed to inherit the right conversion diagnostics based on qualification conversion rules implemented earlier for the pointer type.

Then in order to tests the deduction rules fully, I had to enable some extra features from OpenCL 2.0 that are also valid in C++.

A number of ICEs fired in the CodeGen due to missing addrspacecast. Not convinced the current solution is good though. May be it would be cleaner to add a separate CastKind here - CK_LValueAddressSpaceConversion? Although I am not entirely clear about the benefits yet.

Diff Detail


Event Timeline

Anastasia created this revision.Oct 26 2018, 7:55 AM
Anastasia added inline comments.
1609 ↗(On Diff #171306)

I don't like this assert now. Would adding extra variable be cleaner here?

rjmccall added inline comments.Oct 26 2018, 9:19 PM
1609 ↗(On Diff #171306)

Yeah, this assertion doesn't make any sense like this. It should be checking whether the cast is a gl-value and, if so, requiring the subexpression to also be a gl-value and then asserting the difference between the type. But you can certainly do an address-space conversion on l-values that just happen to be of pointer or block-pointer type.

4252 ↗(On Diff #171306)

Please use the performAddrSpaceCast target hook instead of directly constructing an LLVM addrspacecast.

576 ↗(On Diff #171306)

Please update the comment above this.

7366 ↗(On Diff #171306)

Please update the comment above this.

7614 ↗(On Diff #171306)

Please extract a function to do an l-value qualification conversion just in case we add more non-trivial conversions that we need to represent.

Anastasia updated this revision to Diff 172109.Nov 1 2018, 5:13 AM

Addressed comments from John.

Anastasia marked 4 inline comments as done.Nov 1 2018, 5:15 AM
rjmccall added inline comments.Nov 1 2018, 9:34 AM
1609 ↗(On Diff #171306)

No, if this is a gl-value cast, the assertion must ignore whether there's a pointee type, or it will be messed up on gl-values of pointer types.

That is, if I have a gl-value of type char * __private, I should be able to do an address-space promotion to get a gl-value of type char * __generic. It's okay that the pointers are into the same address space here — in fact, it's more than okay, it's necessary.

Anastasia updated this revision to Diff 173334.Nov 9 2018, 7:41 AM

Changed the assert for address space conversion.

rjmccall added inline comments.Nov 9 2018, 9:42 AM
1609 ↗(On Diff #171306)

Thanks, that's right now. Although please assert that the base has the same value kind; I've seen bugs before where ICEs tried to implicitly materialize their arguments, and it's really frustrating to root out.

4285 ↗(On Diff #173334)

If ToType is a reference type, the address space will be on its pointee type.

Anastasia updated this revision to Diff 173693.Nov 12 2018, 9:23 AM
Anastasia marked an inline comment as done.
  • Extended assert
    • Handled AS of ToType
rjmccall added inline comments.Nov 12 2018, 11:00 AM
4289 ↗(On Diff #173693)

Okay. But if ToType *isn't* a reference type, this will never be an address-space conversion. I feel like this code could be written more clearly to express what it's trying to do.

Anastasia updated this revision to Diff 173873.Nov 13 2018, 9:46 AM

Rewrite how CastKind is set for reference and pointer type.

Anastasia added inline comments.Nov 13 2018, 9:48 AM
4289 ↗(On Diff #173693)

I hope it makes more sense now. Btw, it also applies to pointer type.

rjmccall added inline comments.Nov 13 2018, 9:57 AM
4289 ↗(On Diff #173693)

The logic is wrong for pointer types; if you're converting pointers, you need to be checking the address space of the pointee type of the from type.

It sounds like this is totally inadequately tested; please flesh out the test with all of these cases. While you're at it, please ensure that there are tests verifying that we don't allowing address-space changes in nested positions.

Anastasia updated this revision to Diff 174033.Nov 14 2018, 7:00 AM

Fixed check for AS mismatch of pointer type and added missing test case

Anastasia added inline comments.Nov 14 2018, 7:28 AM
4289 ↗(On Diff #173693)

Thanks for spotting this bug! The generated IR for the test was still correct because AS of FromType happened to correctly mismatch AS of pointee of ToType.

I failed to construct the test case where it would miss classifying addrspacecast due to OpenCL or C++ sema rules but I managed to add a case in which addrspacecast was incorrectly added for pointers where it wasn't needed (see line 36 of the test). I think this code is covered now.

As for the address space position in pointers, the following test checks the address spaces of pointers in addrspacecast. For the other program paths we also have a test with similar checks in test/CodeGenOpenCL/ that we now run for C++ mode too.

BTW, while trying to construct a test case for the bug, I have discovered that multiple pointer indirection casting isn't working correctly. I.e. for the following program:

kernel void foo(){
   __private int** loc;
   int** loc_p = loc;
   **loc_p = 1;

We generate:

bitcast i32* addrspace(4)* %0 to i32 addrspace(4)* addrspace(4)*

in OpenCL C and then perform store over pointer in AS 4 (generic). We have now lost the information that the original pointer was in private AS and that the adjustment of AS segment has to be performed before accessing memory pointed by the pointer. Based on the current specification of addrspacecast in I am not very clear whether it can be used for this case without any modifications or clarifications and also what would happen if there are multiple AS mismatches. I am going to look at this issue separately in more details. In OpenCL C++ an ICE is triggered for this though. Let me know if you have any thoughts on this.

rjmccall added inline comments.Nov 14 2018, 9:57 AM
4289 ↗(On Diff #173693)

Thanks, the check looks good now.

BTW, while trying to construct a test case for the bug, I have discovered that multiple pointer indirection casting isn't working correctly.

This needs to be an error in Sema. The only qualification conversions that should be allowed in general on nested pointers (i.e. on T in T** or T*&) are the basic C qualifiers: const, volatile, and restrict; any other qualification change there is unsound.

Anastasia added inline comments.Nov 14 2018, 10:45 AM
4289 ↗(On Diff #173693)

I see. I guess it's because C++ rules don't cover address spaces.

It feels like it would be a regression for OpenCL C++ vs OpenCL C to reject nested pointers with address spaces because it was allowed before. :(

However, the generation for OpenCL C and C are incorrect currently. I will try to sort that all out as a separate patch though, if it makes sense?

rjmccall added inline comments.Nov 14 2018, 12:12 PM
4289 ↗(On Diff #173693)

C++'s rules assume that qualifiers don't introduce real representation differences and that operations on qualified types are compatible with operations on unqualified types. That's not true of qualifiers in general: address space qualifiers can change representations, ARC qualifiers can have incompatible semantics, etc. There is no way to soundly implement a conversion from __private int ** to __generic int **, just there's no way to soundly implement a conversion from Derived ** to Base **.

If you want to allow this conversion anyway for source-compatibility reasons (and I don't think that's a good idea), it should be a bitcast.

Anastasia marked 5 inline comments as done.Nov 15 2018, 2:24 AM
Anastasia marked 4 inline comments as done.Nov 15 2018, 2:44 AM

Do you think there is anything else to do for this patch?

4289 ↗(On Diff #173693)

Ok, then bitcast is not a good solution because it has an issue of loosing address space information. Perhaps disallowing it completely is a better approach in this case. I have created a bug to investigate it further and may be request some feedback from other OpenCL developers:

rjmccall accepted this revision.Nov 15 2018, 12:58 PM

Thanks, LGTM.

This revision is now accepted and ready to land.Nov 15 2018, 12:58 PM
This revision was automatically updated to reflect the committed changes.
romanovvlad added inline comments.


It seems this code doesn't work correctly(repro at the end). TBAA information is lost here because MakeNaturalAlignPointeeAddrLValue constructs LValue with alignment of poinee type but TBAA info is taken from pointer itself what is strange enough. As a result, for example, memcpy with wrong size is generated for copy constructors.


class P {
  P(const P &Rhs) = default;

  long a;
  long b;

__kernel void foo(__global P* GPtr) {
  P Val = GPtr[0];

As a solution the line could be replaced with the following:

return MakeAddrLValue(Address(V, LV.getAddress().getAlignment()),
                                       E->getType(), LV.getBaseInfo(),
                                       CGM.getTBAAInfoForSubobject(LV, E->getType()));

To take all the information from the original pointer.

What do you think about solution?

rjmccall added inline comments.Nov 29 2018, 10:06 AM

Oh, yes, this should absolutely not be using MakeNaturalAlignPointerAddrLValue; it should be preserving all of the extra information from the original l-value, as you say.

I think TBAA information is independent of address-space qualification and can just be taken from the original LV directly instead of using getTBAAInfoForSubobject.