Page MenuHomePhabricator

[OpenCL][PR42033] Deducing addr space with template parameter types
ClosedPublic

Authored by Anastasia on May 29 2019, 6:32 AM.

Details

Summary

If dependent types appear in pointers or references we have to allow the addr space deduction because the addr space in template argument will belong to the pointee and not the pointer or reference itself. Hence we end up with un-deduced i.e. Default address space.

We also need to change somke diagnostic to account for addr space logic with templates.

Diff Detail

Repository
rL LLVM

Event Timeline

Anastasia created this revision.May 29 2019, 6:32 AM

I think the right approach here is probably to make sure you're applying deduction during instantiation as well.

Anastasia updated this revision to Diff 202753.Jun 3 2019, 10:50 AM
Anastasia retitled this revision from [OpenCL][PR42033] Deducing addr space of pointer/reference with template parameter types to [OpenCL][PR42033] Deducing addr space with template parameter types.
Anastasia edited the summary of this revision. (Show Details)
  • Added more test cases
  • Improved diagnostics to account for use of templates with addr spaces

I think the right approach here is probably to make sure you're applying deduction during instantiation as well.

I agree I think we might need to extend the template instantiation logic to account for some corner cases. However, all local variables in OpenCL are to be deduced to __private, therefore would it be better to deduce them already as we parse the template definition instead of doing it multiple times on each instantiation? It doesn't seem the template argument should affect this rule...

Anastasia marked an inline comment as done.Jun 3 2019, 11:01 AM
Anastasia added inline comments.
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)

I am trying to be a bit more helpful here although I am not sure if we should instead require explicit template parameter and fail the template deduction instead.

Basically, do we want the following code to always require specifying template argument explicitly:

template <class T>
T xxx(T *in) {
  T *i = in;
  return *i;
}

__kernel void test() {
  int foo[10];
  xxx<int>(&foo[0]); // if we deduce type from foo, it ends up being qualified by __private that we currently diagnose. However private is default (implicit) address space for return type so technically there is no danger in just allowing xxx(&foo[0])
}
rjmccall added inline comments.Jun 10 2019, 3:47 PM
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)

Implicitly ignoring all address-space qualifiers on the return type seems like the right thing to do; I don't think it needs to be limited to __private. That's probably also the right thing to do for locals, but I'm less certain about it.

Anastasia marked an inline comment as done.Jun 21 2019, 7:29 AM
Anastasia added inline comments.
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)

Just to clarify by "Implicitly ignoring" you mean ignoring if the templ parameters were deduced?

Although I am a bit concerned about allowing other than __private address spaces in return types as we reject them in return types of functions generally. Would it not be somehow inconsistent?

Anastasia updated this revision to Diff 207562.Jul 2 2019, 7:44 AM
  • Removed diagnostic for address space in return type in tree transforms
Anastasia marked an inline comment as done.Jul 2 2019, 7:45 AM
Anastasia added inline comments.
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)

Ok, I have removed the diagnostic completely. At least we don't seem to generate any incorrect IR.

rjmccall added inline comments.Jul 9 2019, 11:04 PM
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)

They should be diagnosed somehow when written explicitly on a return type, but if you just do that on the parser path you don't have to worry about it during template instantiation. They should probably otherwise be ignored no matter where they came from — if someone typedefs private_int_t to __private int, you should just treat that as int in a return type. Stripping the qualifier from the type is probably the right thing to do so that it doesn't further impact semantic analysis.

I definitely don't think you want a model where the qualifier actually means that the return is somehow done via an object in that address space.

Anastasia marked an inline comment as done.Jul 10 2019, 11:42 AM
Anastasia added inline comments.
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)

They should be diagnosed somehow when written explicitly on a return type, but if you just do that on the parser path you don't have to worry about it during template instantiation.

Ok, this seems to be working currently. The following won't compile:

template <typename T>
struct S {
  __global T f1();     // error: return value cannot be qualified with address space
};

They should probably otherwise be ignored no matter where they came from — if someone typedefs private_int_t to __private int, you should just treat that as int in a return type.

We produce diagnostic for this case currently. I can update this in a separate patch if you think it's more helpful behavior. Although I feel a bit worried about it. However it would align with what we are doing with templates here... so perhaps it's logical change...

Stripping the qualifier from the type is probably the right thing to do so that it doesn't further impact semantic analysis.

I guess you mean stripping quals for the case of typedef or others where we will accept the code and ignore quals on the return type? I have tried to do this just for the template case but there are some assertions firing because we are expecting the original return type to be the same before and after template instantiation. So it feels like I would have to do it for everything together. Maybe it should rather go into separate review?

rjmccall added inline comments.Jul 11 2019, 7:41 PM
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)

We produce diagnostic for this case currently. I can update this in a separate patch if you think it's more helpful behavior. Although I feel a bit worried about it. However it would align with what we are doing with templates here... so perhaps it's logical change...

Well, it's debatable. C doesn't say that qualifiers are erased in function return types in general. On the other hand, qualifiers are semantically meaningless there, and C has few rules that rely on exact type identity. On reflection, we should probably keep diagnosing that at parse-time, even for typedefs, and just ignore it in templates.

Can you point out which assertion is firing?

Anastasia marked an inline comment as done.Jul 12 2019, 8:05 AM
Anastasia added inline comments.
lib/Sema/TreeTransform.h
5363 ↗(On Diff #202753)
On reflection, we should probably keep diagnosing that at parse-time, even for typedefs, and just ignore it in templates.

In that case we don't need to fix anything in parsing.

Can you point out which assertion is firing?

I have tried this:

diff --git a/lib/Sema/TreeTransform.h b/lib/Sema/TreeTransform.h
index 878795bb70..db4ae6e56e 100644
--- a/lib/Sema/TreeTransform.h
+++ b/lib/Sema/TreeTransform.h
@@ -5365,6 +5365,11 @@ QualType TreeTransform<Derived>::TransformFunctionProtoType(
     if (ResultType.isNull())
       return QualType();
 
+    if (ResultType.hasQualifiers()) {
+      ResultType = ResultType.getUnqualifiedType();
+      TLB.TypeWasModifiedSafely(ResultType);
+    }
+
     if (getDerived().TransformFunctionTypeParams(
             TL.getBeginLoc(), TL.getParams(),
             TL.getTypePtr()->param_type_begin(),

and I am getting some errors in tests mainly because of the differences in the return type. I.e.

error: 'error' diagnostics expected but not seen: 
  File llvm/tools/clang/test/CXX/dcl.decl/dcl.init/dcl.init.ref/p5-var.cpp Line 126: binding reference of type 'const Base' to value of type 'const volatile Base' drops 'volatile' qualifier
  File llvm/tools/clang/test/CXX/dcl.decl/dcl.init/dcl.init.ref/p5-var.cpp Line 127: binding reference of type 'const Base' to value of type 'const volatile Derived' drops 'volatile' qualifier

Practically we no longer get diagnostics for cases like:

48 template<typename T> T create();

126  const Base &br5 = create<const volatile Base>();
127  const Base &br6 = create<const volatile Derived>();

Also the qualifiers are not printed in the type of return type.

I guess this is expected. I am however struggling with this assert in a couple tests:

llvm/tools/clang/lib/Sema/TypeLocBuilder.cpp:160: clang::TypeLoc clang::TypeLocBuilder::pushImpl(clang::QualType, size_t, unsigned int): Assertion `Capacity - Index == TypeLoc::getFullDataSizeForType(T) && "incorrect data size provided to CreateTypeSourceInfo!"' failed.

I am not sure if I am modifying TypeLocBuilder in a wrong way or perhaps something isn't working well in the Index calculation. If this is the direction we would like to go I can spend more time debugging this and upload the review. Not sure if that might be too disruptive change for the release but I would at least like to see the current patch in Clang9 because it fixes a couple of useful cases. :)

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

Ok, regarding address space qualifiers in template instantiation - is it still ok that we ignore them here in this patch?

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

Ok, regarding address space qualifiers in template instantiation - is it still ok that we ignore them here in this patch?

I think ignoring them in templates rather than diagnosing is the right thing to do, since it means that reasonable things, e.g. template functions returning T, won't be erroneous if you instantiate T = __global int just like they wouldn't be with T = const int.

Oh, yes, it definitely can't be done to class types. I suppose we should just forget about it.

Ok, regarding address space qualifiers in template instantiation - is it still ok that we ignore them here in this patch?

I think ignoring them in templates rather than diagnosing is the right thing to do, since it means that reasonable things, e.g. template functions returning T, won't be erroneous if you instantiate T = __global int just like they wouldn't be with T = const int.

Ok cool, I have removed the diagnostic btw. :)

Minor comment then LGTM

lib/Sema/SemaType.cpp
7418 ↗(On Diff #207562)

"Except"

Anastasia updated this revision to Diff 210277.Jul 17 2019, 2:42 AM

Fixed typo in the comment

rjmccall accepted this revision.Jul 17 2019, 9:30 AM

Thanks!

This revision is now accepted and ready to land.Jul 17 2019, 9:30 AM
This revision was automatically updated to reflect the committed changes.
Herald added a project: Restricted Project. · View Herald TranscriptJul 18 2019, 2:13 AM