[PATCH] D68031: [CUDA][HIP] Enable kernel function return type deduction.

[Michael Liao via Phabricator via cfe-commits]

hliao added a comment.

In D68031#1682822 <https://reviews.llvm.org/D68031#1682822>, @tra wrote:

> Nice.  I'd mention in the commit message that NVCC does not support deduced return type for kernel functions.


Just tried with NVCC from CUDA 10, except auto-based deduced type is not supported, type deduction in a template is supported, the following test code passes compilation with NVCC

  #include <cuda.h>
  
  template <typename T>
  __global__ T foo() {
  }
  
  void f0() {
    foo<void><<<0, 0>>>();
  #if 0
    foo<int><<<0, 0>>>();
  #endif
  }
  
  template <bool Cond, typename T = void> struct enable_if { typedef T type; };
  template <typename T> struct enable_if<false, T> {};
  
  template <int N>
  __global__
  auto bar() -> typename enable_if<N == 1>::type {
  }
  
  template <int N>
  __global__
  auto bar() -> typename enable_if<N == 2>::type {
  }
  
  void f3() {
    bar<1><<<0, 0>>>();
    bar<2><<<0, 0>>>();
  #if 0
    bar<3><<<0, 0>>>();
  #endif
  }

`s/#if 0/#if 1` also shows NVCC could give the error on the correct position but the message, IMHO, is misleading compared to the one from clang.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D68031