[PATCH] D62584: [OpenCL][PR42033] Deducing addr space with template parameter types

[Anastasia Stulova via Phabricator via cfe-commits]

Anastasia marked an inline comment as done.
Anastasia added inline comments.


================
Comment at: lib/Sema/TreeTransform.h:5363
+    if (ResultType.getAddressSpace() != LangAS::Default &&
+        (ResultType.getAddressSpace() != LangAS::opencl_private)) {
       SemaRef.Diag(TL.getReturnLoc().getBeginLoc(),
----------------
rjmccall wrote:
> Anastasia wrote:
> > Anastasia wrote:
> > > rjmccall wrote:
> > > > Anastasia wrote:
> > > > > 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])
> > > > > }
> > > > > ```
> > > > 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.
> > > 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?
> > Ok, I have removed the diagnostic completely. At least we don't seem to generate any incorrect IR.
> 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.
> 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?


CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D62584/new/

https://reviews.llvm.org/D62584