[PATCH] D154300: [CUDA][HIP] Fix template argument deduction

[Yaxun Liu via Phabricator via cfe-commits]

yaxunl created this revision.
yaxunl added reviewers: tra, rsmith.
Herald added subscribers: mattd, carlosgalvezp.
Herald added a project: All.
yaxunl requested review of this revision.

nvcc allows using std::malloc and std::free in device code.
When std::malloc or std::free is passed as a template
function argument with template argument deduction,
there is no diagnostics. e.g.

  #include <memory>
  
  __global__ void kern() {
      void *p = std::malloc(1);
      std::free(p);
  }
  int main()
  {
  
      std::shared_ptr<float> a;
      a = std::shared_ptr<float>(
        (float*)std::malloc(sizeof(float) * 100),
        std::free
      );
      return 0;
  }

However, the same code fails to compile with clang
(https://godbolt.org/z/1roGvo6YY). The reason is
that clang does not have logic to choose a function
argument from an overloaded set of candidates
based on host/device attributes for template argument
deduction.

Currently, clang does have a logic to choose a candidate
based on the constraints of the candidates. This patch
extends that logic to account for the CUDA host/device-based
preference.


https://reviews.llvm.org/D154300

Files:
  clang/lib/Sema/SemaOverload.cpp
  clang/test/SemaCUDA/template-arg-deduction.cu

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D154300.536586.patch
Type: text/x-patch
Size: 3889 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20230702/5d32cabf/attachment.bin>