[all-commits] [llvm/llvm-project] c77a40: [CUDA][HIP] Fix host/device based overload resolution

Yaxun (Sam) Liu via All-commits all-commits at lists.llvm.org
Fri Apr 24 11:56:11 PDT 2020


  Branch: refs/heads/master
  Home:   https://github.com/llvm/llvm-project
  Commit: c77a4078e01033aa2206c31a579d217c8a07569b
      https://github.com/llvm/llvm-project/commit/c77a4078e01033aa2206c31a579d217c8a07569b
  Author: Yaxun (Sam) Liu <yaxun.liu at amd.com>
  Date:   2020-04-24 (Fri, 24 Apr 2020)

  Changed paths:
    M clang/lib/Sema/SemaOverload.cpp
    M clang/test/SemaCUDA/function-overload.cu

  Log Message:
  -----------
  [CUDA][HIP] Fix host/device based overload resolution

Currently clang fails to compile the following CUDA program in device compilation:

__host__ int foo(int x) {
     return 1;
}

template<class T>
__device__ __host__ int foo(T x) {
    return 2;
}

__device__ __host__ int bar() {
    return foo(1);
}

__global__ void test(int *a) {
    *a = bar();
}

This is due to foo is resolved to the __host__ foo instead of __device__ __host__ foo.
This seems to be a bug since __device__ __host__ foo is a viable callee for foo whereas
clang is unable to choose it.

This patch fixes that.

Differential Revision: https://reviews.llvm.org/D77954




More information about the All-commits mailing list