Thanks for submitting more broken test cases! :)

The fix is in review: https://reviews.llvm.org/D62584

the problem in this test case is fixed, but I see more similar problems. Investigating and creating a new test case...

ok, that was easy, new testcase below, which still produces the failure also with the patch:

template <class T>
void xxx(T* left, T* right)
{
    T* min = left;
    T min2 = *left;
}

__kernel void test()
{
    int foo[10];
    xxx(&foo[0], &foo[9]);
}

The following similar example fails, but with a different error.

template <class T>
T xxx(T* in) {
  T* i = in; // works with the patch
  return *i; // fails
}

__kernel void test(__global int* out) {
  int foo[1] = { *out };
  int x = xxx(&foo[0]);
}

The error is:

../../test/addr_space/template.cl:9:11: error: no matching function for call to 'xxx'
  int x = xxx(&foo[0]);
          ^~~
../../test/addr_space/template.cl:2:3: note: candidate template ignored: substitution failure [with T = int]: function type may not be qualified with an address space
T xxx(T* in) {
~ ^
1 error generated.


When `T xxx(T* in)` is replaced with `int xxx(T* in)` the program compiles.

Ok, I have now updated the review to fix extra test cases.

I confirm that the current test cases are fixed.
Then, I am kind of sorry, but here is another broken test case:

class s
{
 public:
  void test () {
    swap(&foo[0], &foo[1]);
  }

 private:
  int foo[10];
  template <class T> static void swap(T* v1, T* v2);
};

template <class T> void s::swap(T* v1, T* v2)
{
  T tmp = *v1;
  *v1 = *v2;
  *v2 = tmp;
}

__kernel void test()
{
  s obj;
  obj.test();
}

Thanks for providing extra test cases. The last one doesn't seem to be a bug in clang considering our current approach though. I have to say however the documentation is missing currently to allow understand this case deeper. Let me try to explain the concept here but my aim is to provide more elaborate documentation in clang's manual as soon as possible.

In OpenCL every method is implicitly qualified by __generic that means every object of a class is considered to be in generic address space inside the method. Therefore your method 'test' is equivalent to

  void test () {
    swap(&(this->foo[0]), &(this->foo[1]));
  }

Because the type of 'this' is '__generic s*' the template argument of 'swap' is deduced to be 'generic int'. During the template instantiation of 'swap' with an argument 'generic int' we end up declaring automatic variable in generic address space that is not allowed in OpenCL. From OpenCL C v2.0 s6.5: "The address space name for arguments to a function in a program, or local variables of a function is __private."

Does it clarify the issue for you?

There are multiple ways to solve this problem in the current clang.
1. You can qualify the method 'test' by '__private' (then the type of 'this' will become '__private s*').

  void test __private () {
    swap(&foo[0], &foo[1]);
  }

2. You can explicitly add the template parameter. 

  void test () {
    swap<__private int>((__private int*)&foo[0], (__private int*)&foo[1]);
  }

Both are not ideal because you might need to use the same trick in multiple places. In the future we plan to mitigate this range of problems by adding a metaprogramming support that would allow to remove address space qualifier from the type or by allowing to template method qualifiers on address spaces. However, this is not something that is planned for clang9 unfortunately. :(

Thx for the explanation.
I understand the point, but for the normal user this is not really apparent.
Anyhow, I think I will be able to adapt our code accordingly.

From my side, we can close this bug report as soon as the patch in review is merged.

Certainly, a better solution would be useful, also if it comes in clang>9. In general, this entire __generic address space business in OpenCL is a bit clumsy. Handling in CUDA and HIP is much simpler, and I am not so convinced the approach taken by OpenCL is the ideal one. But this is a different story.

Yes, I see your point. It is generally a really valuable input. We will try to make it as usable as possible. One thing to point out is that CUDA/HIP is originally intended for one family of architectures and hence they are able to simply some concepts. OpenCL has to be more generic in terms of architectures supported. We also have other requirements such as support of non optimized mode.

At some point I would like to follow up with you more on your feedback and comparison. Perhaps it will be better to do offline. :)

@Anastasia: FYI, I think I actually found a better solution:

template <class T> void s::swap(__generic T* v1, __generic T* v2)
{
  T tmp = *v1;
  *v1 = *v2;
  *v2 = tmp;
}

deduces the correct template parameter, and is not too much extra typing word.

Cool. Indeed this seems like a nice solution for your case.

However, I think generically providing some metaprograming support would help in the future.

Fixed in r366417